1 // RUN: %clang_cc1 -x hip -emit-llvm -std=c++11 %s -o - \
2 // RUN: -triple x86_64-linux-gnu \
3 // RUN: | FileCheck -check-prefix=HOST %s
4 // RUN: %clang_cc1 -x hip -emit-llvm -std=c++11 %s -o - \
5 // RUN: -triple amdgcn-amd-amdhsa -fcuda-is-device \
6 // RUN: | FileCheck -check-prefix=DEV %s
7
8 #include "Inputs/cuda.h"
9
10 // HOST: %[[T1:.*]] = type <{ i32*, i32, [4 x i8] }>
11 // HOST: %[[T2:.*]] = type { i32*, i32** }
12 // HOST: %[[T3:.*]] = type <{ i32*, i32, [4 x i8] }>
13 // DEV: %[[T1:.*]] = type { i32* }
14 // DEV: %[[T2:.*]] = type { i32** }
15 // DEV: %[[T3:.*]] = type <{ i32*, i32, [4 x i8] }>
16 int global_host_var;
17 __device__ int global_device_var;
18
19 template<class F>
kern(F f)20 __global__ void kern(F f) { f(); }
21
22 // DEV-LABEL: @_ZZ27dev_capture_dev_ref_by_copyPiENKUlvE_clEv(
23 // DEV: %[[VAL:.*]] = load i32, i32* addrspacecast (i32 addrspace(1)* @global_device_var to i32*)
24 // DEV: store i32 %[[VAL]]
dev_capture_dev_ref_by_copy(int * out)25 __device__ void dev_capture_dev_ref_by_copy(int *out) {
26 int &ref = global_device_var;
27 [=](){ *out = ref;}();
28 }
29
30 // DEV-LABEL: @_ZZ28dev_capture_dev_rval_by_copyPiENKUlvE_clEv(
31 // DEV: store i32 3
dev_capture_dev_rval_by_copy(int * out)32 __device__ void dev_capture_dev_rval_by_copy(int *out) {
33 constexpr int a = 1;
34 constexpr int b = 2;
35 constexpr int c = a + b;
36 [=](){ *out = c;}();
37 }
38
39 // DEV-LABEL: @_ZZ26dev_capture_dev_ref_by_refPiENKUlvE_clEv(
40 // DEV: %[[VAL:.*]] = load i32, i32* addrspacecast (i32 addrspace(1)* @global_device_var to i32*)
41 // DEV: %[[VAL2:.*]] = add nsw i32 %[[VAL]], 1
42 // DEV: store i32 %[[VAL2]], i32* addrspacecast (i32 addrspace(1)* @global_device_var to i32*)
43 // DEV: %[[VAL:.*]] = load i32, i32* addrspacecast (i32 addrspace(1)* @global_device_var to i32*)
44 // DEV: store i32 %[[VAL]]
dev_capture_dev_ref_by_ref(int * out)45 __device__ void dev_capture_dev_ref_by_ref(int *out) {
46 int &ref = global_device_var;
47 [&](){ ref++; *out = ref;}();
48 }
49
50 // DEV-LABEL: define void @_Z7dev_refPi(
51 // DEV: %[[VAL:.*]] = load i32, i32* addrspacecast (i32 addrspace(1)* @global_device_var to i32*)
52 // DEV: %[[VAL2:.*]] = add nsw i32 %[[VAL]], 1
53 // DEV: store i32 %[[VAL2]], i32* addrspacecast (i32 addrspace(1)* @global_device_var to i32*)
54 // DEV: %[[VAL:.*]] = load i32, i32* addrspacecast (i32 addrspace(1)* @global_device_var to i32*)
55 // DEV: store i32 %[[VAL]]
dev_ref(int * out)56 __device__ void dev_ref(int *out) {
57 int &ref = global_device_var;
58 ref++;
59 *out = ref;
60 }
61
62 // DEV-LABEL: @_ZZ14dev_lambda_refPiENKUlvE_clEv(
63 // DEV: %[[VAL:.*]] = load i32, i32* addrspacecast (i32 addrspace(1)* @global_device_var to i32*)
64 // DEV: %[[VAL2:.*]] = add nsw i32 %[[VAL]], 1
65 // DEV: store i32 %[[VAL2]], i32* addrspacecast (i32 addrspace(1)* @global_device_var to i32*)
66 // DEV: %[[VAL:.*]] = load i32, i32* addrspacecast (i32 addrspace(1)* @global_device_var to i32*)
67 // DEV: store i32 %[[VAL]]
dev_lambda_ref(int * out)68 __device__ void dev_lambda_ref(int *out) {
69 [=](){
70 int &ref = global_device_var;
71 ref++;
72 *out = ref;
73 }();
74 }
75
76 // HOST-LABEL: @_ZZ29host_capture_host_ref_by_copyPiENKUlvE_clEv(
77 // HOST: %[[VAL:.*]] = load i32, i32* @global_host_var
78 // HOST: store i32 %[[VAL]]
host_capture_host_ref_by_copy(int * out)79 void host_capture_host_ref_by_copy(int *out) {
80 int &ref = global_host_var;
81 [=](){ *out = ref;}();
82 }
83
84 // HOST-LABEL: @_ZZ28host_capture_host_ref_by_refPiENKUlvE_clEv(
85 // HOST: %[[CAP:.*]] = getelementptr inbounds %[[T2]], %[[T2]]* %this1, i32 0, i32 0
86 // HOST: %[[REF:.*]] = load i32*, i32** %[[CAP]]
87 // HOST: %[[VAL:.*]] = load i32, i32* %[[REF]]
88 // HOST: %[[VAL2:.*]] = add nsw i32 %[[VAL]], 1
89 // HOST: store i32 %[[VAL2]], i32* %[[REF]]
90 // HOST: %[[VAL:.*]] = load i32, i32* @global_host_var
91 // HOST: store i32 %[[VAL]]
host_capture_host_ref_by_ref(int * out)92 void host_capture_host_ref_by_ref(int *out) {
93 int &ref = global_host_var;
94 [&](){ ref++; *out = ref;}();
95 }
96
97 // HOST-LABEL: define void @_Z8host_refPi(
98 // HOST: %[[VAL:.*]] = load i32, i32* @global_host_var
99 // HOST: %[[VAL2:.*]] = add nsw i32 %[[VAL]], 1
100 // HOST: store i32 %[[VAL2]], i32* @global_host_var
101 // HOST: %[[VAL:.*]] = load i32, i32* @global_host_var
102 // HOST: store i32 %[[VAL]]
host_ref(int * out)103 void host_ref(int *out) {
104 int &ref = global_host_var;
105 ref++;
106 *out = ref;
107 }
108
109 // HOST-LABEL: @_ZZ15host_lambda_refPiENKUlvE_clEv(
110 // HOST: %[[VAL:.*]] = load i32, i32* @global_host_var
111 // HOST: %[[VAL2:.*]] = add nsw i32 %[[VAL]], 1
112 // HOST: store i32 %[[VAL2]], i32* @global_host_var
113 // HOST: %[[VAL:.*]] = load i32, i32* @global_host_var
114 // HOST: store i32 %[[VAL]]
host_lambda_ref(int * out)115 void host_lambda_ref(int *out) {
116 [=](){
117 int &ref = global_host_var;
118 ref++;
119 *out = ref;
120 }();
121 }
122
123 // HOST-LABEL: define void @_Z28dev_capture_host_ref_by_copyPi(
124 // HOST: %[[CAP:.*]] = getelementptr inbounds %[[T3]], %[[T3]]* %{{.*}}, i32 0, i32 1
125 // HOST: %[[VAL:.*]] = load i32, i32* @global_host_var
126 // HOST: store i32 %[[VAL]], i32* %[[CAP]]
127 // DEV-LABEL: define internal void @_ZZ28dev_capture_host_ref_by_copyPiENKUlvE_clEv(
128 // DEV: %[[CAP:.*]] = getelementptr inbounds %[[T3]], %[[T3]]* %this1, i32 0, i32 1
129 // DEV: %[[VAL:.*]] = load i32, i32* %[[CAP]]
130 // DEV: store i32 %[[VAL]]
dev_capture_host_ref_by_copy(int * out)131 void dev_capture_host_ref_by_copy(int *out) {
132 int &ref = global_host_var;
133 kern<<<1, 1>>>([=]__device__() { *out = ref;});
134 }
135
136