|
221
|
1 // REQUIRES: x86-registered-target
|
|
|
2 // REQUIRES: amdgpu-registered-target
|
|
|
3
|
|
252
|
4 // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device \
|
|
236
|
5 // RUN: -std=c++11 -fgpu-rdc -emit-llvm -o %t.nocuid.dev -x hip %s
|
|
|
6 // RUN: cat %t.nocuid.dev | FileCheck -check-prefixes=DEV,INT-DEV %s
|
|
221
|
7
|
|
252
|
8 // RUN: %clang_cc1 -triple x86_64-gnu-linux \
|
|
236
|
9 // RUN: -std=c++11 -fgpu-rdc -emit-llvm -o %t.nocuid.host -x hip %s
|
|
|
10 // RUN: cat %t.nocuid.host | FileCheck -check-prefixes=HOST,INT-HOST %s
|
|
221
|
11
|
|
252
|
12 // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -cuid=abc \
|
|
221
|
13 // RUN: -std=c++11 -fgpu-rdc -emit-llvm -o - -x hip %s > %t.dev
|
|
|
14 // RUN: cat %t.dev | FileCheck -check-prefixes=DEV,EXT-DEV %s
|
|
|
15
|
|
252
|
16 // RUN: %clang_cc1 -triple x86_64-gnu-linux -cuid=abc \
|
|
221
|
17 // RUN: -std=c++11 -fgpu-rdc -emit-llvm -o - -x hip %s > %t.host
|
|
|
18 // RUN: cat %t.host | FileCheck -check-prefixes=HOST,EXT-HOST %s
|
|
|
19
|
|
|
20 // Check host and device compilations use the same postfixes for static
|
|
|
21 // variable names.
|
|
|
22
|
|
|
23 // RUN: cat %t.dev %t.host | FileCheck -check-prefix=POSTFIX %s
|
|
236
|
24 // RUN: cat %t.nocuid.dev %t.nocuid.host | FileCheck -check-prefix=POSTFIX-ID %s
|
|
221
|
25
|
|
|
26 // Negative tests.
|
|
|
27
|
|
252
|
28 // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device \
|
|
221
|
29 // RUN: -std=c++11 -fgpu-rdc -emit-llvm -o - -x hip %s | FileCheck \
|
|
|
30 // RUN: -check-prefix=DEV-NEG %s
|
|
|
31
|
|
252
|
32 // RUN: %clang_cc1 -triple x86_64-gnu-linux \
|
|
221
|
33 // RUN: -std=c++11 -fgpu-rdc -emit-llvm -o - -x hip %s | FileCheck \
|
|
|
34 // RUN: -check-prefix=HOST-NEG %s
|
|
|
35
|
|
252
|
36 // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -cuid=abc \
|
|
221
|
37 // RUN: -std=c++11 -fgpu-rdc -emit-llvm -o - -x hip %s > %t.dev
|
|
|
38 // RUN: cat %t.dev | FileCheck -check-prefix=DEV-NEG %s
|
|
|
39
|
|
252
|
40 // RUN: %clang_cc1 -triple x86_64-gnu-linux -cuid=abc \
|
|
221
|
41 // RUN: -std=c++11 -fgpu-rdc -emit-llvm -o - -x hip %s > %t.host
|
|
|
42 // RUN: cat %t.host | FileCheck -check-prefix=HOST-NEG %s
|
|
|
43
|
|
236
|
44 // Check postfix for CUDA.
|
|
|
45
|
|
252
|
46 // RUN: %clang_cc1 -triple nvptx -fcuda-is-device -cuid=abc \
|
|
236
|
47 // RUN: -std=c++11 -fgpu-rdc -emit-llvm -o - %s | FileCheck \
|
|
|
48 // RUN: -check-prefixes=CUDA %s
|
|
221
|
49
|
|
|
50 #include "Inputs/cuda.h"
|
|
|
51
|
|
236
|
52 // Make sure we can still mangle with a line directive.
|
|
|
53 #line 0 "-"
|
|
|
54
|
|
221
|
55 // Test function scope static device variable, which should not be externalized.
|
|
|
56 // DEV-DAG: @_ZZ6kernelPiPPKiE1w = internal addrspace(4) constant i32 1
|
|
|
57
|
|
|
58
|
|
|
59 // HOST-DAG: @_ZL1x = internal global i32 undef
|
|
|
60 // HOST-DAG: @_ZL1y = internal global i32 undef
|
|
|
61
|
|
|
62 // Test normal static device variables
|
|
236
|
63 // INT-DEV-DAG: @_ZL1x[[FILEID:.*]] = addrspace(1) externally_initialized global i32 0
|
|
|
64 // INT-HOST-DAG: @[[DEVNAMEX:[0-9]+]] = {{.*}}c"_ZL1x[[FILEID:.*]]\00"
|
|
221
|
65
|
|
|
66 // Test externalized static device variables
|
|
|
67 // EXT-DEV-DAG: @_ZL1x.static.[[HASH:.*]] = addrspace(1) externally_initialized global i32 0
|
|
|
68 // EXT-HOST-DAG: @[[DEVNAMEX:[0-9]+]] = {{.*}}c"_ZL1x.static.[[HASH:.*]]\00"
|
|
236
|
69 // CUDA-DAG: @_ZL1x__static__[[HASH:.*]] = addrspace(1) externally_initialized global i32 0
|
|
221
|
70
|
|
|
71 // POSTFIX: @_ZL1x.static.[[HASH:.*]] = addrspace(1) externally_initialized global i32 0
|
|
|
72 // POSTFIX: @[[DEVNAMEX:[0-9]+]] = {{.*}}c"_ZL1x.static.[[HASH]]\00"
|
|
236
|
73 // POSTFIX-ID: @_ZL1x.static.[[FILEID:.*]] = addrspace(1) externally_initialized global i32 0
|
|
|
74 // POSTFIX-ID: @[[DEVNAMEX:[0-9]+]] = {{.*}}c"_ZL1x.static.[[FILEID]]\00"
|
|
221
|
75
|
|
|
76 static __device__ int x;
|
|
|
77
|
|
|
78 // Test static device variables not used by host code should not be externalized
|
|
|
79 // DEV-DAG: @_ZL2x2 = internal addrspace(1) global i32 0
|
|
|
80
|
|
|
81 static __device__ int x2;
|
|
|
82
|
|
|
83 // Test normal static device variables
|
|
236
|
84 // INT-DEV-DAG: @_ZL1y[[FILEID:.*]] = addrspace(4) externally_initialized global i32 0
|
|
|
85 // INT-HOST-DAG: @[[DEVNAMEY:[0-9]+]] = {{.*}}c"_ZL1y[[FILEID:.*]]\00"
|
|
221
|
86
|
|
|
87 // Test externalized static device variables
|
|
|
88 // EXT-DEV-DAG: @_ZL1y.static.[[HASH]] = addrspace(4) externally_initialized global i32 0
|
|
|
89 // EXT-HOST-DAG: @[[DEVNAMEY:[0-9]+]] = {{.*}}c"_ZL1y.static.[[HASH]]\00"
|
|
|
90
|
|
|
91 static __constant__ int y;
|
|
|
92
|
|
|
93 // Test static host variable, which should not be externalized nor registered.
|
|
|
94 // HOST-DAG: @_ZL1z = internal global i32 0
|
|
|
95 // DEV-NEG-NOT: @_ZL1z
|
|
|
96 static int z;
|
|
|
97
|
|
|
98 // Test non-ODR-use of static device variable is not emitted or registered.
|
|
|
99 // DEV-NEG-NOT: @_ZL1u
|
|
|
100 // HOST-NEG-NOT: @_ZL1u
|
|
|
101 static __device__ int u;
|
|
|
102
|
|
|
103 // Test static device variable in inline function, which should not be
|
|
|
104 // externalized nor registered.
|
|
|
105 // DEV-DAG: @_ZZ6devfunPPKiE1p = linkonce_odr addrspace(4) constant i32 2, comdat
|
|
|
106
|
|
|
107 inline __device__ void devfun(const int ** b) {
|
|
|
108 const static int p = 2;
|
|
|
109 b[0] = &p;
|
|
|
110 }
|
|
|
111
|
|
|
112 __global__ void kernel(int *a, const int **b) {
|
|
|
113 const static int w = 1;
|
|
|
114 a[0] = x;
|
|
|
115 a[1] = y;
|
|
|
116 a[2] = sizeof(u);
|
|
|
117 b[0] = &w;
|
|
|
118 b[1] = &x2;
|
|
|
119 devfun(b);
|
|
|
120 }
|
|
|
121
|
|
|
122 int* getDeviceSymbol(int *x);
|
|
|
123
|
|
|
124 void foo() {
|
|
|
125 getDeviceSymbol(&x);
|
|
|
126 getDeviceSymbol(&y);
|
|
|
127 z = 123;
|
|
|
128 decltype(u) tmp;
|
|
|
129 }
|
|
|
130
|
|
252
|
131 // HOST-DAG: __hipRegisterVar({{.*}}@_ZL1x, {{.*}}@[[DEVNAMEX]]
|
|
|
132 // HOST-DAG: __hipRegisterVar({{.*}}@_ZL1y, {{.*}}@[[DEVNAMEY]]
|
|
221
|
133 // HOST-NEG-NOT: __hipRegisterVar({{.*}}@_ZL2x2
|
|
|
134 // HOST-NEG-NOT: __hipRegisterVar({{.*}}@_ZZ6kernelPiPPKiE1w
|
|
|
135 // HOST-NEG-NOT: __hipRegisterVar({{.*}}@_ZZ6devfunPPKiE1p
|
|
|
136 // HOST-NEG-NOT: __hipRegisterVar({{.*}}@_ZL1u
|
