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