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