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 // RUN: %clang_cc1 -fcuda-is-device -x hip \ 9 // RUN: -triple amdgcn-amd-amdhsa -target-cpu gfx900 -emit-llvm -o - %s | \ 10 // RUN: FileCheck %s -check-prefix CHECK -check-prefix AMDNOFTZ 11 // RUN: %clang_cc1 -fcuda-is-device -x hip -fcuda-flush-denormals-to-zero \ 12 // RUN: -triple amdgcn-amd-amdhsa -target-cpu gfx900 -emit-llvm -o - %s | \ 13 // RUN: FileCheck %s -check-prefix CHECK -check-prefix AMDFTZ 14 15 #include "Inputs/cuda.h" 16 17 // Checks that device function calls get emitted with the "ntpvx-f32ftz" 18 // attribute set to "true" when we compile CUDA device code with 19 // -fcuda-flush-denormals-to-zero. Further, check that we reflect the presence 20 // or absence of -fcuda-flush-denormals-to-zero in a module flag. 21 22 // AMDGCN targets always have +fp64-fp16-denormals. 23 // AMDGCN targets without fast FMAF (e.g. gfx803) always have +fp32-denormals. 24 // For AMDGCN target with fast FMAF (e.g. gfx900), it has +fp32-denormals 25 // by default and -fp32-denormals when there is option 26 // -fcuda-flush-denormals-to-zero. 27 28 // CHECK-LABEL: define void @foo() #0 29 extern "C" __device__ void foo() {} 30 31 // FTZ: attributes #0 = {{.*}} "nvptx-f32ftz"="true" 32 // NOFTZ-NOT: attributes #0 = {{.*}} "nvptx-f32ftz" 33 // AMDNOFTZ: attributes #0 = {{.*}}+fp32-denormals{{.*}}+fp64-fp16-denormals 34 // AMDFTZ: attributes #0 = {{.*}}+fp64-fp16-denormals{{.*}}-fp32-denormals 35 36 // FTZ:!llvm.module.flags = !{{{.*}}[[MODFLAG:![0-9]+]]} 37 // FTZ:[[MODFLAG]] = !{i32 4, !"nvvm-reflect-ftz", i32 1} 38 39 // NOFTZ:!llvm.module.flags = !{{{.*}}[[MODFLAG:![0-9]+]]} 40 // NOFTZ:[[MODFLAG]] = !{i32 4, !"nvvm-reflect-ftz", i32 0} 41