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 // In addition, we set -ftrapping-math on the bitcode library, but then set
9 // -fno-trapping-math on the main compilations, and ensure that the latter flag
10 // overrides the flag on the bitcode library.
11
12 // Build the bitcode library. This is not built in CUDA mode, otherwise it
13 // might have incompatible attributes. This mirrors how libdevice is built.
14 // RUN: %clang_cc1 -x c++ -fconvergent-functions -emit-llvm-bc -ftrapping-math -DLIB \
15 // RUN: %s -o %t.bc -triple nvptx-unknown-unknown
16
17 // RUN: %clang_cc1 -x cuda %s -emit-llvm -mlink-builtin-bitcode %t.bc -o - \
18 // RUN: -fno-trapping-math -fcuda-is-device -triple nvptx-unknown-unknown \
19 // RUN: | FileCheck %s --check-prefix=CHECK --check-prefix=NOFTZ --check-prefix=NOFAST
20
21 // RUN: %clang_cc1 -x cuda %s -emit-llvm -mlink-builtin-bitcode %t.bc \
22 // RUN: -fno-trapping-math -fdenormal-fp-math-f32=preserve-sign -o - \
23 // RUN: -fcuda-is-device -triple nvptx-unknown-unknown \
24 // RUN: | FileCheck %s --check-prefix=CHECK --check-prefix=FTZ \
25 // RUN: --check-prefix=NOFAST
26
27 // RUN: %clang_cc1 -x cuda %s -emit-llvm -mlink-builtin-bitcode %t.bc \
28 // RUN: -fno-trapping-math -fdenormal-fp-math-f32=preserve-sign -o - \
29 // RUN: -fcuda-is-device -menable-unsafe-fp-math -triple nvptx-unknown-unknown \
30 // RUN: | FileCheck %s --check-prefix=CHECK --check-prefix=FAST
31
32 // Wrap everything in extern "C" so we don't have to worry about name mangling
33 // in the IR.
34 extern "C" {
35 #ifdef LIB
36
37 // This function is defined in the library and only declared in the main
38 // compilation.
lib_fn()39 void lib_fn() {}
40
41 #else
42
43 #include "Inputs/cuda.h"
44 __device__ void lib_fn();
45 __global__ void kernel() { lib_fn(); }
46
47 #endif
48 }
49
50 // The kernel and lib function should have the same attributes.
51 // CHECK: define void @kernel() [[kattr:#[0-9]+]]
52 // CHECK: define internal void @lib_fn() [[fattr:#[0-9]+]]
53
54 // FIXME: These -NOT checks do not work as intended and do not check on the same
55 // line.
56
57 // Check the attribute list for kernel.
58 // CHECK: attributes [[kattr]] = {
59
60 // CHECK-SAME: convergent
61 // CHECK-SAME: norecurse
62
63 // FTZ-NOT: "denormal-fp-math"
64 // FTZ-SAME: "denormal-fp-math-f32"="preserve-sign,preserve-sign"
65 // NOFTZ-NOT: "denormal-fp-math-f32"
66
67 // CHECK-SAME: "no-trapping-math"="true"
68
69 // FAST-SAME: "unsafe-fp-math"="true"
70 // NOFAST-NOT: "unsafe-fp-math"="true"
71
72 // Check the attribute list for lib_fn.
73 // CHECK: attributes [[fattr]] = {
74
75 // CHECK-SAME: convergent
76 // CHECK-NOT: norecurse
77
78 // FTZ-NOT: "denormal-fp-math"
79 // NOFTZ-NOT: "denormal-fp-math"
80
81 // FTZ-SAME: "denormal-fp-math-f32"="preserve-sign,preserve-sign"
82 // NOFTZ-NOT: "denormal-fp-math-f32"
83
84 // CHECK-SAME: "no-trapping-math"="true"
85
86 // FAST-SAME: "unsafe-fp-math"="true"
87 // NOFAST-NOT: "unsafe-fp-math"="true"
88