1 // REQUIRES: x86-registered-target
2 // REQUIRES: amdgpu-registered-target
4 // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device \
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
8 // RUN: %clang_cc1 -triple x86_64-gnu-linux \
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
12 // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -cuid=abc \
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
16 // RUN: %clang_cc1 -triple x86_64-gnu-linux -cuid=abc \
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
20 // Check host and device compilations use the same postfixes for static
23 // RUN: cat %t.dev %t.host | FileCheck -check-prefix=POSTFIX %s
24 // RUN: cat %t.nocuid.dev %t.nocuid.host | FileCheck -check-prefix=POSTFIX-ID %s
28 // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device \
29 // RUN: -std=c++11 -fgpu-rdc -emit-llvm -o - -x hip %s | FileCheck \
30 // RUN: -check-prefix=DEV-NEG %s
32 // RUN: %clang_cc1 -triple x86_64-gnu-linux \
33 // RUN: -std=c++11 -fgpu-rdc -emit-llvm -o - -x hip %s | FileCheck \
34 // RUN: -check-prefix=HOST-NEG %s
36 // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -cuid=abc \
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
40 // RUN: %clang_cc1 -triple x86_64-gnu-linux -cuid=abc \
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
44 // Check postfix for CUDA.
46 // RUN: %clang_cc1 -triple nvptx -fcuda-is-device -cuid=abc \
47 // RUN: -std=c++11 -fgpu-rdc -emit-llvm -o - %s | FileCheck \
48 // RUN: -check-prefixes=CUDA %s
50 #include "Inputs/cuda.h"
52 // Make sure we can still mangle with a line directive.
55 // Test function scope static device variable, which should not be externalized.
56 // DEV-DAG: @_ZZ6kernelPiPPKiE1w = internal addrspace(4) constant i32 1
59 // HOST-DAG: @_ZL1x = internal global i32 undef
60 // HOST-DAG: @_ZL1y = internal global i32 undef
62 // Test normal static device variables
63 // INT-DEV-DAG: @_ZL1x[[FILEID:.*]] = addrspace(1) externally_initialized global i32 0
64 // INT-HOST-DAG: @[[DEVNAMEX:[0-9]+]] = {{.*}}c"_ZL1x[[FILEID:.*]]\00"
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"
69 // CUDA-DAG: @_ZL1x__static__[[HASH:.*]] = addrspace(1) externally_initialized global i32 0
71 // POSTFIX: @_ZL1x.static.[[HASH:.*]] = addrspace(1) externally_initialized global i32 0
72 // POSTFIX: @[[DEVNAMEX:[0-9]+]] = {{.*}}c"_ZL1x.static.[[HASH]]\00"
73 // POSTFIX-ID: @_ZL1x.static.[[FILEID:.*]] = addrspace(1) externally_initialized global i32 0
74 // POSTFIX-ID: @[[DEVNAMEX:[0-9]+]] = {{.*}}c"_ZL1x.static.[[FILEID]]\00"
76 static __device__ int x;
78 // Test static device variables not used by host code should not be externalized
79 // DEV-DAG: @_ZL2x2 = internal addrspace(1) global i32 0
81 static __device__ int x2;
83 // Test normal static device variables
84 // INT-DEV-DAG: @_ZL1y[[FILEID:.*]] = addrspace(4) externally_initialized constant i32 0
85 // INT-HOST-DAG: @[[DEVNAMEY:[0-9]+]] = {{.*}}c"_ZL1y[[FILEID:.*]]\00"
87 // Test externalized static device variables
88 // EXT-DEV-DAG: @_ZL1y.static.[[HASH]] = addrspace(4) externally_initialized constant i32 0
89 // EXT-HOST-DAG: @[[DEVNAMEY:[0-9]+]] = {{.*}}c"_ZL1y.static.[[HASH]]\00"
91 static __constant__ int y;
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
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;
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
107 inline __device__ void devfun(const int ** b) {
108 const static int p = 2;
112 __global__ void kernel(int *a, const int **b) {
113 const static int w = 1;
122 int* getDeviceSymbol(int *x);
131 // HOST-DAG: __hipRegisterVar({{.*}}@_ZL1x, {{.*}}@[[DEVNAMEX]]
132 // HOST-DAG: __hipRegisterVar({{.*}}@_ZL1y, {{.*}}@[[DEVNAMEY]]
133 // HOST-NEG-NOT: __hipRegisterVar({{.*}}@_ZL2x2
134 // HOST-NEG-NOT: __hipRegisterVar({{.*}}@_ZZ6kernelPiPPKiE1w
135 // HOST-NEG-NOT: __hipRegisterVar({{.*}}@_ZZ6devfunPPKiE1p
136 // HOST-NEG-NOT: __hipRegisterVar({{.*}}@_ZL1u