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