1 // Check that when we link a bitcode module into a file using
2 // -mlink-builtin-bitcode, we apply the same attributes to the functions in that
3 // bitcode module as we apply to functions we generate.
5 // In particular, we check that ftz and unsafe-math are propagated into the
6 // bitcode library as appropriate.
8 // Build the bitcode library. This is not built in CUDA mode, otherwise it
9 // might have incompatible attributes. This mirrors how libdevice is built.
10 // RUN: %clang_cc1 -x c++ -fconvergent-functions -emit-llvm-bc -DLIB \
11 // RUN: %s -o %t.bc -triple nvptx-unknown-unknown
13 // RUN: %clang_cc1 -x cuda %s -emit-llvm -mlink-builtin-bitcode %t.bc -o - \
14 // RUN: -fcuda-is-device -triple nvptx-unknown-unknown \
15 // RUN: | FileCheck %s --check-prefix=CHECK --check-prefix=NOFTZ --check-prefix=NOFAST
17 // RUN: %clang_cc1 -x cuda %s -emit-llvm -mlink-builtin-bitcode %t.bc \
18 // RUN: -fdenormal-fp-math-f32=preserve-sign -o - \
19 // RUN: -fcuda-is-device -triple nvptx-unknown-unknown \
20 // RUN: | FileCheck %s --check-prefix=CHECK --check-prefix=FTZ \
21 // RUN: --check-prefix=NOFAST
23 // RUN: %clang_cc1 -x cuda %s -emit-llvm -mlink-builtin-bitcode %t.bc \
24 // RUN: -fdenormal-fp-math-f32=preserve-sign -o - \
25 // RUN: -fcuda-is-device -menable-unsafe-fp-math -triple nvptx-unknown-unknown \
26 // RUN: | FileCheck %s --check-prefix=CHECK --check-prefix=FAST
28 // Wrap everything in extern "C" so we don't have to worry about name mangling
33 // This function is defined in the library and only declared in the main
39 #include "Inputs/cuda.h"
40 __device__ void lib_fn();
41 __global__ void kernel() { lib_fn(); }
46 // The kernel and lib function should have the same attributes.
47 // CHECK: define{{.*}} void @kernel() [[kattr:#[0-9]+]]
48 // CHECK: define internal void @lib_fn() [[fattr:#[0-9]+]]
50 // FIXME: These -NOT checks do not work as intended and do not check on the same
53 // Check the attribute list for kernel.
54 // CHECK: attributes [[kattr]] = {
56 // CHECK-SAME: convergent
57 // CHECK-SAME: norecurse
59 // FTZ-NOT: "denormal-fp-math"
60 // FTZ-SAME: "denormal-fp-math-f32"="preserve-sign,preserve-sign"
61 // NOFTZ-NOT: "denormal-fp-math-f32"
63 // CHECK-SAME: "no-trapping-math"="true"
65 // FAST-SAME: "unsafe-fp-math"="true"
66 // NOFAST-NOT: "unsafe-fp-math"="true"
68 // Check the attribute list for lib_fn.
69 // CHECK: attributes [[fattr]] = {
71 // CHECK-SAME: convergent
72 // CHECK-NOT: norecurse
74 // FTZ-NOT: "denormal-fp-math"
75 // NOFTZ-NOT: "denormal-fp-math"
77 // FTZ-SAME: "denormal-fp-math-f32"="preserve-sign,preserve-sign"
78 // NOFTZ-NOT: "denormal-fp-math-f32"
80 // CHECK-SAME: "no-trapping-math"="true"
82 // FAST-SAME: "unsafe-fp-math"="true"
83 // NOFAST-NOT: "unsafe-fp-math"="true"