1 // RUN: %clang_cc1 -fcuda-is-device \ 2 // RUN: -triple nvptx-nvidia-cuda -emit-llvm -o - %s | FileCheck %s -check-prefix NOFTZ 3 // RUN: %clang_cc1 -fcuda-is-device -fcuda-flush-denormals-to-zero \ 4 // RUN: -triple nvptx-nvidia-cuda -emit-llvm -o - %s | FileCheck %s -check-prefix FTZ 5 6 #include "Inputs/cuda.h" 7 8 // Checks that device function calls get emitted with the "ntpvx-f32ftz" 9 // attribute set to "true" when we compile CUDA device code with 10 // -fcuda-flush-denormals-to-zero. Further, check that we reflect the presence 11 // or absence of -fcuda-flush-denormals-to-zero in a module flag. 12 13 // CHECK: define void @foo() #0 14 extern "C" __device__ void foo() {} 15 16 // FTZ: attributes #0 = {{.*}} "nvptx-f32ftz"="true" 17 // NOFTZ-NOT: attributes #0 = {{.*}} "nvptx-f32ftz" 18 19 // FTZ:!llvm.module.flags = !{[[MODFLAG:![0-9]+]]} 20 // FTZ:[[MODFLAG]] = !{i32 4, !"nvvm-reflect-ftz", i32 1} 21 22 // NOFTZ:!llvm.module.flags = !{[[MODFLAG:![0-9]+]]} 23 // NOFTZ:[[MODFLAG]] = !{i32 4, !"nvvm-reflect-ftz", i32 0} 24