Skip to content

Commit 423d34f

Browse files
shiltianshiltian
committed
[OpenMP][Offloading] Change bool IsSPMD to int8_t Mode in __kmpc_target_init and __kmpc_target_deinit
This is a follow-up of D110029, which uses bitset to indicate execution mode. This patches makes the changes in the function call. Reviewed By: jdoerfert Differential Revision: https://reviews.llvm.org/D110279
1 parent b875343 commit 423d34f

File tree

58 files changed

+1042
-1023
lines changed

Some content is hidden

Large Commits have some content hidden by default. Use the searchbox below for content that may be hidden.

58 files changed

+1042
-1023
lines changed

‎clang/test/OpenMP/amdgcn_target_codegen.cpp

+2-2
Original file line numberDiff line numberDiff line change
@@ -13,7 +13,7 @@ int test_amdgcn_target_tid_threads() {
1313

1414
int arr[N];
1515

16-
// CHECK: call i32 @__kmpc_target_init(%struct.ident_t* addrspacecast (%struct.ident_t addrspace(1)* @1 to %struct.ident_t*), i1 false, i1 true, i1 true)
16+
// CHECK: call i32 @__kmpc_target_init(%struct.ident_t* addrspacecast (%struct.ident_t addrspace(1)* @1 to %struct.ident_t*), i8 1, i1 true, i1 true)
1717
#pragma omp target
1818
for (int i = 0; i < N; i++) {
1919
arr[i] = 1;
@@ -27,7 +27,7 @@ int test_amdgcn_target_tid_threads_simd() {
2727

2828
int arr[N];
2929

30-
// CHECK: call i32 @__kmpc_target_init(%struct.ident_t* addrspacecast (%struct.ident_t addrspace(1)* @1 to %struct.ident_t*), i1 true, i1 false, i1 false)
30+
// CHECK: call i32 @__kmpc_target_init(%struct.ident_t* addrspacecast (%struct.ident_t addrspace(1)* @1 to %struct.ident_t*), i8 2, i1 false, i1 false)
3131
#pragma omp target simd
3232
for (int i = 0; i < N; i++) {
3333
arr[i] = 1;

‎clang/test/OpenMP/declare_target_codegen_globalization.cpp

+2-2
Original file line numberDiff line numberDiff line change
@@ -31,7 +31,7 @@ int maini1() {
3131
// CHECK1-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [1 x i8*], align 8
3232
// CHECK1-NEXT: store i32* [[A]], i32** [[A_ADDR]], align 8
3333
// CHECK1-NEXT: [[TMP0:%.*]] = load i32*, i32** [[A_ADDR]], align 8
34-
// CHECK1-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1:[0-9]+]], i1 true, i1 false, i1 true)
34+
// CHECK1-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1:[0-9]+]], i8 2, i1 false, i1 true)
3535
// CHECK1-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1
3636
// CHECK1-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
3737
// CHECK1: user_code.entry:
@@ -41,7 +41,7 @@ int maini1() {
4141
// CHECK1-NEXT: store i8* [[TMP4]], i8** [[TMP3]], align 8
4242
// CHECK1-NEXT: [[TMP5:%.*]] = bitcast [1 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8**
4343
// CHECK1-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* @[[GLOB2]], i32 [[TMP2]], i32 1, i32 -1, i32 -1, i8* bitcast (void (i32*, i32*, i32*)* @__omp_outlined__ to i8*), i8* null, i8** [[TMP5]], i64 1)
44-
// CHECK1-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i1 true, i1 true)
44+
// CHECK1-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i8 2, i1 true)
4545
// CHECK1-NEXT: ret void
4646
// CHECK1: worker.exit:
4747
// CHECK1-NEXT: ret void

‎clang/test/OpenMP/nvptx_SPMD_codegen.cpp

+62-62
Original file line numberDiff line numberDiff line change
@@ -21,28 +21,28 @@ int a;
2121
// CHECK-NOT: @__omp_offloading_{{.+}}_exec_mode = weak constant i8 1
2222

2323
void foo() {
24-
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 false)
24+
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 false)
2525
// CHECK-DAG: [[DISTR_LIGHT]]
2626
// CHECK-DAG: [[FOR_LIGHT]]
2727
// CHECK-DAG: [[LIGHT]]
28-
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 false)
28+
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 false)
2929
// CHECK-DAG: [[DISTR_LIGHT]]
3030
// CHECK-DAG: [[FOR_LIGHT]]
3131
// CHECK-DAG: [[LIGHT]]
32-
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 false)
32+
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 false)
3333
// CHECK-DAG: [[DISTR_LIGHT]]
3434
// CHECK-DAG: [[FOR_LIGHT]]
3535
// CHECK-DAG: [[LIGHT]]
36-
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
36+
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
3737
// CHECK-DAG: [[DISTR_FULL]]
3838
// CHECK-DAG: [[FULL]]
39-
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
39+
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
4040
// CHECK-DAG: [[DISTR_FULL]]
4141
// CHECK-DAG: [[FULL]]
42-
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
42+
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
4343
// CHECK-DAG: [[DISTR_FULL]]
4444
// CHECK-DAG: [[FULL]]
45-
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
45+
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
4646
// CHECK-DAG: [[DISTR_FULL]]
4747
// CHECK-DAG: [[FULL]]
4848
#pragma omp target teams distribute parallel for simd if(a)
@@ -67,28 +67,28 @@ void foo() {
6767
for (int i = 0; i < 10; ++i)
6868
;
6969
int a;
70-
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 false)
70+
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 false)
7171
// CHECK-DAG: [[DISTR_LIGHT]]
7272
// CHECK-DAG: [[FOR_LIGHT]]
7373
// CHECK-DAG: [[LIGHT]]
74-
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 false)
74+
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 false)
7575
// CHECK-DAG: [[DISTR_LIGHT]]
7676
// CHECK-DAG: [[FOR_LIGHT]]
7777
// CHECK-DAG: [[LIGHT]]
78-
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 false)
78+
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 false)
7979
// CHECK-DAG: [[DISTR_LIGHT]]
8080
// CHECK-DAG: [[FOR_LIGHT]]
8181
// CHECK-DAG: [[LIGHT]]
82-
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
82+
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
8383
// CHECK-DAG: [[DISTR_FULL]]
8484
// CHECK-DAG: [[FULL]]
85-
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
85+
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
8686
// CHECK-DAG: [[DISTR_FULL]]
8787
// CHECK-DAG: [[FULL]]
88-
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
88+
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
8989
// CHECK-DAG: [[DISTR_FULL]]
9090
// CHECK-DAG: [[FULL]]
91-
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
91+
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
9292
// CHECK-DAG: [[DISTR_FULL]]
9393
// CHECK-DAG: [[FULL]]
9494
#pragma omp target teams distribute parallel for lastprivate(a)
@@ -112,25 +112,25 @@ int a;
112112
#pragma omp target teams distribute parallel for schedule(guided)
113113
for (int i = 0; i < 10; ++i)
114114
;
115-
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 false)
115+
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 false)
116116
// CHECK-DAG: [[DISTR_LIGHT]]
117117
// CHECK-DAG: [[FOR_LIGHT]]
118118
// CHECK-DAG: [[LIGHT]]
119-
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 false)
119+
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 false)
120120
// CHECK-DAG: [[DISTR_LIGHT]]
121121
// CHECK-DAG: [[FOR_LIGHT]]
122122
// CHECK-DAG: [[LIGHT]]
123123
// CHECK: call i32 @__kmpc_target_init(
124-
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
124+
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
125125
// CHECK-DAG: [[DISTR_FULL]]
126126
// CHECK-DAG: [[FULL]]
127-
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
127+
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
128128
// CHECK-DAG: [[DISTR_FULL]]
129129
// CHECK-DAG: [[FULL]]
130-
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
130+
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
131131
// CHECK-DAG: [[DISTR_FULL]]
132132
// CHECK-DAG: [[FULL]]
133-
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
133+
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
134134
// CHECK-DAG: [[DISTR_FULL]]
135135
// CHECK-DAG: [[FULL]]
136136
#pragma omp target teams
@@ -172,28 +172,28 @@ int a;
172172
#pragma omp distribute parallel for simd schedule(guided)
173173
for (int i = 0; i < 10; ++i)
174174
;
175-
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 false)
175+
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 false)
176176
// CHECK-DAG: [[DISTR_LIGHT]]
177177
// CHECK-DAG: [[FOR_LIGHT]]
178178
// CHECK-DAG: [[LIGHT]]
179-
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 false)
179+
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 false)
180180
// CHECK-DAG: [[DISTR_LIGHT]]
181181
// CHECK-DAG: [[FOR_LIGHT]]
182182
// CHECK-DAG: [[LIGHT]]
183-
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 false)
183+
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 false)
184184
// CHECK-DAG: [[DISTR_LIGHT]]
185185
// CHECK-DAG: [[FOR_LIGHT]]
186186
// CHECK-DAG: [[LIGHT]]
187-
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
187+
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
188188
// CHECK-DAG: [[DISTR_FULL]]
189189
// CHECK-DAG: [[FULL]]
190-
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
190+
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
191191
// CHECK-DAG: [[DISTR_FULL]]
192192
// CHECK-DAG: [[FULL]]
193-
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
193+
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
194194
// CHECK-DAG: [[DISTR_FULL]]
195195
// CHECK-DAG: [[FULL]]
196-
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
196+
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
197197
// CHECK-DAG: [[DISTR_FULL]]
198198
// CHECK-DAG: [[FULL]]
199199
#pragma omp target teams
@@ -224,28 +224,28 @@ int a;
224224
#pragma omp distribute parallel for schedule(guided)
225225
for (int i = 0; i < 10; ++i)
226226
;
227-
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 false)
227+
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 false)
228228
// CHECK-DAG: [[DISTR_LIGHT]]
229229
// CHECK-DAG: [[FOR_LIGHT]]
230230
// CHECK-DAG: [[LIGHT]]
231-
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 false)
231+
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 false)
232232
// CHECK-DAG: [[DISTR_LIGHT]]
233233
// CHECK-DAG: [[FOR_LIGHT]]
234234
// CHECK-DAG: [[LIGHT]]
235-
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 false)
235+
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 false)
236236
// CHECK-DAG: [[DISTR_LIGHT]]
237237
// CHECK-DAG: [[FOR_LIGHT]]
238238
// CHECK-DAG: [[LIGHT]]
239-
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
239+
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
240240
// CHECK-DAG: [[DISTR_FULL]]
241241
// CHECK-DAG: [[FULL]]
242-
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
242+
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
243243
// CHECK-DAG: [[DISTR_FULL]]
244244
// CHECK-DAG: [[FULL]]
245-
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
245+
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
246246
// CHECK-DAG: [[DISTR_FULL]]
247247
// CHECK-DAG: [[FULL]]
248-
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
248+
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
249249
// CHECK-DAG: [[DISTR_FULL]]
250250
// CHECK-DAG: [[FULL]]
251251
#pragma omp target
@@ -283,22 +283,22 @@ int a;
283283
#pragma omp distribute parallel for schedule(guided)
284284
for (int i = 0; i < 10; ++i)
285285
;
286-
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 false)
286+
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 false)
287287
// CHECK-DAG: [[FOR_LIGHT]]
288288
// CHECK-DAG: [[LIGHT]]
289-
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 false)
289+
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 false)
290290
// CHECK-DAG: [[FOR_LIGHT]]
291291
// CHECK-DAG: [[LIGHT]]
292-
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 false)
292+
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 false)
293293
// CHECK-DAG: [[FOR_LIGHT]]
294294
// CHECK-DAG: [[LIGHT]]
295-
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
295+
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
296296
// CHECK-DAG: [[FULL]]
297-
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
297+
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
298298
// CHECK-DAG: [[FULL]]
299-
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
299+
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
300300
// CHECK-DAG: [[FULL]]
301-
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
301+
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
302302
// CHECK-DAG: [[FULL]]
303303
#pragma omp target parallel for if(a)
304304
for (int i = 0; i < 10; ++i)
@@ -321,28 +321,28 @@ int a;
321321
#pragma omp target parallel for schedule(guided)
322322
for (int i = 0; i < 10; ++i)
323323
;
324-
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 false)
324+
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 false)
325325
// CHECK-DAG: [[FOR_LIGHT]]
326326
// CHECK-DAG: [[LIGHT]]
327327
// CHECK-DAG: [[BAR_LIGHT]]
328-
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 false)
328+
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 false)
329329
// CHECK-DAG: [[FOR_LIGHT]]
330330
// CHECK-DAG: [[LIGHT]]
331331
// CHECK-DAG: [[BAR_LIGHT]]
332-
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 false)
332+
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 false)
333333
// CHECK-DAG: [[FOR_LIGHT]]
334334
// CHECK-DAG: [[LIGHT]]
335335
// CHECK-DAG: [[BAR_LIGHT]]
336-
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
336+
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
337337
// CHECK-DAG: [[FULL]]
338338
// CHECK-DAG: [[BAR_FULL]]
339-
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
339+
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
340340
// CHECK-DAG: [[FULL]]
341341
// CHECK-DAG: [[BAR_FULL]]
342-
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
342+
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
343343
// CHECK-DAG: [[FULL]]
344344
// CHECK-DAG: [[BAR_FULL]]
345-
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
345+
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
346346
// CHECK-DAG: [[FULL]]
347347
// CHECK-DAG: [[BAR_FULL]]
348348
#pragma omp target parallel if(a)
@@ -373,27 +373,27 @@ int a;
373373
#pragma omp for simd schedule(guided)
374374
for (int i = 0; i < 10; ++i)
375375
;
376-
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
376+
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
377377
// CHECK-DAG: [[FULL]]
378378
// CHECK-DAG: [[BAR_FULL]]
379-
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 false)
379+
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 false)
380380
// CHECK-DAG: [[FOR_LIGHT]]
381381
// CHECK-DAG: [[LIGHT]]
382382
// CHECK-DAG: [[BAR_LIGHT]]
383-
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 false)
383+
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 false)
384384
// CHECK-DAG: [[FOR_LIGHT]]
385385
// CHECK-DAG: [[LIGHT]]
386386
// CHECK-DAG: [[BAR_LIGHT]]
387-
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
387+
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
388388
// CHECK-DAG: [[FULL]]
389389
// CHECK-DAG: [[BAR_FULL]]
390-
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
390+
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
391391
// CHECK-DAG: [[FULL]]
392392
// CHECK-DAG: [[BAR_FULL]]
393-
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
393+
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
394394
// CHECK-DAG: [[FULL]]
395395
// CHECK-DAG: [[BAR_FULL]]
396-
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
396+
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
397397
// CHECK-DAG: [[FULL]]
398398
// CHECK-DAG: [[BAR_FULL]]
399399
#pragma omp target
@@ -431,22 +431,22 @@ int a;
431431
#pragma omp for simd schedule(guided)
432432
for (int i = 0; i < 10; ++i)
433433
;
434-
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 false)
434+
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 false)
435435
// CHECK-DAG: [[FOR_LIGHT]]
436436
// CHECK-DAG: [[LIGHT]]
437-
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 false)
437+
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 false)
438438
// CHECK-DAG: [[FOR_LIGHT]]
439439
// CHECK-DAG: [[LIGHT]]
440-
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 false)
440+
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 false)
441441
// CHECK-DAG: [[FOR_LIGHT]]
442442
// CHECK-DAG: [[LIGHT]]
443-
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
443+
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
444444
// CHECK-DAG: [[FULL]]
445-
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
445+
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
446446
// CHECK-DAG: [[FULL]]
447-
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
447+
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
448448
// CHECK-DAG: [[FULL]]
449-
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
449+
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
450450
// CHECK-DAG: [[FULL]]
451451
#pragma omp target
452452
#pragma omp parallel for

‎clang/test/OpenMP/nvptx_data_sharing.cpp

+2-2
Original file line numberDiff line numberDiff line change
@@ -393,7 +393,7 @@ void test_ds(){
393393
// CHECK-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [1 x i8*], align 8
394394
// CHECK-NEXT: [[C:%.*]] = alloca i32, align 4
395395
// CHECK-NEXT: [[CAPTURED_VARS_ADDRS1:%.*]] = alloca [2 x i8*], align 8
396-
// CHECK-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1:[0-9]+]], i1 false, i1 true, i1 true)
396+
// CHECK-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1:[0-9]+]], i8 1, i1 true, i1 true)
397397
// CHECK-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1
398398
// CHECK-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
399399
// CHECK: user_code.entry:
@@ -420,7 +420,7 @@ void test_ds(){
420420
// CHECK-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* @[[GLOB1]], i32 [[TMP1]], i32 1, i32 -1, i32 -1, i8* bitcast (void (i32*, i32*, i32*, i32*)* @__omp_outlined__1 to i8*), i8* bitcast (void (i16, i32)* @__omp_outlined__1_wrapper to i8*), i8** [[TMP9]], i64 2)
421421
// CHECK-NEXT: call void @__kmpc_free_shared(i8* [[B]], i64 4)
422422
// CHECK-NEXT: call void @__kmpc_free_shared(i8* [[A]], i64 4)
423-
// CHECK-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i1 false, i1 true)
423+
// CHECK-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i8 1, i1 true)
424424
// CHECK-NEXT: ret void
425425
// CHECK: worker.exit:
426426
// CHECK-NEXT: ret void

0 commit comments

Comments
 (0)