// RUN: %clang_cc1 -fcuda-is-device \
// RUN:   -triple nvptx-nvidia-cuda -emit-llvm -o - %s | \
// RUN:   FileCheck %s -check-prefix CHECK -check-prefix NOFTZ
// RUN: %clang_cc1 -fcuda-is-device -fcuda-flush-denormals-to-zero \
// RUN:   -triple nvptx-nvidia-cuda -emit-llvm -o - %s | \
// RUN:   FileCheck %s -check-prefix CHECK -check-prefix FTZ

// RUN: %clang_cc1 -fcuda-is-device -x hip \
// RUN:   -triple amdgcn-amd-amdhsa -target-cpu gfx900 -emit-llvm -o - %s | \
// RUN:   FileCheck %s -check-prefix CHECK -check-prefix AMDNOFTZ
// RUN: %clang_cc1 -fcuda-is-device -x hip -fcuda-flush-denormals-to-zero \
// RUN:   -triple amdgcn-amd-amdhsa -target-cpu gfx900 -emit-llvm -o - %s | \
// RUN:   FileCheck %s -check-prefix CHECK -check-prefix AMDFTZ

#include "Inputs/cuda.h"

// Checks that device function calls get emitted with the "ntpvx-f32ftz"
// attribute set to "true" when we compile CUDA device code with
// -fcuda-flush-denormals-to-zero.  Further, check that we reflect the presence
// or absence of -fcuda-flush-denormals-to-zero in a module flag.

// AMDGCN targets always have +fp64-fp16-denormals.
// AMDGCN targets without fast FMAF (e.g. gfx803) always have +fp32-denormals.
// For AMDGCN target with fast FMAF (e.g. gfx900), it has +fp32-denormals
// by default and -fp32-denormals when there is option
// -fcuda-flush-denormals-to-zero.

// CHECK-LABEL: define void @foo() #0
extern "C" __device__ void foo() {}

// FTZ: attributes #0 = {{.*}} "nvptx-f32ftz"="true"
// NOFTZ-NOT: attributes #0 = {{.*}} "nvptx-f32ftz"
// AMDNOFTZ: attributes #0 = {{.*}}+fp32-denormals{{.*}}+fp64-fp16-denormals
// AMDFTZ: attributes #0 = {{.*}}+fp64-fp16-denormals{{.*}}-fp32-denormals

// FTZ:!llvm.module.flags = !{{{.*}}[[MODFLAG:![0-9]+]]}
// FTZ:[[MODFLAG]] = !{i32 4, !"nvvm-reflect-ftz", i32 1}

// NOFTZ:!llvm.module.flags = !{{{.*}}[[MODFLAG:![0-9]+]]}
// NOFTZ:[[MODFLAG]] = !{i32 4, !"nvvm-reflect-ftz", i32 0}