Skip to content

Commit

Permalink
[OpenMP][Offloading] Change bool IsSPMD to int8_t Mode in `__kmpc…
Browse files Browse the repository at this point in the history 
…_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
  • Loading branch information
@shiltian
shiltian committed Sep 22, 2021
1 parent b875343 commit 423d34f
Show file tree
Hide file tree
Showing 58 changed files with 1,042 additions and 1,023 deletions.
4 changes: 2 additions & 2 deletions clang/test/OpenMP/amdgcn_target_codegen.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -13,7 +13,7 @@ int test_amdgcn_target_tid_threads() {

int arr[N];

// 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)
// 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)
#pragma omp target
for (int i = 0; i < N; i++) {
arr[i] = 1;
Expand All @@ -27,7 +27,7 @@ int test_amdgcn_target_tid_threads_simd() {

int arr[N];

// 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)
// 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)
#pragma omp target simd
for (int i = 0; i < N; i++) {
arr[i] = 1;
Expand Down
4 changes: 2 additions & 2 deletions clang/test/OpenMP/declare_target_codegen_globalization.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -31,7 +31,7 @@ int maini1() {
// CHECK1-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [1 x i8*], align 8
// CHECK1-NEXT: store i32* [[A]], i32** [[A_ADDR]], align 8
// CHECK1-NEXT: [[TMP0:%.*]] = load i32*, i32** [[A_ADDR]], align 8
// CHECK1-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1:[0-9]+]], i1 true, i1 false, i1 true)
// CHECK1-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1:[0-9]+]], i8 2, i1 false, i1 true)
// CHECK1-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1
// CHECK1-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
// CHECK1: user_code.entry:
Expand All @@ -41,7 +41,7 @@ int maini1() {
// CHECK1-NEXT: store i8* [[TMP4]], i8** [[TMP3]], align 8
// CHECK1-NEXT: [[TMP5:%.*]] = bitcast [1 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8**
// 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)
// CHECK1-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i1 true, i1 true)
// CHECK1-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i8 2, i1 true)
// CHECK1-NEXT: ret void
// CHECK1: worker.exit:
// CHECK1-NEXT: ret void
Expand Down
124 changes: 62 additions & 62 deletions clang/test/OpenMP/nvptx_SPMD_codegen.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -21,28 +21,28 @@ int a;
// CHECK-NOT: @__omp_offloading_{{.+}}_exec_mode = weak constant i8 1

void foo() {
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 false)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 false)
// CHECK-DAG: [[DISTR_LIGHT]]
// CHECK-DAG: [[FOR_LIGHT]]
// CHECK-DAG: [[LIGHT]]
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 false)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 false)
// CHECK-DAG: [[DISTR_LIGHT]]
// CHECK-DAG: [[FOR_LIGHT]]
// CHECK-DAG: [[LIGHT]]
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 false)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 false)
// CHECK-DAG: [[DISTR_LIGHT]]
// CHECK-DAG: [[FOR_LIGHT]]
// CHECK-DAG: [[LIGHT]]
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK-DAG: [[DISTR_FULL]]
// CHECK-DAG: [[FULL]]
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK-DAG: [[DISTR_FULL]]
// CHECK-DAG: [[FULL]]
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK-DAG: [[DISTR_FULL]]
// CHECK-DAG: [[FULL]]
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK-DAG: [[DISTR_FULL]]
// CHECK-DAG: [[FULL]]
#pragma omp target teams distribute parallel for simd if(a)
Expand All @@ -67,28 +67,28 @@ void foo() {
for (int i = 0; i < 10; ++i)
;
int a;
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 false)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 false)
// CHECK-DAG: [[DISTR_LIGHT]]
// CHECK-DAG: [[FOR_LIGHT]]
// CHECK-DAG: [[LIGHT]]
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 false)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 false)
// CHECK-DAG: [[DISTR_LIGHT]]
// CHECK-DAG: [[FOR_LIGHT]]
// CHECK-DAG: [[LIGHT]]
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 false)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 false)
// CHECK-DAG: [[DISTR_LIGHT]]
// CHECK-DAG: [[FOR_LIGHT]]
// CHECK-DAG: [[LIGHT]]
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK-DAG: [[DISTR_FULL]]
// CHECK-DAG: [[FULL]]
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK-DAG: [[DISTR_FULL]]
// CHECK-DAG: [[FULL]]
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK-DAG: [[DISTR_FULL]]
// CHECK-DAG: [[FULL]]
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK-DAG: [[DISTR_FULL]]
// CHECK-DAG: [[FULL]]
#pragma omp target teams distribute parallel for lastprivate(a)
Expand All @@ -112,25 +112,25 @@ int a;
#pragma omp target teams distribute parallel for schedule(guided)
for (int i = 0; i < 10; ++i)
;
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 false)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 false)
// CHECK-DAG: [[DISTR_LIGHT]]
// CHECK-DAG: [[FOR_LIGHT]]
// CHECK-DAG: [[LIGHT]]
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 false)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 false)
// CHECK-DAG: [[DISTR_LIGHT]]
// CHECK-DAG: [[FOR_LIGHT]]
// CHECK-DAG: [[LIGHT]]
// CHECK: call i32 @__kmpc_target_init(
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK-DAG: [[DISTR_FULL]]
// CHECK-DAG: [[FULL]]
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK-DAG: [[DISTR_FULL]]
// CHECK-DAG: [[FULL]]
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK-DAG: [[DISTR_FULL]]
// CHECK-DAG: [[FULL]]
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK-DAG: [[DISTR_FULL]]
// CHECK-DAG: [[FULL]]
#pragma omp target teams
Expand Down Expand Up @@ -172,28 +172,28 @@ int a;
#pragma omp distribute parallel for simd schedule(guided)
for (int i = 0; i < 10; ++i)
;
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 false)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 false)
// CHECK-DAG: [[DISTR_LIGHT]]
// CHECK-DAG: [[FOR_LIGHT]]
// CHECK-DAG: [[LIGHT]]
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 false)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 false)
// CHECK-DAG: [[DISTR_LIGHT]]
// CHECK-DAG: [[FOR_LIGHT]]
// CHECK-DAG: [[LIGHT]]
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 false)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 false)
// CHECK-DAG: [[DISTR_LIGHT]]
// CHECK-DAG: [[FOR_LIGHT]]
// CHECK-DAG: [[LIGHT]]
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK-DAG: [[DISTR_FULL]]
// CHECK-DAG: [[FULL]]
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK-DAG: [[DISTR_FULL]]
// CHECK-DAG: [[FULL]]
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK-DAG: [[DISTR_FULL]]
// CHECK-DAG: [[FULL]]
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK-DAG: [[DISTR_FULL]]
// CHECK-DAG: [[FULL]]
#pragma omp target teams
Expand Down Expand Up @@ -224,28 +224,28 @@ int a;
#pragma omp distribute parallel for schedule(guided)
for (int i = 0; i < 10; ++i)
;
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 false)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 false)
// CHECK-DAG: [[DISTR_LIGHT]]
// CHECK-DAG: [[FOR_LIGHT]]
// CHECK-DAG: [[LIGHT]]
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 false)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 false)
// CHECK-DAG: [[DISTR_LIGHT]]
// CHECK-DAG: [[FOR_LIGHT]]
// CHECK-DAG: [[LIGHT]]
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 false)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 false)
// CHECK-DAG: [[DISTR_LIGHT]]
// CHECK-DAG: [[FOR_LIGHT]]
// CHECK-DAG: [[LIGHT]]
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK-DAG: [[DISTR_FULL]]
// CHECK-DAG: [[FULL]]
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK-DAG: [[DISTR_FULL]]
// CHECK-DAG: [[FULL]]
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK-DAG: [[DISTR_FULL]]
// CHECK-DAG: [[FULL]]
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK-DAG: [[DISTR_FULL]]
// CHECK-DAG: [[FULL]]
#pragma omp target
Expand Down Expand Up @@ -283,22 +283,22 @@ int a;
#pragma omp distribute parallel for schedule(guided)
for (int i = 0; i < 10; ++i)
;
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 false)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 false)
// CHECK-DAG: [[FOR_LIGHT]]
// CHECK-DAG: [[LIGHT]]
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 false)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 false)
// CHECK-DAG: [[FOR_LIGHT]]
// CHECK-DAG: [[LIGHT]]
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 false)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 false)
// CHECK-DAG: [[FOR_LIGHT]]
// CHECK-DAG: [[LIGHT]]
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK-DAG: [[FULL]]
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK-DAG: [[FULL]]
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK-DAG: [[FULL]]
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK-DAG: [[FULL]]
#pragma omp target parallel for if(a)
for (int i = 0; i < 10; ++i)
Expand All @@ -321,28 +321,28 @@ int a;
#pragma omp target parallel for schedule(guided)
for (int i = 0; i < 10; ++i)
;
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 false)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 false)
// CHECK-DAG: [[FOR_LIGHT]]
// CHECK-DAG: [[LIGHT]]
// CHECK-DAG: [[BAR_LIGHT]]
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 false)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 false)
// CHECK-DAG: [[FOR_LIGHT]]
// CHECK-DAG: [[LIGHT]]
// CHECK-DAG: [[BAR_LIGHT]]
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 false)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 false)
// CHECK-DAG: [[FOR_LIGHT]]
// CHECK-DAG: [[LIGHT]]
// CHECK-DAG: [[BAR_LIGHT]]
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK-DAG: [[FULL]]
// CHECK-DAG: [[BAR_FULL]]
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK-DAG: [[FULL]]
// CHECK-DAG: [[BAR_FULL]]
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK-DAG: [[FULL]]
// CHECK-DAG: [[BAR_FULL]]
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK-DAG: [[FULL]]
// CHECK-DAG: [[BAR_FULL]]
#pragma omp target parallel if(a)
Expand Down Expand Up @@ -373,27 +373,27 @@ int a;
#pragma omp for simd schedule(guided)
for (int i = 0; i < 10; ++i)
;
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK-DAG: [[FULL]]
// CHECK-DAG: [[BAR_FULL]]
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 false)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 false)
// CHECK-DAG: [[FOR_LIGHT]]
// CHECK-DAG: [[LIGHT]]
// CHECK-DAG: [[BAR_LIGHT]]
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 false)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 false)
// CHECK-DAG: [[FOR_LIGHT]]
// CHECK-DAG: [[LIGHT]]
// CHECK-DAG: [[BAR_LIGHT]]
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK-DAG: [[FULL]]
// CHECK-DAG: [[BAR_FULL]]
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK-DAG: [[FULL]]
// CHECK-DAG: [[BAR_FULL]]
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK-DAG: [[FULL]]
// CHECK-DAG: [[BAR_FULL]]
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK-DAG: [[FULL]]
// CHECK-DAG: [[BAR_FULL]]
#pragma omp target
Expand Down Expand Up @@ -431,22 +431,22 @@ int a;
#pragma omp for simd schedule(guided)
for (int i = 0; i < 10; ++i)
;
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 false)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 false)
// CHECK-DAG: [[FOR_LIGHT]]
// CHECK-DAG: [[LIGHT]]
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 false)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 false)
// CHECK-DAG: [[FOR_LIGHT]]
// CHECK-DAG: [[LIGHT]]
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 false)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 false)
// CHECK-DAG: [[FOR_LIGHT]]
// CHECK-DAG: [[LIGHT]]
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK-DAG: [[FULL]]
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK-DAG: [[FULL]]
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK-DAG: [[FULL]]
// CHECK: call i32 @__kmpc_target_init({{.*}}, i1 true, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK-DAG: [[FULL]]
#pragma omp target
#pragma omp parallel for
Expand Down
4 changes: 2 additions & 2 deletions clang/test/OpenMP/nvptx_data_sharing.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -393,7 +393,7 @@ void test_ds(){
// CHECK-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [1 x i8*], align 8
// CHECK-NEXT: [[C:%.*]] = alloca i32, align 4
// CHECK-NEXT: [[CAPTURED_VARS_ADDRS1:%.*]] = alloca [2 x i8*], align 8
// CHECK-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1:[0-9]+]], i1 false, i1 true, i1 true)
// CHECK-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1:[0-9]+]], i8 1, i1 true, i1 true)
// CHECK-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1
// CHECK-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
// CHECK: user_code.entry:
Expand All @@ -420,7 +420,7 @@ void test_ds(){
// 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)
// CHECK-NEXT: call void @__kmpc_free_shared(i8* [[B]], i64 4)
// CHECK-NEXT: call void @__kmpc_free_shared(i8* [[A]], i64 4)
// CHECK-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i1 false, i1 true)
// CHECK-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i8 1, i1 true)
// CHECK-NEXT: ret void
// CHECK: worker.exit:
// CHECK-NEXT: ret void
Expand Down
Loading

0 comments on commit 423d34f

Please sign in to comment.