|
252
|
1 // RUN: %clang_cc1 -x hip -emit-llvm -std=c++11 %s -o - \
|
|
221
|
2 // RUN: -triple x86_64-linux-gnu \
|
|
|
3 // RUN: | FileCheck -check-prefix=HOST %s
|
|
252
|
4 // RUN: %clang_cc1 -x hip -emit-llvm -std=c++11 %s -o - \
|
|
221
|
5 // RUN: -triple amdgcn-amd-amdhsa -fcuda-is-device \
|
|
|
6 // RUN: | FileCheck -check-prefix=DEV %s
|
|
|
7
|
|
|
8 #include "Inputs/cuda.h"
|
|
|
9
|
|
252
|
10 // HOST: %[[T1:.*]] = type <{ ptr, i32, [4 x i8] }>
|
|
|
11 // HOST: %[[T2:.*]] = type { ptr, ptr }
|
|
|
12 // HOST: %[[T3:.*]] = type <{ ptr, i32, [4 x i8] }>
|
|
|
13 // DEV: %[[T1:.*]] = type { ptr }
|
|
|
14 // DEV: %[[T2:.*]] = type { ptr }
|
|
|
15 // DEV: %[[T3:.*]] = type <{ ptr, i32, [4 x i8] }>
|
|
221
|
16 int global_host_var;
|
|
|
17 __device__ int global_device_var;
|
|
|
18
|
|
|
19 template<class F>
|
|
|
20 __global__ void kern(F f) { f(); }
|
|
|
21
|
|
|
22 // DEV-LABEL: @_ZZ27dev_capture_dev_ref_by_copyPiENKUlvE_clEv(
|
|
252
|
23 // DEV: %[[VAL:.*]] = load i32, ptr addrspacecast (ptr addrspace(1) @global_device_var to ptr)
|
|
221
|
24 // DEV: store i32 %[[VAL]]
|
|
|
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
|
|
|
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(
|
|
252
|
40 // DEV: %[[VAL:.*]] = load i32, ptr addrspacecast (ptr addrspace(1) @global_device_var to ptr)
|
|
221
|
41 // DEV: %[[VAL2:.*]] = add nsw i32 %[[VAL]], 1
|
|
252
|
42 // DEV: store i32 %[[VAL2]], ptr addrspacecast (ptr addrspace(1) @global_device_var to ptr)
|
|
|
43 // DEV: %[[VAL:.*]] = load i32, ptr addrspacecast (ptr addrspace(1) @global_device_var to ptr)
|
|
221
|
44 // DEV: store i32 %[[VAL]]
|
|
|
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(
|
|
252
|
51 // DEV: %[[VAL:.*]] = load i32, ptr addrspacecast (ptr addrspace(1) @global_device_var to ptr)
|
|
221
|
52 // DEV: %[[VAL2:.*]] = add nsw i32 %[[VAL]], 1
|
|
252
|
53 // DEV: store i32 %[[VAL2]], ptr addrspacecast (ptr addrspace(1) @global_device_var to ptr)
|
|
|
54 // DEV: %[[VAL:.*]] = load i32, ptr addrspacecast (ptr addrspace(1) @global_device_var to ptr)
|
|
221
|
55 // DEV: store i32 %[[VAL]]
|
|
|
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(
|
|
252
|
63 // DEV: %[[VAL:.*]] = load i32, ptr addrspacecast (ptr addrspace(1) @global_device_var to ptr)
|
|
221
|
64 // DEV: %[[VAL2:.*]] = add nsw i32 %[[VAL]], 1
|
|
252
|
65 // DEV: store i32 %[[VAL2]], ptr addrspacecast (ptr addrspace(1) @global_device_var to ptr)
|
|
|
66 // DEV: %[[VAL:.*]] = load i32, ptr addrspacecast (ptr addrspace(1) @global_device_var to ptr)
|
|
221
|
67 // DEV: store i32 %[[VAL]]
|
|
|
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(
|
|
252
|
77 // HOST: %[[VAL:.*]] = load i32, ptr @global_host_var
|
|
221
|
78 // HOST: store i32 %[[VAL]]
|
|
|
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(
|
|
252
|
85 // HOST: %[[CAP:.*]] = getelementptr inbounds %[[T2]], ptr %this1, i32 0, i32 0
|
|
|
86 // HOST: %[[REF:.*]] = load ptr, ptr %[[CAP]]
|
|
|
87 // HOST: %[[VAL:.*]] = load i32, ptr %[[REF]]
|
|
221
|
88 // HOST: %[[VAL2:.*]] = add nsw i32 %[[VAL]], 1
|
|
252
|
89 // HOST: store i32 %[[VAL2]], ptr %[[REF]]
|
|
|
90 // HOST: %[[VAL:.*]] = load i32, ptr @global_host_var
|
|
221
|
91 // HOST: store i32 %[[VAL]]
|
|
|
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(
|
|
252
|
98 // HOST: %[[VAL:.*]] = load i32, ptr @global_host_var
|
|
221
|
99 // HOST: %[[VAL2:.*]] = add nsw i32 %[[VAL]], 1
|
|
252
|
100 // HOST: store i32 %[[VAL2]], ptr @global_host_var
|
|
|
101 // HOST: %[[VAL:.*]] = load i32, ptr @global_host_var
|
|
221
|
102 // HOST: store i32 %[[VAL]]
|
|
|
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(
|
|
252
|
110 // HOST: %[[VAL:.*]] = load i32, ptr @global_host_var
|
|
221
|
111 // HOST: %[[VAL2:.*]] = add nsw i32 %[[VAL]], 1
|
|
252
|
112 // HOST: store i32 %[[VAL2]], ptr @global_host_var
|
|
|
113 // HOST: %[[VAL:.*]] = load i32, ptr @global_host_var
|
|
221
|
114 // HOST: store i32 %[[VAL]]
|
|
|
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(
|
|
252
|
124 // HOST: %[[CAP:.*]] = getelementptr inbounds %[[T3]], ptr %{{.*}}, i32 0, i32 1
|
|
|
125 // HOST: %[[VAL:.*]] = load i32, ptr @global_host_var
|
|
|
126 // HOST: store i32 %[[VAL]], ptr %[[CAP]]
|
|
221
|
127 // DEV-LABEL: define internal void @_ZZ28dev_capture_host_ref_by_copyPiENKUlvE_clEv(
|
|
252
|
128 // DEV: %[[CAP:.*]] = getelementptr inbounds %[[T3]], ptr %this1, i32 0, i32 1
|
|
|
129 // DEV: %[[VAL:.*]] = load i32, ptr %[[CAP]]
|
|
221
|
130 // DEV: store i32 %[[VAL]]
|
|
|
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
|
