|
236
|
1 // RUN: %clang_cc1 -no-opaque-pointers -triple amdgcn-amd-amdhsa -fcuda-is-device -std=c++11 \
|
|
221
|
2 // RUN: -emit-llvm -o - -x hip %s | FileCheck \
|
|
|
3 // RUN: -check-prefixes=COMMON,DEV,NORDC-D %s
|
|
|
4
|
|
236
|
5 // RUN: %clang_cc1 -no-opaque-pointers -triple amdgcn-amd-amdhsa -fcuda-is-device -std=c++11 \
|
|
221
|
6 // RUN: -emit-llvm -fgpu-rdc -cuid=abc -o - -x hip %s > %t.dev
|
|
|
7 // RUN: cat %t.dev | FileCheck -check-prefixes=COMMON,DEV,RDC-D %s
|
|
|
8
|
|
236
|
9 // RUN: %clang_cc1 -no-opaque-pointers -triple x86_64-gnu-linux -std=c++11 \
|
|
221
|
10 // RUN: -emit-llvm -o - -x hip %s | FileCheck \
|
|
|
11 // RUN: -check-prefixes=COMMON,HOST,NORDC %s
|
|
|
12
|
|
236
|
13 // RUN: %clang_cc1 -no-opaque-pointers -triple x86_64-gnu-linux -std=c++11 \
|
|
221
|
14 // RUN: -emit-llvm -fgpu-rdc -cuid=abc -o - -x hip %s > %t.host
|
|
|
15 // RUN: cat %t.host | FileCheck -check-prefixes=COMMON,HOST,RDC %s
|
|
|
16
|
|
|
17 // Check device and host compilation use the same postfix for static
|
|
|
18 // variable name.
|
|
|
19
|
|
|
20 // RUN: cat %t.dev %t.host | FileCheck -check-prefix=POSTFIX %s
|
|
|
21
|
|
|
22 #include "Inputs/cuda.h"
|
|
|
23
|
|
|
24 struct vec {
|
|
|
25 float x,y,z;
|
|
|
26 };
|
|
|
27
|
|
|
28 // DEV-DAG: @x.managed = addrspace(1) externally_initialized global i32 1, align 4
|
|
|
29 // DEV-DAG: @x = addrspace(1) externally_initialized global i32 addrspace(1)* null
|
|
|
30 // NORDC-DAG: @x.managed = internal global i32 1
|
|
|
31 // RDC-DAG: @x.managed = global i32 1
|
|
|
32 // NORDC-DAG: @x = internal externally_initialized global i32* null
|
|
|
33 // RDC-DAG: @x = externally_initialized global i32* null
|
|
|
34 // HOST-DAG: @[[DEVNAMEX:[0-9]+]] = {{.*}}c"x\00"
|
|
|
35 __managed__ int x = 1;
|
|
|
36
|
|
|
37 // DEV-DAG: @v.managed = addrspace(1) externally_initialized global [100 x %struct.vec] zeroinitializer, align 4
|
|
|
38 // DEV-DAG: @v = addrspace(1) externally_initialized global [100 x %struct.vec] addrspace(1)* null
|
|
|
39 __managed__ vec v[100];
|
|
|
40
|
|
|
41 // DEV-DAG: @v2.managed = addrspace(1) externally_initialized global <{ %struct.vec, [99 x %struct.vec] }> <{ %struct.vec { float 1.000000e+00, float 1.000000e+00, float 1.000000e+00 }, [99 x %struct.vec] zeroinitializer }>, align 4
|
|
|
42 // DEV-DAG: @v2 = addrspace(1) externally_initialized global <{ %struct.vec, [99 x %struct.vec] }> addrspace(1)* null
|
|
|
43 __managed__ vec v2[100] = {{1, 1, 1}};
|
|
|
44
|
|
|
45 // DEV-DAG: @ex.managed = external addrspace(1) global i32, align 4
|
|
|
46 // DEV-DAG: @ex = external addrspace(1) externally_initialized global i32 addrspace(1)*
|
|
|
47 // HOST-DAG: @ex.managed = external global i32
|
|
|
48 // HOST-DAG: @ex = external externally_initialized global i32*
|
|
|
49 extern __managed__ int ex;
|
|
|
50
|
|
|
51 // NORDC-D-DAG: @_ZL2sx.managed = addrspace(1) externally_initialized global i32 1, align 4
|
|
|
52 // NORDC-D-DAG: @_ZL2sx = addrspace(1) externally_initialized global i32 addrspace(1)* null
|
|
|
53 // RDC-D-DAG: @_ZL2sx.static.[[HASH:.*]].managed = addrspace(1) externally_initialized global i32 1, align 4
|
|
|
54 // RDC-D-DAG: @_ZL2sx.static.[[HASH]] = addrspace(1) externally_initialized global i32 addrspace(1)* null
|
|
|
55 // HOST-DAG: @_ZL2sx.managed = internal global i32 1
|
|
|
56 // HOST-DAG: @_ZL2sx = internal externally_initialized global i32* null
|
|
|
57 // NORDC-DAG: @[[DEVNAMESX:[0-9]+]] = {{.*}}c"_ZL2sx\00"
|
|
|
58 // RDC-DAG: @[[DEVNAMESX:[0-9]+]] = {{.*}}c"_ZL2sx.static.[[HASH:.*]]\00"
|
|
|
59
|
|
|
60 // POSTFIX: @_ZL2sx.static.[[HASH:.*]] = addrspace(1) externally_initialized global i32 addrspace(1)* null
|
|
|
61 // POSTFIX: @[[DEVNAMESX:[0-9]+]] = {{.*}}c"_ZL2sx.static.[[HASH]]\00"
|
|
|
62 static __managed__ int sx = 1;
|
|
|
63
|
|
|
64 // DEV-DAG: @llvm.compiler.used
|
|
|
65 // DEV-SAME-DAG: @x.managed
|
|
|
66 // DEV-SAME-DAG: @x
|
|
|
67 // DEV-SAME-DAG: @v.managed
|
|
|
68 // DEV-SAME-DAG: @v
|
|
|
69 // DEV-SAME-DAG: @_ZL2sx.managed
|
|
|
70 // DEV-SAME-DAG: @_ZL2sx
|
|
|
71
|
|
|
72 // Force ex and sx mitted in device compilation.
|
|
|
73 __global__ void foo(int *z) {
|
|
|
74 *z = x + ex + sx;
|
|
|
75 v[1].x = 2;
|
|
|
76 }
|
|
|
77
|
|
|
78 // Force ex and sx emitted in host compilatioin.
|
|
|
79 int foo2() {
|
|
|
80 return ex + sx;
|
|
|
81 }
|
|
|
82
|
|
|
83 // COMMON-LABEL: define {{.*}}@_Z4loadv()
|
|
|
84 // DEV: %ld.managed = load i32 addrspace(1)*, i32 addrspace(1)* addrspace(1)* @x, align 4
|
|
|
85 // DEV: %0 = addrspacecast i32 addrspace(1)* %ld.managed to i32*
|
|
|
86 // DEV: %1 = load i32, i32* %0, align 4
|
|
|
87 // DEV: ret i32 %1
|
|
|
88 // HOST: %ld.managed = load i32*, i32** @x, align 4
|
|
|
89 // HOST: %0 = load i32, i32* %ld.managed, align 4
|
|
|
90 // HOST: ret i32 %0
|
|
|
91 __device__ __host__ int load() {
|
|
|
92 return x;
|
|
|
93 }
|
|
|
94
|
|
|
95 // COMMON-LABEL: define {{.*}}@_Z5storev()
|
|
|
96 // DEV: %ld.managed = load i32 addrspace(1)*, i32 addrspace(1)* addrspace(1)* @x, align 4
|
|
|
97 // DEV: %0 = addrspacecast i32 addrspace(1)* %ld.managed to i32*
|
|
|
98 // DEV: store i32 2, i32* %0, align 4
|
|
|
99 // HOST: %ld.managed = load i32*, i32** @x, align 4
|
|
|
100 // HOST: store i32 2, i32* %ld.managed, align 4
|
|
|
101 __device__ __host__ void store() {
|
|
|
102 x = 2;
|
|
|
103 }
|
|
|
104
|
|
|
105 // COMMON-LABEL: define {{.*}}@_Z10addr_takenv()
|
|
|
106 // DEV: %0 = addrspacecast i32 addrspace(1)* %ld.managed to i32*
|
|
|
107 // DEV: store i32* %0, i32** %p.ascast, align 8
|
|
|
108 // DEV: %1 = load i32*, i32** %p.ascast, align 8
|
|
|
109 // DEV: store i32 3, i32* %1, align 4
|
|
|
110 // HOST: %ld.managed = load i32*, i32** @x, align 4
|
|
|
111 // HOST: store i32* %ld.managed, i32** %p, align 8
|
|
|
112 // HOST: %0 = load i32*, i32** %p, align 8
|
|
|
113 // HOST: store i32 3, i32* %0, align 4
|
|
|
114 __device__ __host__ void addr_taken() {
|
|
|
115 int *p = &x;
|
|
|
116 *p = 3;
|
|
|
117 }
|
|
|
118
|
|
|
119 // HOST-LABEL: define {{.*}}@_Z5load2v()
|
|
|
120 // HOST: %ld.managed = load [100 x %struct.vec]*, [100 x %struct.vec]** @v, align 16
|
|
|
121 // HOST: %0 = getelementptr inbounds [100 x %struct.vec], [100 x %struct.vec]* %ld.managed, i64 0, i64 1, i32 0
|
|
|
122 // HOST: %1 = load float, float* %0, align 4
|
|
|
123 // HOST: ret float %1
|
|
|
124 __device__ __host__ float load2() {
|
|
|
125 return v[1].x;
|
|
|
126 }
|
|
|
127
|
|
|
128 // HOST-LABEL: define {{.*}}@_Z5load3v()
|
|
|
129 // HOST: %ld.managed = load <{ %struct.vec, [99 x %struct.vec] }>*, <{ %struct.vec, [99 x %struct.vec] }>** @v2, align 16
|
|
|
130 // HOST: %0 = bitcast <{ %struct.vec, [99 x %struct.vec] }>* %ld.managed to [100 x %struct.vec]*
|
|
|
131 // HOST: %1 = getelementptr inbounds [100 x %struct.vec], [100 x %struct.vec]* %0, i64 0, i64 1, i32 1
|
|
|
132 // HOST: %2 = load float, float* %1, align 4
|
|
|
133 // HOST: ret float %2
|
|
|
134 float load3() {
|
|
|
135 return v2[1].y;
|
|
|
136 }
|
|
|
137
|
|
|
138 // HOST-LABEL: define {{.*}}@_Z11addr_taken2v()
|
|
|
139 // HOST: %ld.managed = load [100 x %struct.vec]*, [100 x %struct.vec]** @v, align 16
|
|
|
140 // HOST: %0 = getelementptr inbounds [100 x %struct.vec], [100 x %struct.vec]* %ld.managed, i64 0, i64 1, i32 0
|
|
|
141 // HOST: %1 = ptrtoint float* %0 to i64
|
|
|
142 // HOST: %ld.managed1 = load <{ %struct.vec, [99 x %struct.vec] }>*, <{ %struct.vec, [99 x %struct.vec] }>** @v2, align 16
|
|
|
143 // HOST: %2 = bitcast <{ %struct.vec, [99 x %struct.vec] }>* %ld.managed1 to [100 x %struct.vec]*
|
|
|
144 // HOST: %3 = getelementptr inbounds [100 x %struct.vec], [100 x %struct.vec]* %2, i64 0, i64 1, i32 1
|
|
|
145 // HOST: %4 = ptrtoint float* %3 to i64
|
|
|
146 // HOST: %5 = sub i64 %4, %1
|
|
236
|
147 // HOST: %sub.ptr.div = sdiv exact i64 %5, 4
|
|
|
148 // HOST: %conv = sitofp i64 %sub.ptr.div to float
|
|
|
149 // HOST: ret float %conv
|
|
221
|
150 float addr_taken2() {
|
|
|
151 return (float)reinterpret_cast<long>(&(v2[1].y)-&(v[1].x));
|
|
|
152 }
|
|
|
153
|
|
|
154 // COMMON-LABEL: define {{.*}}@_Z5load4v()
|
|
|
155 // DEV: %ld.managed = load i32 addrspace(1)*, i32 addrspace(1)* addrspace(1)* @ex, align 4
|
|
|
156 // DEV: %0 = addrspacecast i32 addrspace(1)* %ld.managed to i32*
|
|
|
157 // DEV: %1 = load i32, i32* %0, align 4
|
|
|
158 // DEV: ret i32 %1
|
|
|
159 // HOST: %ld.managed = load i32*, i32** @ex, align 4
|
|
|
160 // HOST: %0 = load i32, i32* %ld.managed, align 4
|
|
|
161 // HOST: ret i32 %0
|
|
|
162 __device__ __host__ int load4() {
|
|
|
163 return ex;
|
|
|
164 }
|
|
|
165
|
|
|
166 // HOST-DAG: __hipRegisterManagedVar({{.*}}@x {{.*}}@x.managed {{.*}}@[[DEVNAMEX]]{{.*}}, i64 4, i32 4)
|
|
|
167 // HOST-DAG: __hipRegisterManagedVar({{.*}}@_ZL2sx {{.*}}@_ZL2sx.managed {{.*}}@[[DEVNAMESX]]
|
|
|
168 // HOST-NOT: __hipRegisterManagedVar({{.*}}@ex {{.*}}@ex.managed
|
|
|
169 // HOST-DAG: declare void @__hipRegisterManagedVar(i8**, i8*, i8*, i8*, i64, i32)
|
