-
Notifications
You must be signed in to change notification settings - Fork 10.8k
Commit
This commit does not belong to any branch on this repository, and may belong to a fork outside of the repository.
[OpenMP][5.1] Fix parallel masked is ignored #59939
Code generation support for 'parallel masked' directive. The `EmitOMPParallelMaskedDirective` was implemented. In addition, the appropiate device functions were added. Fix #59939. Reviewed By: jdoerfert Differential Revision: https://reviews.llvm.org/D143527
- Loading branch information
1 parent
475dd6f
commit 64549f0
Showing
8 changed files
with
263 additions
and
3 deletions.
There are no files selected for viewing
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Original file line number | Diff line number | Diff line change |
---|---|---|
@@ -0,0 +1,109 @@ | ||
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --function-signature --include-generated-funcs --prefix-filecheck-ir-name _ | ||
// RUN: %clang_cc1 -triple x86_64-unknown-unknown -fopenmp -fopenmp-version=52 -x c -emit-llvm %s -o - | FileCheck %s | ||
// expected-no-diagnostics | ||
|
||
void foo(); | ||
|
||
void masked() { | ||
#pragma omp parallel masked | ||
{ | ||
foo(); | ||
} | ||
} | ||
|
||
void maskedFilter() { | ||
const int tid = 1; | ||
#pragma omp parallel masked filter(tid) | ||
{ | ||
foo(); | ||
} | ||
} | ||
|
||
void master() { | ||
#pragma omp parallel master | ||
{ | ||
foo(); | ||
} | ||
} | ||
// CHECK-LABEL: define {{[^@]+}}@masked | ||
// CHECK-SAME: () #[[ATTR0:[0-9]+]] { | ||
// CHECK-NEXT: entry: | ||
// CHECK-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1:[0-9]+]], i32 0, ptr @.omp_outlined.) | ||
// CHECK-NEXT: ret void | ||
// | ||
// | ||
// CHECK-LABEL: define {{[^@]+}}@.omp_outlined. | ||
// CHECK-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR1:[0-9]+]] { | ||
// CHECK-NEXT: entry: | ||
// CHECK-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8 | ||
// CHECK-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8 | ||
// CHECK-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8 | ||
// CHECK-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8 | ||
// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8 | ||
// CHECK-NEXT: [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 4 | ||
// CHECK-NEXT: [[TMP2:%.*]] = call i32 @__kmpc_masked(ptr @[[GLOB1]], i32 [[TMP1]], i32 0) | ||
// CHECK-NEXT: [[TMP3:%.*]] = icmp ne i32 [[TMP2]], 0 | ||
// CHECK-NEXT: br i1 [[TMP3]], label [[OMP_IF_THEN:%.*]], label [[OMP_IF_END:%.*]] | ||
// CHECK: omp_if.then: | ||
// CHECK-NEXT: call void (...) @foo() | ||
// CHECK-NEXT: call void @__kmpc_end_masked(ptr @[[GLOB1]], i32 [[TMP1]]) | ||
// CHECK-NEXT: br label [[OMP_IF_END]] | ||
// CHECK: omp_if.end: | ||
// CHECK-NEXT: ret void | ||
// | ||
// | ||
// CHECK-LABEL: define {{[^@]+}}@maskedFilter | ||
// CHECK-SAME: () #[[ATTR0]] { | ||
// CHECK-NEXT: entry: | ||
// CHECK-NEXT: [[TID:%.*]] = alloca i32, align 4 | ||
// CHECK-NEXT: store i32 1, ptr [[TID]], align 4 | ||
// CHECK-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 0, ptr @.omp_outlined..1) | ||
// CHECK-NEXT: ret void | ||
// | ||
// | ||
// CHECK-LABEL: define {{[^@]+}}@.omp_outlined..1 | ||
// CHECK-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR1]] { | ||
// CHECK-NEXT: entry: | ||
// CHECK-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8 | ||
// CHECK-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8 | ||
// CHECK-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8 | ||
// CHECK-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8 | ||
// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8 | ||
// CHECK-NEXT: [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 4 | ||
// CHECK-NEXT: [[TMP2:%.*]] = call i32 @__kmpc_masked(ptr @[[GLOB1]], i32 [[TMP1]], i32 1) | ||
// CHECK-NEXT: [[TMP3:%.*]] = icmp ne i32 [[TMP2]], 0 | ||
// CHECK-NEXT: br i1 [[TMP3]], label [[OMP_IF_THEN:%.*]], label [[OMP_IF_END:%.*]] | ||
// CHECK: omp_if.then: | ||
// CHECK-NEXT: call void (...) @foo() | ||
// CHECK-NEXT: call void @__kmpc_end_masked(ptr @[[GLOB1]], i32 [[TMP1]]) | ||
// CHECK-NEXT: br label [[OMP_IF_END]] | ||
// CHECK: omp_if.end: | ||
// CHECK-NEXT: ret void | ||
// | ||
// | ||
// CHECK-LABEL: define {{[^@]+}}@master | ||
// CHECK-SAME: () #[[ATTR0]] { | ||
// CHECK-NEXT: entry: | ||
// CHECK-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 0, ptr @.omp_outlined..2) | ||
// CHECK-NEXT: ret void | ||
// | ||
// | ||
// CHECK-LABEL: define {{[^@]+}}@.omp_outlined..2 | ||
// CHECK-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR1]] { | ||
// CHECK-NEXT: entry: | ||
// CHECK-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8 | ||
// CHECK-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8 | ||
// CHECK-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8 | ||
// CHECK-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8 | ||
// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8 | ||
// CHECK-NEXT: [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 4 | ||
// CHECK-NEXT: [[TMP2:%.*]] = call i32 @__kmpc_master(ptr @[[GLOB1]], i32 [[TMP1]]) | ||
// CHECK-NEXT: [[TMP3:%.*]] = icmp ne i32 [[TMP2]], 0 | ||
// CHECK-NEXT: br i1 [[TMP3]], label [[OMP_IF_THEN:%.*]], label [[OMP_IF_END:%.*]] | ||
// CHECK: omp_if.then: | ||
// CHECK-NEXT: call void (...) @foo() | ||
// CHECK-NEXT: call void @__kmpc_end_master(ptr @[[GLOB1]], i32 [[TMP1]]) | ||
// CHECK-NEXT: br label [[OMP_IF_END]] | ||
// CHECK: omp_if.end: | ||
// CHECK-NEXT: ret void | ||
// |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Original file line number | Diff line number | Diff line change |
---|---|---|
@@ -0,0 +1,112 @@ | ||
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --function-signature --include-generated-funcs --prefix-filecheck-ir-name _ | ||
// RUN: %clang_cc1 -triple x86_64-unknown-unknown -fopenmp -fopenmp-version=52 -fopenmp-targets=nvptx64 -offload-device-only -x c -emit-llvm %s -o - | FileCheck %s | ||
// expected-no-diagnostics | ||
|
||
void foo(); | ||
|
||
void masked() { | ||
#pragma target | ||
#pragma omp parallel masked | ||
{ | ||
foo(); | ||
} | ||
} | ||
|
||
void maskedFilter() { | ||
const int tid = 1; | ||
#pragma target | ||
#pragma omp parallel masked filter(tid) | ||
{ | ||
foo(); | ||
} | ||
} | ||
|
||
void master() { | ||
#pragma target | ||
#pragma omp parallel master | ||
{ | ||
foo(); | ||
} | ||
} | ||
// CHECK-LABEL: define {{[^@]+}}@masked | ||
// CHECK-SAME: () #[[ATTR0:[0-9]+]] { | ||
// CHECK-NEXT: entry: | ||
// CHECK-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1:[0-9]+]], i32 0, ptr @.omp_outlined.) | ||
// CHECK-NEXT: ret void | ||
// | ||
// | ||
// CHECK-LABEL: define {{[^@]+}}@.omp_outlined. | ||
// CHECK-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR1:[0-9]+]] { | ||
// CHECK-NEXT: entry: | ||
// CHECK-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8 | ||
// CHECK-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8 | ||
// CHECK-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8 | ||
// CHECK-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8 | ||
// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8 | ||
// CHECK-NEXT: [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 4 | ||
// CHECK-NEXT: [[TMP2:%.*]] = call i32 @__kmpc_masked(ptr @[[GLOB1]], i32 [[TMP1]], i32 0) | ||
// CHECK-NEXT: [[TMP3:%.*]] = icmp ne i32 [[TMP2]], 0 | ||
// CHECK-NEXT: br i1 [[TMP3]], label [[OMP_IF_THEN:%.*]], label [[OMP_IF_END:%.*]] | ||
// CHECK: omp_if.then: | ||
// CHECK-NEXT: call void (...) @foo() | ||
// CHECK-NEXT: call void @__kmpc_end_masked(ptr @[[GLOB1]], i32 [[TMP1]]) | ||
// CHECK-NEXT: br label [[OMP_IF_END]] | ||
// CHECK: omp_if.end: | ||
// CHECK-NEXT: ret void | ||
// | ||
// | ||
// CHECK-LABEL: define {{[^@]+}}@maskedFilter | ||
// CHECK-SAME: () #[[ATTR0]] { | ||
// CHECK-NEXT: entry: | ||
// CHECK-NEXT: [[TID:%.*]] = alloca i32, align 4 | ||
// CHECK-NEXT: store i32 1, ptr [[TID]], align 4 | ||
// CHECK-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 0, ptr @.omp_outlined..1) | ||
// CHECK-NEXT: ret void | ||
// | ||
// | ||
// CHECK-LABEL: define {{[^@]+}}@.omp_outlined..1 | ||
// CHECK-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR1]] { | ||
// CHECK-NEXT: entry: | ||
// CHECK-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8 | ||
// CHECK-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8 | ||
// CHECK-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8 | ||
// CHECK-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8 | ||
// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8 | ||
// CHECK-NEXT: [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 4 | ||
// CHECK-NEXT: [[TMP2:%.*]] = call i32 @__kmpc_masked(ptr @[[GLOB1]], i32 [[TMP1]], i32 1) | ||
// CHECK-NEXT: [[TMP3:%.*]] = icmp ne i32 [[TMP2]], 0 | ||
// CHECK-NEXT: br i1 [[TMP3]], label [[OMP_IF_THEN:%.*]], label [[OMP_IF_END:%.*]] | ||
// CHECK: omp_if.then: | ||
// CHECK-NEXT: call void (...) @foo() | ||
// CHECK-NEXT: call void @__kmpc_end_masked(ptr @[[GLOB1]], i32 [[TMP1]]) | ||
// CHECK-NEXT: br label [[OMP_IF_END]] | ||
// CHECK: omp_if.end: | ||
// CHECK-NEXT: ret void | ||
// | ||
// | ||
// CHECK-LABEL: define {{[^@]+}}@master | ||
// CHECK-SAME: () #[[ATTR0]] { | ||
// CHECK-NEXT: entry: | ||
// CHECK-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 0, ptr @.omp_outlined..2) | ||
// CHECK-NEXT: ret void | ||
// | ||
// | ||
// CHECK-LABEL: define {{[^@]+}}@.omp_outlined..2 | ||
// CHECK-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR1]] { | ||
// CHECK-NEXT: entry: | ||
// CHECK-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8 | ||
// CHECK-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8 | ||
// CHECK-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8 | ||
// CHECK-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8 | ||
// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8 | ||
// CHECK-NEXT: [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 4 | ||
// CHECK-NEXT: [[TMP2:%.*]] = call i32 @__kmpc_master(ptr @[[GLOB1]], i32 [[TMP1]]) | ||
// CHECK-NEXT: [[TMP3:%.*]] = icmp ne i32 [[TMP2]], 0 | ||
// CHECK-NEXT: br i1 [[TMP3]], label [[OMP_IF_THEN:%.*]], label [[OMP_IF_END:%.*]] | ||
// CHECK: omp_if.then: | ||
// CHECK-NEXT: call void (...) @foo() | ||
// CHECK-NEXT: call void @__kmpc_end_master(ptr @[[GLOB1]], i32 [[TMP1]]) | ||
// CHECK-NEXT: br label [[OMP_IF_END]] | ||
// CHECK: omp_if.end: | ||
// CHECK-NEXT: ret void | ||
// |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters