52 lines · plain
1// RUN: %clang_cc1 -fcuda-is-device \2// RUN: -triple nvptx-nvidia-cuda -emit-llvm -o - %s | \3// RUN: FileCheck -check-prefixes=NOFTZ,PTXNOFTZ %s4 5// RUN: %clang_cc1 -fcuda-is-device -fdenormal-fp-math-f32=ieee \6// RUN: -triple nvptx-nvidia-cuda -emit-llvm -o - %s | \7// RUN: FileCheck -check-prefixes=NOFTZ,PTXNOFTZ %s8 9// RUN: %clang_cc1 -fcuda-is-device -fdenormal-fp-math-f32=preserve-sign \10// RUN: -triple nvptx-nvidia-cuda -emit-llvm -o - %s | \11// RUN: FileCheck -check-prefixes=FTZ,PTXFTZ %s12 13// RUN: %clang_cc1 -fcuda-is-device -x hip \14// RUN: -triple amdgcn-amd-amdhsa -target-cpu gfx900 -emit-llvm -o - %s | \15// RUN: FileCheck -check-prefix=NOFTZ %s16 17// RUN: %clang_cc1 -fcuda-is-device -x hip \18// RUN: -triple amdgcn-amd-amdhsa -target-cpu gfx900 -fdenormal-fp-math-f32=ieee -emit-llvm -o - %s | \19// RUN: FileCheck -check-prefix=NOFTZ %s20 21// RUN: %clang_cc1 -fcuda-is-device -x hip -fdenormal-fp-math-f32=preserve-sign \22// RUN: -triple amdgcn-amd-amdhsa -target-cpu gfx900 -emit-llvm -o - %s | \23// RUN: FileCheck -check-prefix=FTZ %s24 25#include "Inputs/cuda.h"26 27// Checks that device function calls get emitted with the "denormal-fp-math-f32"28// attribute set when we compile CUDA device code with29// -fdenormal-fp-math-f32. Further, check that we reflect the presence or30// absence of -fgpu-flush-denormals-to-zero in a module flag.31 32// AMDGCN targets always have f64/f16 denormals enabled.33//34// AMDGCN targets without fast FMAF (e.g. gfx803) always have f32 denormal35// flushing by default.36//37// For AMDGCN target with fast FMAF (e.g. gfx900), it has ieee denormals by38// default and preserve-sign when there with the option39// -fgpu-flush-denormals-to-zero.40 41// CHECK-LABEL: define void @foo() #042extern "C" __device__ void foo() {}43 44// FTZ: attributes #0 = {{.*}} "denormal-fp-math-f32"="preserve-sign,preserve-sign"45// NOFTZ-NOT: "denormal-fp-math-f32"46 47// PTXFTZ:!llvm.module.flags = !{{{.*}}[[MODFLAG:![0-9]+]]}48// PTXFTZ:[[MODFLAG]] = !{i32 4, !"nvvm-reflect-ftz", i32 1}49 50// PTXNOFTZ:!llvm.module.flags = !{{{.*}}[[MODFLAG:![0-9]+]]}51// PTXNOFTZ:[[MODFLAG]] = !{i32 4, !"nvvm-reflect-ftz", i32 0}52