diff options
| author | Yaxun Liu <Yaxun.Liu@amd.com> | 2018-07-21 02:02:22 +0000 |
|---|---|---|
| committer | Yaxun Liu <Yaxun.Liu@amd.com> | 2018-07-21 02:02:22 +0000 |
| commit | e1bfbc589fea7e9e18e073b331f40869610cbf6c (patch) | |
| tree | 2a8c2d26a29c950de240b72be71a18ff077e215d /clang/test | |
| parent | 37f25a24bbf1fdb384c1e2559f033952f934a6e1 (diff) | |
| download | bcm5719-llvm-e1bfbc589fea7e9e18e073b331f40869610cbf6c.tar.gz bcm5719-llvm-e1bfbc589fea7e9e18e073b331f40869610cbf6c.zip | |
[HIP] Support -fcuda-flush-denormals-to-zero for amdgcn
Differential Revision: https://reviews.llvm.org/D48287
llvm-svn: 337639
Diffstat (limited to 'clang/test')
| -rw-r--r-- | clang/test/CodeGenCUDA/flush-denormals.cu | 15 |
1 files changed, 15 insertions, 0 deletions
diff --git a/clang/test/CodeGenCUDA/flush-denormals.cu b/clang/test/CodeGenCUDA/flush-denormals.cu index 94285f1f33c..13f2e3b2871 100644 --- a/clang/test/CodeGenCUDA/flush-denormals.cu +++ b/clang/test/CodeGenCUDA/flush-denormals.cu @@ -5,6 +5,13 @@ // RUN: -triple nvptx-nvidia-cuda -emit-llvm -o - %s | \ // RUN: FileCheck %s -check-prefix CHECK -check-prefix FTZ +// RUN: %clang_cc1 -fcuda-is-device -x hip \ +// RUN: -triple amdgcn-amd-amdhsa -target-cpu gfx900 -emit-llvm -o - %s | \ +// RUN: FileCheck %s -check-prefix CHECK -check-prefix AMDNOFTZ +// RUN: %clang_cc1 -fcuda-is-device -x hip -fcuda-flush-denormals-to-zero \ +// RUN: -triple amdgcn-amd-amdhsa -target-cpu gfx900 -emit-llvm -o - %s | \ +// RUN: FileCheck %s -check-prefix CHECK -check-prefix AMDFTZ + #include "Inputs/cuda.h" // Checks that device function calls get emitted with the "ntpvx-f32ftz" @@ -12,11 +19,19 @@ // -fcuda-flush-denormals-to-zero. Further, check that we reflect the presence // or absence of -fcuda-flush-denormals-to-zero in a module flag. +// AMDGCN targets always have +fp64-fp16-denormals. +// AMDGCN targets without fast FMAF (e.g. gfx803) always have +fp32-denormals. +// For AMDGCN target with fast FMAF (e.g. gfx900), it has +fp32-denormals +// by default and -fp32-denormals when there is option +// -fcuda-flush-denormals-to-zero. + // CHECK-LABEL: define void @foo() #0 extern "C" __device__ void foo() {} // FTZ: attributes #0 = {{.*}} "nvptx-f32ftz"="true" // NOFTZ-NOT: attributes #0 = {{.*}} "nvptx-f32ftz" +// AMDNOFTZ: attributes #0 = {{.*}}+fp32-denormals{{.*}}+fp64-fp16-denormals +// AMDFTZ: attributes #0 = {{.*}}+fp64-fp16-denormals{{.*}}-fp32-denormals // FTZ:!llvm.module.flags = !{{{.*}}[[MODFLAG:![0-9]+]]} // FTZ:[[MODFLAG]] = !{i32 4, !"nvvm-reflect-ftz", i32 1} |
