[clang][OpenMP] 6.0: Add default clause support for 'target' directive #162910
+1,605
−15
Conversation
This file contains hidden or 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
Per OpenMP 6.0 specification, section 7.5.1, default Clause Page 224, lines 3-5 default Clause, Semantics If data-sharing-attribute is shared then the clause has no effect on a target construct; otherwise, its effect on a target construct is equivalent to specifying the defaultmap clause with the same data-sharing-attribute and variable-category. Testing: OpenMP LIT tests check-all
llvmbot
added
clang
|
@llvm/pr-subscribers-clang Author: David Pagan (ddpagan) ChangesPer OpenMP 6.0 specification, section 7.5.1, default Clause Page 224, lines 3-5 default Clause, Semantics Testing: Patch is 112.44 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/162910.diff 4 Files Affected:
diff --git a/clang/docs/ReleaseNotes.rst b/clang/docs/ReleaseNotes.rst
index 65b086caf3652..15e50d39a706b 100644
--- a/clang/docs/ReleaseNotes.rst
+++ b/clang/docs/ReleaseNotes.rst
@@ -635,6 +635,7 @@ OpenMP Support
- Added support for ``defaultmap`` directive implicit-behavior ``private``.
- Added parsing and semantic analysis support for ``groupprivate`` directive.
- Added support for 'omp fuse' directive.
+- Added support for ``default`` clause on ``target`` directive.
Improvements
^^^^^^^^^^^^
diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp
index 0fa21e89b1236..3416ff2632db6 100644
--- a/clang/lib/Sema/SemaOpenMP.cpp
+++ b/clang/lib/Sema/SemaOpenMP.cpp
@@ -17297,6 +17297,43 @@ OMPClause *SemaOpenMP::ActOnOpenMPSimpleClause(
return Res;
}
+static std::pair<OpenMPDefaultmapClauseModifier, OpenMPDefaultmapClauseKind>
+getDefaultmapModifierAndKind(llvm::omp::DefaultKind M,
+ OpenMPDefaultClauseVariableCategory VCKind) {
+ OpenMPDefaultmapClauseModifier DefMapMod;
+ OpenMPDefaultmapClauseKind DefMapKind;
+ switch (M) {
+ case OMP_DEFAULT_none:
+ DefMapMod = OMPC_DEFAULTMAP_MODIFIER_none;
+ break;
+ case OMP_DEFAULT_firstprivate:
+ DefMapMod = OMPC_DEFAULTMAP_MODIFIER_firstprivate;
+ break;
+ case OMP_DEFAULT_private:
+ DefMapMod = OMPC_DEFAULTMAP_MODIFIER_private;
+ break;
+ default:
+ llvm_unreachable("unexpected DSA in OpenMP default clause");
+ }
+ switch (VCKind) {
+ case OMPC_DEFAULT_VC_aggregate:
+ DefMapKind = OMPC_DEFAULTMAP_aggregate;
+ break;
+ case OMPC_DEFAULT_VC_pointer:
+ DefMapKind = OMPC_DEFAULTMAP_pointer;
+ break;
+ case OMPC_DEFAULT_VC_scalar:
+ DefMapKind = OMPC_DEFAULTMAP_scalar;
+ break;
+ case OMPC_DEFAULT_VC_all:
+ DefMapKind = OMPC_DEFAULTMAP_all;
+ break;
+ default:
+ llvm_unreachable("unexpected variable category in OpenMP default clause");
+ }
+ return std::make_pair(DefMapMod, DefMapKind);
+}
+
OMPClause *SemaOpenMP::ActOnOpenMPDefaultClause(
llvm::omp::DefaultKind M, SourceLocation MLoc,
OpenMPDefaultClauseVariableCategory VCKind, SourceLocation VCKindLoc,
@@ -17309,21 +17346,40 @@ OMPClause *SemaOpenMP::ActOnOpenMPDefaultClause(
return nullptr;
}
- switch (M) {
- case OMP_DEFAULT_none:
- DSAStack->setDefaultDSANone(MLoc);
- break;
- case OMP_DEFAULT_shared:
- DSAStack->setDefaultDSAShared(MLoc);
- break;
- case OMP_DEFAULT_firstprivate:
- DSAStack->setDefaultDSAFirstPrivate(MLoc);
- break;
- case OMP_DEFAULT_private:
- DSAStack->setDefaultDSAPrivate(MLoc);
- break;
- default:
- llvm_unreachable("DSA unexpected in OpenMP default clause");
+ if (getLangOpts().OpenMP >= 60 &&
+ DSAStack->getCurrentDirective() == OMPD_target) {
+ // OpenMP 6.0 (see page 224, lines 3-5) default Clause, Semantics
+ // If data-sharing-attribute is shared then the clause has no effect
+ // on a target construct; otherwise, its effect on a target construct is
+ // equivalent to specifying the defaultmap clause with the same
+ // data-sharing-attribute and variable-category.
+ if (M != OMP_DEFAULT_shared) {
+ auto [DefMapMod, DefMapKind] = getDefaultmapModifierAndKind(M, VCKind);
+ if (DefMapKind == OMPC_DEFAULTMAP_all) {
+ DSAStack->setDefaultDMAAttr(DefMapMod, OMPC_DEFAULTMAP_aggregate, MLoc);
+ DSAStack->setDefaultDMAAttr(DefMapMod, OMPC_DEFAULTMAP_scalar, MLoc);
+ DSAStack->setDefaultDMAAttr(DefMapMod, OMPC_DEFAULTMAP_pointer, MLoc);
+ } else {
+ DSAStack->setDefaultDMAAttr(DefMapMod, DefMapKind, MLoc);
+ }
+ }
+ } else {
+ switch (M) {
+ case OMP_DEFAULT_none:
+ DSAStack->setDefaultDSANone(MLoc);
+ break;
+ case OMP_DEFAULT_shared:
+ DSAStack->setDefaultDSAShared(MLoc);
+ break;
+ case OMP_DEFAULT_firstprivate:
+ DSAStack->setDefaultDSAFirstPrivate(MLoc);
+ break;
+ case OMP_DEFAULT_private:
+ DSAStack->setDefaultDSAPrivate(MLoc);
+ break;
+ default:
+ llvm_unreachable("DSA unexpected in OpenMP default clause");
+ }
}
switch (VCKind) {
diff --git a/clang/test/OpenMP/target_default_codegen.cpp b/clang/test/OpenMP/target_default_codegen.cpp
new file mode 100644
index 0000000000000..a3ee569e343ef
--- /dev/null
+++ b/clang/test/OpenMP/target_default_codegen.cpp
@@ -0,0 +1,1531 @@
+// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --include-generated-funcs --replace-value-regex "__omp_offloading_[0-9a-z]+_[0-9a-z]+" "reduction_size[.].+[.]" "pl_cond[.].+[.|,]" --prefix-filecheck-ir-name _ --version 5
+// expected-no-diagnostics
+
+// RUN: %clang_cc1 -no-enable-noundef-analysis -verify -Wno-vla -fopenmp -fopenmp-version=60 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK-64
+// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-version=60 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-version=60 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK-64
+// RUN: %clang_cc1 -no-enable-noundef-analysis -verify -Wno-vla -fopenmp -fopenmp-version=60 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK-32
+// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-version=60 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-version=60 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK-32
+
+// RUN: %clang_cc1 -no-enable-noundef-analysis -verify -Wno-vla -fopenmp-simd -fopenmp-version=60 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY-64 %s
+// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-version=60 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-version=60 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY-64 %s
+// RUN: %clang_cc1 -no-enable-noundef-analysis -verify -Wno-vla -fopenmp-simd -fopenmp-version=60 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY-32 %s
+// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-version=60 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-version=60 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY-32 %s
+
+#ifndef HEADER
+#define HEADER
+void foo1(int a) {
+ double d = (double)a;
+
+ #pragma omp target default(private: scalar)
+ {
+ d += 1.0;
+ }
+}
+
+void foo2() {
+ int pvtArr[10];
+
+ #pragma omp target default(private: aggregate)
+ {
+ pvtArr[5]++;
+ }
+}
+
+void foo3() {
+ int *pa;
+
+ #pragma omp target default(private: pointer)
+ {
+ pa[50]++;
+ }
+}
+
+// Specified variable-category doesn't apply to referenced variable, so
+// normal implicitly determined data-sharing applies.
+void foo4() {
+ int p;
+
+ #pragma omp target default(private: pointer)
+ {
+ p++;
+ }
+}
+
+// Verify default clause with variable-category 'all' is equivalent to no
+// variable-category. IR checks generated with 'all' but test runs without
+// variable-category.
+void foo5(int a) {
+ double d = (double)a;
+ int pvtArr[10];
+ int *pa;
+
+ #pragma omp target default(private)
+ {
+ d += 1.0;
+ pvtArr[5]++;
+ pa[50]++;
+ }
+}
+
+// Verify default clause with 'shared' DSA is ignored. This makes it
+// equivalent to target with no default clause. IR checks generated with
+// no default clause but test runs with default 'shared'.
+void foo6(int a) {
+ double d = (double)a;
+ int pvtArr[10];
+ int *pa;
+
+ #pragma omp target default(shared)
+ {
+ d += 1.0;
+ pvtArr[5]++;
+ pa[50]++;
+ }
+}
+
+// Verify default clause with 'firstprivate' DSA is equivalent to specifying
+// defaultmap with 'firstprivate'. IR checks generated with
+// defaultmap(firstprivate) but test runs with default(firstprivate).
+void foo7(int a) {
+ double d = (double)a;
+ int pvtArr[10];
+ int *pa;
+
+ #pragma omp target default(firstprivate)
+ {
+ d += 1.0;
+ pvtArr[5]++;
+ pa[50]++;
+ }
+}
+#endif // HEADER
+// CK-64-LABEL: define dso_local void @_Z4foo1i(
+// CK-64-SAME: i32 signext [[A:%.*]]) #[[ATTR0:[0-9]+]] {
+// CK-64-NEXT: [[ENTRY:.*:]]
+// CK-64-NEXT: [[A_ADDR:%.*]] = alloca i32, align 4
+// CK-64-NEXT: [[D:%.*]] = alloca double, align 8
+// CK-64-NEXT: [[D_CASTED:%.*]] = alloca i64, align 8
+// CK-64-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 8
+// CK-64-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 8
+// CK-64-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 8
+// CK-64-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8
+// CK-64-NEXT: store i32 [[A]], ptr [[A_ADDR]], align 4
+// CK-64-NEXT: [[TMP0:%.*]] = load i32, ptr [[A_ADDR]], align 4
+// CK-64-NEXT: [[CONV:%.*]] = sitofp i32 [[TMP0]] to double
+// CK-64-NEXT: store double [[CONV]], ptr [[D]], align 8
+// CK-64-NEXT: [[TMP1:%.*]] = load double, ptr [[D]], align 8
+// CK-64-NEXT: store double [[TMP1]], ptr [[D_CASTED]], align 8
+// CK-64-NEXT: [[TMP2:%.*]] = load i64, ptr [[D_CASTED]], align 8
+// CK-64-NEXT: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
+// CK-64-NEXT: store i64 [[TMP2]], ptr [[TMP3]], align 8
+// CK-64-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
+// CK-64-NEXT: store i64 [[TMP2]], ptr [[TMP4]], align 8
+// CK-64-NEXT: [[TMP5:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0
+// CK-64-NEXT: store ptr null, ptr [[TMP5]], align 8
+// CK-64-NEXT: [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
+// CK-64-NEXT: [[TMP7:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
+// CK-64-NEXT: [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0
+// CK-64-NEXT: store i32 3, ptr [[TMP8]], align 4
+// CK-64-NEXT: [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1
+// CK-64-NEXT: store i32 1, ptr [[TMP9]], align 4
+// CK-64-NEXT: [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2
+// CK-64-NEXT: store ptr [[TMP6]], ptr [[TMP10]], align 8
+// CK-64-NEXT: [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3
+// CK-64-NEXT: store ptr [[TMP7]], ptr [[TMP11]], align 8
+// CK-64-NEXT: [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4
+// CK-64-NEXT: store ptr @.offload_sizes, ptr [[TMP12]], align 8
+// CK-64-NEXT: [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5
+// CK-64-NEXT: store ptr @.offload_maptypes, ptr [[TMP13]], align 8
+// CK-64-NEXT: [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6
+// CK-64-NEXT: store ptr null, ptr [[TMP14]], align 8
+// CK-64-NEXT: [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7
+// CK-64-NEXT: store ptr null, ptr [[TMP15]], align 8
+// CK-64-NEXT: [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8
+// CK-64-NEXT: store i64 0, ptr [[TMP16]], align 8
+// CK-64-NEXT: [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9
+// CK-64-NEXT: store i64 0, ptr [[TMP17]], align 8
+// CK-64-NEXT: [[TMP18:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10
+// CK-64-NEXT: store [3 x i32] [i32 -1, i32 0, i32 0], ptr [[TMP18]], align 4
+// CK-64-NEXT: [[TMP19:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11
+// CK-64-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP19]], align 4
+// CK-64-NEXT: [[TMP20:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12
+// CK-64-NEXT: store i32 0, ptr [[TMP20]], align 4
+// CK-64-NEXT: [[TMP21:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1:[0-9]+]], i64 -1, i32 -1, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4foo1i_l23.region_id, ptr [[KERNEL_ARGS]])
+// CK-64-NEXT: [[TMP22:%.*]] = icmp ne i32 [[TMP21]], 0
+// CK-64-NEXT: br i1 [[TMP22]], label %[[OMP_OFFLOAD_FAILED:.*]], label %[[OMP_OFFLOAD_CONT:.*]]
+// CK-64: [[OMP_OFFLOAD_FAILED]]:
+// CK-64-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4foo1i_l23(i64 [[TMP2]]) #[[ATTR2:[0-9]+]]
+// CK-64-NEXT: br label %[[OMP_OFFLOAD_CONT]]
+// CK-64: [[OMP_OFFLOAD_CONT]]:
+// CK-64-NEXT: ret void
+//
+//
+// CK-64-LABEL: define internal void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4foo1i_l23(
+// CK-64-SAME: i64 [[D:%.*]]) #[[ATTR1:[0-9]+]] {
+// CK-64-NEXT: [[ENTRY:.*:]]
+// CK-64-NEXT: [[D_ADDR:%.*]] = alloca i64, align 8
+// CK-64-NEXT: [[D1:%.*]] = alloca double, align 8
+// CK-64-NEXT: store i64 [[D]], ptr [[D_ADDR]], align 8
+// CK-64-NEXT: [[TMP0:%.*]] = load double, ptr [[D1]], align 8
+// CK-64-NEXT: [[ADD:%.*]] = fadd double [[TMP0]], 1.000000e+00
+// CK-64-NEXT: store double [[ADD]], ptr [[D1]], align 8
+// CK-64-NEXT: ret void
+//
+//
+// CK-64-LABEL: define dso_local void @_Z4foo2v(
+// CK-64-SAME: ) #[[ATTR0]] {
+// CK-64-NEXT: [[ENTRY:.*:]]
+// CK-64-NEXT: [[PVTARR:%.*]] = alloca [10 x i32], align 4
+// CK-64-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 8
+// CK-64-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 8
+// CK-64-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 8
+// CK-64-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8
+// CK-64-NEXT: [[TMP0:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
+// CK-64-NEXT: store ptr [[PVTARR]], ptr [[TMP0]], align 8
+// CK-64-NEXT: [[TMP1:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
+// CK-64-NEXT: store ptr [[PVTARR]], ptr [[TMP1]], align 8
+// CK-64-NEXT: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0
+// CK-64-NEXT: store ptr null, ptr [[TMP2]], align 8
+// CK-64-NEXT: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
+// CK-64-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
+// CK-64-NEXT: [[TMP5:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0
+// CK-64-NEXT: store i32 3, ptr [[TMP5]], align 4
+// CK-64-NEXT: [[TMP6:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1
+// CK-64-NEXT: store i32 1, ptr [[TMP6]], align 4
+// CK-64-NEXT: [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2
+// CK-64-NEXT: store ptr [[TMP3]], ptr [[TMP7]], align 8
+// CK-64-NEXT: [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3
+// CK-64-NEXT: store ptr [[TMP4]], ptr [[TMP8]], align 8
+// CK-64-NEXT: [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4
+// CK-64-NEXT: store ptr @.offload_sizes.1, ptr [[TMP9]], align 8
+// CK-64-NEXT: [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5
+// CK-64-NEXT: store ptr @.offload_maptypes.2, ptr [[TMP10]], align 8
+// CK-64-NEXT: [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6
+// CK-64-NEXT: store ptr null, ptr [[TMP11]], align 8
+// CK-64-NEXT: [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7
+// CK-64-NEXT: store ptr null, ptr [[TMP12]], align 8
+// CK-64-NEXT: [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8
+// CK-64-NEXT: store i64 0, ptr [[TMP13]], align 8
+// CK-64-NEXT: [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9
+// CK-64-NEXT: store i64 0, ptr [[TMP14]], align 8
+// CK-64-NEXT: [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10
+// CK-64-NEXT: store [3 x i32] [i32 -1, i32 0, i32 0], ptr [[TMP15]], align 4
+// CK-64-NEXT: [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11
+// CK-64-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP16]], align 4
+// CK-64-NEXT: [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12
+// CK-64-NEXT: store i32 0, ptr [[TMP17]], align 4
+// CK-64-NEXT: [[TMP18:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 -1, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4foo2v_l32.region_id, ptr [[KERNEL_ARGS]])
+// CK-64-NEXT: [[TMP19:%.*]] = icmp ne i32 [[TMP18]], 0
+// CK-64-NEXT: br i1 [[TMP19]], label %[[OMP_OFFLOAD_FAILED:.*]], label %[[OMP_OFFLOAD_CONT:.*]]
+// CK-64: [[OMP_OFFLOAD_FAILED]]:
+// CK-64-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4foo2v_l32(ptr [[PVTARR]]) #[[ATTR2]]
+// CK-64-NEXT: br label %[[OMP_OFFLOAD_CONT]]
+// CK-64: [[OMP_OFFLOAD_CONT]]:
+// CK-64-NEXT: ret void
+//
+//
+// CK-64-LABEL: define internal void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z4foo2v_l32(
+// CK-64-SAME: ptr nonnull align 4 dereferenceable(40) [[PVTARR:%.*]]) #[[ATTR1]] {
+// CK-64-NEXT: [[ENTRY:.*:]]
+// CK-64-NEXT: [[PVTARR_ADDR:%.*]] = alloca ptr, align 8
+// CK-64-NEXT: [[PVTARR1:%.*]] = alloca [10 x i32], align 4
+// CK-64-NEXT: store ptr [[PVTARR]], ptr [[PVTARR_ADDR]], align 8
+// CK-64-NEXT: [[TMP0:%.*]] = load ptr, ptr [[PVTARR_ADDR]], align 8, !nonnull [[META17:![0-9]+]], !align [[META18:![0-9]+]]
+// CK-64-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x i32], ptr [[PVTARR1]], i64 0, i64 5
+// CK-64-NEXT: [[TMP1:%.*]] = load i32, ptr [[ARRAYIDX]], align 4
+// CK-64-NEXT: [[INC:%.*]] = add nsw i32 [[TMP1]], 1
+// CK-64-NEXT: store i32 [[INC]], ptr [[ARRAYIDX]], align 4
+// CK-64-NEXT: ret void
+//
+//
+// CK-64-LABEL: define dso_local void @_Z4foo3v(
+// CK-64-SAME: ) #[[ATTR0]] {
+// CK-64-NEXT: [[ENTRY:.*:]]
+// CK-64-NEXT: [[PA:%.*]] = alloca ptr, align 8
+// CK-64-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 8
+// CK-64-NEXT: [[DOTOFFLOAD_PTRS...
[truncated]
|
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
2 participants
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
Per OpenMP 6.0 specification, section 7.5.1, default Clause
Page 224, lines 3-5 default Clause, Semantics
If data-sharing-attribute is shared then the clause has no effect
on a target construct; otherwise, its effect on a target construct is
equivalent to specifying the defaultmap clause with the same
data-sharing-attribute and variable-category.
Testing:
OpenMP LIT tests
check-all