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.
4 //
5 // In particular, we check that ftz and unsafe-math are propagated into the
6 // bitcode library as appropriate.
7
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
12
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
16
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
22
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 -funsafe-math-optimizations -triple nvptx-unknown-unknown \
26 // RUN: | FileCheck %s --check-prefix=CHECK --check-prefix=FAST
27
28 // Wrap everything in extern "C" so we don't have to worry about name mangling
29 // in the IR.
30 extern "C" {
31 #ifdef LIB
32
33 // This function is defined in the library and only declared in the main
34 // compilation.
lib_fn()35 void lib_fn() {}
36
37 #else
38
39 #include "Inputs/cuda.h"
40 __device__ void lib_fn();
41 __global__ void kernel() { lib_fn(); }
42
43 #endif
44 }
45
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]+]]
49
50 // FIXME: These -NOT checks do not work as intended and do not check on the same
51 // line.
52
53 // Check the attribute list for kernel.
54 // CHECK: attributes [[kattr]] = {
55
56 // CHECK-SAME: convergent
57 // CHECK-SAME: norecurse
58
59 // FTZ-NOT: "denormal-fp-math"
60 // FTZ-SAME: "denormal-fp-math-f32"="preserve-sign,preserve-sign"
61 // NOFTZ-NOT: "denormal-fp-math-f32"
62
63 // CHECK-SAME: "no-trapping-math"="true"
64
65 // FAST-SAME: "unsafe-fp-math"="true"
66 // NOFAST-NOT: "unsafe-fp-math"="true"
67
68 // Check the attribute list for lib_fn.
69 // CHECK: attributes [[fattr]] = {
70
71 // CHECK-SAME: convergent
72 // CHECK-NOT: norecurse
73
74 // FTZ-NOT: "denormal-fp-math"
75 // NOFTZ-NOT: "denormal-fp-math"
76
77 // FTZ-SAME: "denormal-fp-math-f32"="preserve-sign,preserve-sign"
78 // NOFTZ-NOT: "denormal-fp-math-f32"
79
80 // CHECK-SAME: "no-trapping-math"="true"
81
82 // FAST-SAME: "unsafe-fp-math"="true"
83 // NOFAST-NOT: "unsafe-fp-math"="true"
84