diff options
author | Justin Lebar <jlebar@google.com> | 2016-04-05 18:26:20 +0000 |
---|---|---|
committer | Justin Lebar <jlebar@google.com> | 2016-04-05 18:26:20 +0000 |
commit | 92b0ff91a2faa573f171d41835303114a24d6729 (patch) | |
tree | 942570bf189f00086b58bd8307b6b4118dc8d2c8 /test/CodeGenCUDA | |
parent | 1abb3092f1f962ad6c3380f9da6e916faf90c53c (diff) |
[CUDA] Add -fcuda-flush-denormals-to-zero.
Summary:
Setting this flag causes all functions are annotated with the
"nvvm-f32ftz" = "true" attribute.
In addition, we annotate the module with "nvvm-reflect-ftz" set
to 0 or 1, depending on whether -cuda-flush-denormals-to-zero is set.
This is read by the NVVMReflect pass.
Reviewers: tra, rnk
Subscribers: cfe-commits
Differential Revision: http://reviews.llvm.org/D18671
git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@265435 91177308-0d34-0410-b5e6-96231b3b80d8
Diffstat (limited to 'test/CodeGenCUDA')
-rw-r--r-- | test/CodeGenCUDA/flush-denormals.cu | 23 |
1 files changed, 23 insertions, 0 deletions
diff --git a/test/CodeGenCUDA/flush-denormals.cu b/test/CodeGenCUDA/flush-denormals.cu new file mode 100644 index 0000000000..cab660254d --- /dev/null +++ b/test/CodeGenCUDA/flush-denormals.cu @@ -0,0 +1,23 @@ +// RUN: %clang_cc1 -fcuda-is-device \ +// RUN: -triple nvptx-nvidia-cuda -emit-llvm -o - %s | FileCheck %s -check-prefix NOFTZ +// RUN: %clang_cc1 -fcuda-is-device -fcuda-flush-denormals-to-zero \ +// RUN: -triple nvptx-nvidia-cuda -emit-llvm -o - %s | FileCheck %s -check-prefix FTZ + +#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. + +// CHECK: define void @foo() #0 +extern "C" __device__ void foo() {} + +// FTZ: attributes #0 = {{.*}} "nvptx-f32ftz"="true" +// NOFTZ-NOT: attributes #0 = {{.*}} "nvptx-f32ftz" + +// 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} |