- Notifications
You must be signed in to change notification settings - Fork14.5k
[OpenMP] Don't emit redundant zero-sized mapping nodes for overlapped structs#148947
New issue
Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.
By clicking “Sign up for GitHub”, you agree to ourterms of service andprivacy statement. We’ll occasionally send you account related emails.
Already on GitHub?Sign in to your account
Open
jtb20 wants to merge1 commit intollvm:mainChoose a base branch
base:main
Could not load branches
Branch not found:{{ refName }}
Loading
Could not load tags
Nothing to show
Loading
Are you sure you want to change the base?
Some commits from the old base branch may be removed from the timeline, and old review comments may become outdated.
+449 −72
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
Member
llvmbot commentedJul 15, 2025 • edited
Loading Uh oh!
There was an error while loading.Please reload this page.
edited
Uh oh!
There was an error while loading.Please reload this page.
@llvm/pr-subscribers-clang @llvm/pr-subscribers-openmp Author: Julian Brown (jtb20) ChangesThe handling of overlapped structure mapping in CGOpenMPRuntime.cpp can lead to redundant zero-sized mapping nodes at runtime. This patch fixes it using a combination of approaches: trivially adjacent struct members won't have a mapping node created between them, and for more complicated cases (inheritance) the physical layout of the struct/class is used to make sure that elements aren't missed. I've introduced a new class to track the state whilst iterating over the struct. This reduces a bit of redundancy in the code (accumulating CombinedInfo both during and after the loop), which I think is a bit neater. Before: After: For code: Patch is 25.84 KiB, truncated to 20.00 KiB below, full version:https://github.com/llvm/llvm-project/pull/148947.diff 8 Files Affected:
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cppindex ce2dd4d76368a..039210aa4342f 100644--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp@@ -7080,6 +7080,110 @@ class MappableExprsHandler { return ConstLength.getSExtValue() != 1; }+ /// A helper class to copy structures with overlapped elements, i.e. those+ /// which have mappings of both "s" and "s.mem". Consecutive elements that+ /// are not explicitly copied have mapping nodes synthesized for them,+ /// taking care to avoid generating zero-sized copies.+ class CopyOverlappedEntryGaps {+ CodeGenFunction &CGF;+ MapCombinedInfoTy &CombinedInfo;+ OpenMPOffloadMappingFlags Flags;+ const ValueDecl *MapDecl;+ const Expr *MapExpr;+ Address BP;+ bool IsNonContiguous;+ uint64_t DimSize;+ // These elements track the position as the struct is iterated over+ // (in order of increasing element address).+ const RecordDecl *LastParent = nullptr;+ uint64_t Cursor = 0;+ unsigned LastIndex = -1u;+ Address LB;++ public:+ CopyOverlappedEntryGaps(CodeGenFunction &_CGF,+ MapCombinedInfoTy &_CombinedInfo,+ OpenMPOffloadMappingFlags _Flags,+ const ValueDecl *_MapDecl, const Expr *_MapExpr,+ Address _BP, Address _LB, bool _IsNonContiguous,+ uint64_t _DimSize)+ : CGF(_CGF), CombinedInfo(_CombinedInfo), Flags(_Flags), MapDecl(_MapDecl),+ MapExpr(_MapExpr), BP(_BP), LB(_LB),+ IsNonContiguous(_IsNonContiguous), DimSize(_DimSize) { }++ void ProcessField(const OMPClauseMappableExprCommon::MappableComponent &MC,+ const FieldDecl *FD,+ llvm::function_ref<LValue(CodeGenFunction &, const MemberExpr *)> EmitMemberExprBase) {+ const RecordDecl *RD = FD->getParent();+ const ASTRecordLayout &RL = CGF.getContext().getASTRecordLayout(RD);+ uint64_t FieldOffset = RL.getFieldOffset(FD->getFieldIndex());+ uint64_t FieldSize = CGF.getContext().getTypeSize(FD->getType().getCanonicalType());+ Address ComponentLB = Address::invalid();++ if (FD->getType()->isLValueReferenceType()) {+ const auto *ME =+ cast<MemberExpr>(MC.getAssociatedExpression());+ LValue BaseLVal = EmitMemberExprBase(CGF, ME);+ ComponentLB =+ CGF.EmitLValueForFieldInitialization(BaseLVal, FD)+ .getAddress();+ } else {+ ComponentLB =+ CGF.EmitOMPSharedLValue(MC.getAssociatedExpression())+ .getAddress();+ }++ if (LastParent == nullptr) {+ LastParent = RD;+ }+ if (FD->getParent() == LastParent) {+ if (FD->getFieldIndex() != LastIndex + 1)+ CopyUntilField(FD, ComponentLB);+ } else {+ LastParent = FD->getParent();+ if (((int64_t)FieldOffset - (int64_t)Cursor) > 0)+ CopyUntilField(FD, ComponentLB);+ }+ Cursor = FieldOffset + FieldSize;+ LastIndex = FD->getFieldIndex();+ LB = CGF.Builder.CreateConstGEP(ComponentLB, 1);+ }++ void CopyUntilField(const FieldDecl *FD, Address ComponentLB) {+ llvm::Value *ComponentLBPtr = ComponentLB.emitRawPointer(CGF);+ llvm::Value *LBPtr = LB.emitRawPointer(CGF);+ llvm::Value *Size = CGF.Builder.CreatePtrDiff(CGF.Int8Ty, ComponentLBPtr,+ LBPtr);+ CopySizedChunk(LBPtr, Size);+ }++ void CopyUntilEnd(Address HB) {+ if (LastParent) {+ const ASTRecordLayout &RL = CGF.getContext().getASTRecordLayout(LastParent);+ if ((uint64_t)CGF.getContext().toBits(RL.getSize()) <= Cursor)+ return;+ }+ llvm::Value *LBPtr = LB.emitRawPointer(CGF);+ llvm::Value *Size = CGF.Builder.CreatePtrDiff(+ CGF.Int8Ty, CGF.Builder.CreateConstGEP(HB, 1).emitRawPointer(CGF),+ LBPtr);+ CopySizedChunk(LBPtr, Size);+ }++ void CopySizedChunk(llvm::Value *Base, llvm::Value *Size) {+ CombinedInfo.Exprs.emplace_back(MapDecl, MapExpr);+ CombinedInfo.BasePointers.push_back(BP.emitRawPointer(CGF));+ CombinedInfo.DevicePtrDecls.push_back(nullptr);+ CombinedInfo.DevicePointers.push_back(DeviceInfoTy::None);+ CombinedInfo.Pointers.push_back(Base);+ CombinedInfo.Sizes.push_back(CGF.Builder.CreateIntCast(+ Size, CGF.Int64Ty, /*isSigned=*/true));+ CombinedInfo.Types.push_back(Flags);+ CombinedInfo.Mappers.push_back(nullptr);+ CombinedInfo.NonContigInfo.Dims.push_back(IsNonContiguous ? DimSize : 1);+ }+ };+ /// Generate the base pointers, section pointers, sizes, map type bits, and /// user-defined mappers (all included in \a CombinedInfo) for the provided /// map type, map or motion modifiers, and expression components.@@ -7570,63 +7674,22 @@ class MappableExprsHandler { getMapTypeBits(MapType, MapModifiers, MotionModifiers, IsImplicit, /*AddPtrFlag=*/false, /*AddIsTargetParamFlag=*/false, IsNonContiguous);- llvm::Value *Size = nullptr;+ CopyOverlappedEntryGaps CopyGaps(CGF, CombinedInfo, Flags, MapDecl,+ MapExpr, BP, LB, IsNonContiguous,+ DimSize); // Do bitcopy of all non-overlapped structure elements. for (OMPClauseMappableExprCommon::MappableExprComponentListRef Component : OverlappedElements) {- Address ComponentLB = Address::invalid(); for (const OMPClauseMappableExprCommon::MappableComponent &MC : Component) { if (const ValueDecl *VD = MC.getAssociatedDeclaration()) {- const auto *FD = dyn_cast<FieldDecl>(VD);- if (FD && FD->getType()->isLValueReferenceType()) {- const auto *ME =- cast<MemberExpr>(MC.getAssociatedExpression());- LValue BaseLVal = EmitMemberExprBase(CGF, ME);- ComponentLB =- CGF.EmitLValueForFieldInitialization(BaseLVal, FD)- .getAddress();- } else {- ComponentLB =- CGF.EmitOMPSharedLValue(MC.getAssociatedExpression())- .getAddress();+ if (const auto *FD = dyn_cast<FieldDecl>(VD)) {+ CopyGaps.ProcessField(MC, FD, EmitMemberExprBase); }- llvm::Value *ComponentLBPtr = ComponentLB.emitRawPointer(CGF);- llvm::Value *LBPtr = LB.emitRawPointer(CGF);- Size = CGF.Builder.CreatePtrDiff(CGF.Int8Ty, ComponentLBPtr,- LBPtr);- break; } }- assert(Size && "Failed to determine structure size");- CombinedInfo.Exprs.emplace_back(MapDecl, MapExpr);- CombinedInfo.BasePointers.push_back(BP.emitRawPointer(CGF));- CombinedInfo.DevicePtrDecls.push_back(nullptr);- CombinedInfo.DevicePointers.push_back(DeviceInfoTy::None);- CombinedInfo.Pointers.push_back(LB.emitRawPointer(CGF));- CombinedInfo.Sizes.push_back(CGF.Builder.CreateIntCast(- Size, CGF.Int64Ty, /*isSigned=*/true));- CombinedInfo.Types.push_back(Flags);- CombinedInfo.Mappers.push_back(nullptr);- CombinedInfo.NonContigInfo.Dims.push_back(IsNonContiguous ? DimSize- : 1);- LB = CGF.Builder.CreateConstGEP(ComponentLB, 1); }- CombinedInfo.Exprs.emplace_back(MapDecl, MapExpr);- CombinedInfo.BasePointers.push_back(BP.emitRawPointer(CGF));- CombinedInfo.DevicePtrDecls.push_back(nullptr);- CombinedInfo.DevicePointers.push_back(DeviceInfoTy::None);- CombinedInfo.Pointers.push_back(LB.emitRawPointer(CGF));- llvm::Value *LBPtr = LB.emitRawPointer(CGF);- Size = CGF.Builder.CreatePtrDiff(- CGF.Int8Ty, CGF.Builder.CreateConstGEP(HB, 1).emitRawPointer(CGF),- LBPtr);- CombinedInfo.Sizes.push_back(- CGF.Builder.CreateIntCast(Size, CGF.Int64Ty, /*isSigned=*/true));- CombinedInfo.Types.push_back(Flags);- CombinedInfo.Mappers.push_back(nullptr);- CombinedInfo.NonContigInfo.Dims.push_back(IsNonContiguous ? DimSize- : 1);+ CopyGaps.CopyUntilEnd(HB); break; } llvm::Value *Size = getExprTypeSize(I->getAssociatedExpression());diff --git a/clang/test/OpenMP/copy-gaps-1.cpp b/clang/test/OpenMP/copy-gaps-1.cppnew file mode 100644index 0000000000000..3d4fae352eed6--- /dev/null+++ b/clang/test/OpenMP/copy-gaps-1.cpp@@ -0,0 +1,52 @@+// RUN: %clang_cc1 -verify -triple x86_64-pc-linux-gnu -fopenmp-targets=amdgcn-amd-amdhsa -fopenmp -emit-llvm %s -o - | FileCheck %s+// expected-no-diagnostics++struct S {+ int x;+ int y;+ int z;+ int *p1;+ int *p2;+};++struct T : public S {+ int a;+ int b;+ int c;+};++int main() {+ T v;++#pragma omp target map(tofrom: v, v.x, v.y, v.z, v.p1[:8], v.a, v.b, v.c)+ {+ v.x++;+ v.y += 2;+ v.z += 3;+ v.p1[0] += 4;+ v.a += 7;+ v.b += 5;+ v.c += 6;+ }++ return 0;+}++// CHECK: [[CSTSZ:@.+]] = private {{.*}}constant [10 x i64] [i64 0, i64 0, i64 0, i64 4, i64 4, i64 4, i64 32, i64 4, i64 4, i64 4]+// CHECK: [[CSTTY:@.+]] = private {{.*}}constant [10 x i64] [i64 [[#0x20]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000013]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]]]++// CHECK-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])+// CHECK-DAG: [[KSIZE:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4+// CHECK-DAG: store ptr [[SZBASE:%.+]], ptr [[KSIZE]], align 8+// CHECK-DAG: [[SZBASE]] = getelementptr inbounds [10 x i64], ptr [[SIZES:%[^,]*]], i32 0, i32 0++// Check for filling of four non-constant size elements here: the whole struct+// size, the (padded) region covering p1 & p2, and the padding at the end of+// struct T.++// CHECK-DAG: [[STR:%.+]] = getelementptr inbounds [10 x i64], ptr [[SIZES]], i32 0, i32 0+// CHECK-DAG: store i64 %{{.+}}, ptr [[STR]], align 8+// CHECK-DAG: [[P1P2:%.+]] = getelementptr inbounds [10 x i64], ptr [[SIZES]], i32 0, i32 1+// CHECK-DAG: store i64 %{{.+}}, ptr [[P1P2]], align 8+// CHECK-DAG: [[PAD:%.+]] = getelementptr inbounds [10 x i64], ptr [[SIZES]], i32 0, i32 2+// CHECK-DAG: store i64 %{{.+}}, ptr [[PAD]], align 8diff --git a/clang/test/OpenMP/copy-gaps-2.cpp b/clang/test/OpenMP/copy-gaps-2.cppnew file mode 100644index 0000000000000..5bf603a3d9edb--- /dev/null+++ b/clang/test/OpenMP/copy-gaps-2.cpp@@ -0,0 +1,52 @@+// RUN: %clang_cc1 -verify -triple x86_64-pc-linux-gnu -fopenmp-targets=amdgcn-amd-amdhsa -fopenmp -emit-llvm %s -o - | FileCheck %s+// expected-no-diagnostics++struct S {+ int x;+ int y;+ int z;+};++struct M : public S {+ int mid;+};++struct T : public M {+ int a;+ int b;+ int c;+};++int main() {+ T v;++#pragma omp target map(tofrom: v, v.y, v.z, v.a)+ {+ v.y++;+ v.z += 2;+ v.a += 3;+ v.mid += 5;+ }++ return 0;+}++// CHECK: [[CSTSZ:@.+]] = private {{.*}}constant [7 x i64] [i64 0, i64 0, i64 0, i64 0, i64 4, i64 4, i64 4]+// CHECK: [[CSTTY:@.+]] = private {{.*}}constant [7 x i64] [i64 [[#0x20]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]]]++// CHECK-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])+// CHECK-DAG: [[KSIZE:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4+// CHECK-DAG: store ptr [[SZBASE:%.+]], ptr [[KSIZE]], align 8+// CHECK-DAG: [[SZBASE]] = getelementptr inbounds [7 x i64], ptr [[SIZES:%[^,]*]], i32 0, i32 0++// Fill four non-constant size elements here: the whole struct size, the region+// covering v.x, the region covering v.mid and the region covering v.b and v.c.++// CHECK-DAG: [[STR:%.+]] = getelementptr inbounds [7 x i64], ptr [[SIZES]], i32 0, i32 0+// CHECK-DAG: store i64 %{{.+}}, ptr [[STR]], align 8+// CHECK-DAG: [[X:%.+]] = getelementptr inbounds [7 x i64], ptr [[SIZES]], i32 0, i32 1+// CHECK-DAG: store i64 %{{.+}}, ptr [[X]], align 8+// CHECK-DAG: [[MID:%.+]] = getelementptr inbounds [7 x i64], ptr [[SIZES]], i32 0, i32 2+// CHECK-DAG: store i64 %{{.+}}, ptr [[MID]], align 8+// CHECK-DAG: [[BC:%.+]] = getelementptr inbounds [7 x i64], ptr [[SIZES]], i32 0, i32 3+// CHECK-DAG: store i64 %{{.+}}, ptr [[BC]], align 8diff --git a/clang/test/OpenMP/copy-gaps-3.cpp b/clang/test/OpenMP/copy-gaps-3.cppnew file mode 100644index 0000000000000..5febb181ca1c5--- /dev/null+++ b/clang/test/OpenMP/copy-gaps-3.cpp@@ -0,0 +1,46 @@+// RUN: %clang_cc1 -verify -triple x86_64-pc-linux-gnu -fopenmp-targets=amdgcn-amd-amdhsa -fopenmp -emit-llvm %s -o - | FileCheck %s+// expected-no-diagnostics++struct S {+ int x;+ int y;+ int z;+};++struct T : public S {+ int a;+ int b;+ int c;+};++int main() {+ T v;++ // This one should have no gap between v.z & v.a.+#pragma omp target map(tofrom: v, v.y, v.z, v.a)+ {+ v.y++;+ v.z += 2;+ v.a += 3;+ }++ return 0;+}++// CHECK: [[CSTSZ:@.+]] = private {{.*}}constant [6 x i64] [i64 0, i64 0, i64 0, i64 4, i64 4, i64 4]+// CHECK: [[CSTTY:@.+]] = private {{.*}}constant [6 x i64] [i64 [[#0x20]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]]]++// CHECK-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])+// CHECK-DAG: [[KSIZE:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4+// CHECK-DAG: store ptr [[SZBASE:%.+]], ptr [[KSIZE]], align 8+// CHECK-DAG: [[SZBASE]] = getelementptr inbounds [6 x i64], ptr [[SIZES:%[^,]*]], i32 0, i32 0++// Fill three non-constant size elements here: the whole struct size, the region+// covering v.x, and the region covering v.b and v.c.++// CHECK-DAG: [[STR:%.+]] = getelementptr inbounds [6 x i64], ptr [[SIZES]], i32 0, i32 0+// CHECK-DAG: store i64 %{{.+}}, ptr [[STR]], align 8+// CHECK-DAG: [[X:%.+]] = getelementptr inbounds [6 x i64], ptr [[SIZES]], i32 0, i32 1+// CHECK-DAG: store i64 %{{.+}}, ptr [[X]], align 8+// CHECK-DAG: [[BC:%.+]] = getelementptr inbounds [6 x i64], ptr [[SIZES]], i32 0, i32 2+// CHECK-DAG: store i64 %{{.+}}, ptr [[BC]], align 8diff --git a/clang/test/OpenMP/copy-gaps-4.cpp b/clang/test/OpenMP/copy-gaps-4.cppnew file mode 100644index 0000000000000..7060fe3ea2a01--- /dev/null+++ b/clang/test/OpenMP/copy-gaps-4.cpp@@ -0,0 +1,48 @@+// RUN: %clang_cc1 -verify -triple x86_64-pc-linux-gnu -fopenmp-targets=amdgcn-amd-amdhsa -fopenmp -emit-llvm %s -o - | FileCheck %s+// expected-no-diagnostics++struct S {+ int x;+ int y;+ char z; // Hidden padding after here...+};++struct T : public S {+ int a;+ int b;+ int c;+};++int main() {+ T v;++#pragma omp target map(tofrom: v, v.y, v.z, v.a)+ {+ v.y++;+ v.z += 2;+ v.a += 3;+ }++ return 0;+}++// CHECK: [[CSTSZ:@.+]] = private {{.*}}constant [7 x i64] [i64 0, i64 0, i64 0, i64 0, i64 4, i64 1, i64 4]+// CHECK: [[CSTTY:@.+]] = private {{.*}}constant [7 x i64] [i64 [[#0x20]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]]]++// CHECK-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])+// CHECK-DAG: [[KSIZE:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4+// CHECK-DAG: store ptr [[SZBASE:%.+]], ptr [[KSIZE]], align 8+// CHECK-DAG: [[SZBASE]] = getelementptr inbounds [7 x i64], ptr [[SIZES:%[^,]*]], i32 0, i32 0++// Fill four non-constant size elements here: the whole struct size, the region+// covering v.x, the region covering padding after v.z and the region covering+// v.b and v.c.++// CHECK-DAG: [[STR:%.+]] = getelementptr inbounds [7 x i64], ptr [[SIZES]], i32 0, i32 0+// CHECK-DAG: store i64 %{{.+}}, ptr [[STR]], align 8+// CHECK-DAG: [[X:%.+]] = getelementptr inbounds [7 x i64], ptr [[SIZES]], i32 0, i32 1+// CHECK-DAG: store i64 %{{.+}}, ptr [[X]], align 8+// CHECK-DAG: [[PAD:%.+]] = getelementptr inbounds [7 x i64], ptr [[SIZES]], i32 0, i32 2+// CHECK-DAG: store i64 %{{.+}}, ptr [[PAD]], align 8+// CHECK-DAG: [[BC:%.+]] = getelementptr inbounds [7 x i64], ptr [[SIZES]], i32 0, i32 3+// CHECK-DAG: store i64 %{{.+}}, ptr [[BC]], align 8diff --git a/clang/test/OpenMP/copy-gaps-5.cpp b/clang/test/OpenMP/copy-gaps-5.cppnew file mode 100644index 0000000000000..fae675dc2f505--- /dev/null+++ b/clang/test/OpenMP/copy-gaps-5.cpp@@ -0,0 +1,50 @@+// RUN: %clang_cc1 -verify -triple x86_64-pc-linux-gnu -fopenmp-targets=amdgcn-amd-amdhsa -fopenmp -emit-llvm %s -o - | FileCheck %s+// expected-no-diagnostics++template<typename C>+struct S {+ C x;+ C y;+ char z; // Hidden padding after here...+};++template<typename C>+struct T : public S<C> {+ C a;+ C b;+ C c;+};++int main() {+ T<int> v;++#pragma omp target map(tofrom: v, v.y, v.z, v.a)+ {+ v.y++;+ v.z += 2;+ v.a += 3;+ }++ return 0;+}++// CHECK: [[CSTSZ:@.+]] = private {{.*}}constant [7 x i64] [i64 0, i64 0, i64 0, i64 0, i64 4, i64 1, i64 4]+// CHECK: [[CSTTY:@.+]] = private {{.*}}constant [7 x i64] [i64 [[#0x20]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]], i64 [[#0x1000000000003]]]++// CHECK-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])+// CHECK-DAG: [[KSIZE:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4+// CHECK-DAG: store ptr [[SZBASE:%.+]], ptr [[KSIZE]], align 8+// CHECK-DAG: [[SZBASE]] = getelementptr inbounds [7 x i64], ptr [[SIZES:%[^,]*]], i32 0, i32 0++// Fill four non-constant size elements here: the whole struct size, the region+// covering v.x, the region covering padding after v.z and the region covering+// v.b and v.c.++// CHECK-DAG: [[STR:%.+]] = getelementptr inbounds [7 x i64], ptr [[SIZES]], i32 0, i32 0+// CHECK-DAG: store i64 %{{.+}}, ptr [[STR]], align 8+// CHECK-DAG: [[X:%.+]] = getelementptr inbounds [7 x i64], ptr [[SIZES]], i32 0, i32 1+// CHECK-DAG: store i64 %{{.+}}, ptr [[X]], align 8+// CHECK-DAG: [[PAD:%.+]] = getelementptr inbounds [7 x i64], ptr [[SIZES]], i32 0, i32 2+// CHECK-DAG: store i64 %{{.+}}, ptr [[PAD]], align 8+// CHECK-DAG: [[BC:%.+]] = getelementptr inbounds [7 x i64], ptr [[SIZES]], i32 0, i32 3+// CHECK-DAG: store i64 %{{.+}}, ptr [[BC]], align 8diff --git a/clang/test/OpenMP/copy-gaps-6.cpp b/clang/test/OpenMP/copy-gaps-6.cppnew file mode 100644index 0000000000000..9c62fde1c3762--- /dev/null+++ b/clang/test/OpenMP/copy-gaps-6.cpp@@ -0,0 +1,87 @@+// RUN: %clang_cc1 -verify -triple x86_64-pc-linux-gnu -fopenmp-targets=amdgcn-amd-amdhsa -fopenmp -emit-llvm %s -o - | FileCheck %s+// expected-no-diagnostics++struct S {+ int x;+ int *arr;+ int y;+ int z;+};++int main() {+ S v;++#pragma omp target map(tofrom: v, v.x, v.z)+ {+ v.x++;+ v.y += 2;+ v.z += 3;+ }++#pragma omp target map(tofrom: v, v.x, v.arr[:1])+ {+ v.x++;+ v.y += 2;+ v.arr[0] += 2;+ v.z += 4;+ }++#pragma omp target map(tofrom: v, v.arr[:1])+ {+ v.x++;+ v.y += 2;+ v.arr[0] += 2;+ v.z += 4;+ }++ return 0;+}++// CHECK: [[CSTSZ0:@.+]] = private {{.*}}constant [4 x i64] [i64 0, i64 0, i64 4...[truncated] |
github-actionsbot commentedJul 15, 2025 • edited
Loading Uh oh!
There was an error while loading.Please reload this page.
edited
Uh oh!
There was an error while loading.Please reload this page.
✅ With the latest revision this PR passed the C/C++ code formatter. |
849f220 tod6b9747CompareUh oh!
There was an error while loading.Please reload this page.
Uh oh!
There was an error while loading.Please reload this page.
Uh oh!
There was an error while loading.Please reload this page.
Uh oh!
There was an error while loading.Please reload this page.
| @@ -0,0 +1,52 @@ | |||
| // RUN: %clang_cc1 -verify -triple x86_64-pc-linux-gnu -fopenmp-targets=amdgcn-amd-amdhsa -fopenmp -emit-llvm %s -o - | FileCheck %s | |||
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others.Learn more.
auto generate check lines pls
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others.Learn more.
This one though I'm not sure how to do -- the check lines for the new copy-gaps tests are manually written, as it appears were the check lines for the altered test target_map_codegen_35.cpp (which seems to be the only extant test that this patch affects). Would it be better to have auto-generated check lines for the whole generated output for the new tests? That seems like it'd be more brittle, to me.
… structsThe handling of overlapped structure mapping in CGOpenMPRuntime.cpp canlead to redundant zero-sized mapping nodes at runtime. This patch fixesit using a combination of approaches: trivially adjacent struct memberswon't have a mapping node created between them, and for more complicatedcases (inheritance) the physical layout of the struct/class is used tomake sure that elements aren't missed.I've introduced a new class to track the state whilst iterating overthe struct. This reduces a bit of redundancy in the code (accumulatingCombinedInfo both during and after the loop), which I think is a bitneater.Before:omptarget --> Entry 0: Base=0x00007fff8d483830, Begin=0x00007fff8d483830, Size=48, Type=0x20, Name=unknownomptarget --> Entry 1: Base=0x00007fff8d483830, Begin=0x00007fff8d483830, Size=0, Type=0x1000000000003, Name=unknownomptarget --> Entry 2: Base=0x00007fff8d483830, Begin=0x00007fff8d483834, Size=0, Type=0x1000000000003, Name=unknownomptarget --> Entry 3: Base=0x00007fff8d483830, Begin=0x00007fff8d483838, Size=0, Type=0x1000000000003, Name=unknownomptarget --> Entry 4: Base=0x00007fff8d483830, Begin=0x00007fff8d48383c, Size=20, Type=0x1000000000003, Name=unknownomptarget --> Entry 5: Base=0x00007fff8d483830, Begin=0x00007fff8d483854, Size=0, Type=0x1000000000003, Name=unknownomptarget --> Entry 6: Base=0x00007fff8d483830, Begin=0x00007fff8d483858, Size=0, Type=0x1000000000003, Name=unknownomptarget --> Entry 7: Base=0x00007fff8d483830, Begin=0x00007fff8d48385c, Size=4, Type=0x1000000000003, Name=unknownomptarget --> Entry 8: Base=0x00007fff8d483830, Begin=0x00007fff8d483830, Size=4, Type=0x1000000000003, Name=unknownomptarget --> Entry 9: Base=0x00007fff8d483830, Begin=0x00007fff8d483834, Size=4, Type=0x1000000000003, Name=unknownomptarget --> Entry 10: Base=0x00007fff8d483830, Begin=0x00007fff8d483838, Size=4, Type=0x1000000000003, Name=unknownomptarget --> Entry 11: Base=0x00007fff8d483840, Begin=0x00005e7665275130, Size=32, Type=0x1000000000013, Name=unknownomptarget --> Entry 12: Base=0x00007fff8d483830, Begin=0x00007fff8d483850, Size=4, Type=0x1000000000003, Name=unknownomptarget --> Entry 13: Base=0x00007fff8d483830, Begin=0x00007fff8d483854, Size=4, Type=0x1000000000003, Name=unknownomptarget --> Entry 14: Base=0x00007fff8d483830, Begin=0x00007fff8d483858, Size=4, Type=0x1000000000003, Name=unknownAfter:omptarget --> Entry 0: Base=0x00007fffd0f562e0, Begin=0x00007fffd0f562e0, Size=48, Type=0x20, Name=unknownomptarget --> Entry 1: Base=0x00007fffd0f562e0, Begin=0x00007fffd0f562ec, Size=20, Type=0x1000000000003, Name=unknownomptarget --> Entry 2: Base=0x00007fffd0f562e0, Begin=0x00007fffd0f5630c, Size=4, Type=0x1000000000003, Name=unknownomptarget --> Entry 3: Base=0x00007fffd0f562e0, Begin=0x00007fffd0f562e0, Size=4, Type=0x1000000000003, Name=unknownomptarget --> Entry 4: Base=0x00007fffd0f562e0, Begin=0x00007fffd0f562e4, Size=4, Type=0x1000000000003, Name=unknownomptarget --> Entry 5: Base=0x00007fffd0f562e0, Begin=0x00007fffd0f562e8, Size=4, Type=0x1000000000003, Name=unknownomptarget --> Entry 6: Base=0x00007fffd0f562f0, Begin=0x000058b6013fb130, Size=32, Type=0x1000000000013, Name=unknownomptarget --> Entry 7: Base=0x00007fffd0f562e0, Begin=0x00007fffd0f56300, Size=4, Type=0x1000000000003, Name=unknownomptarget --> Entry 8: Base=0x00007fffd0f562e0, Begin=0x00007fffd0f56304, Size=4, Type=0x1000000000003, Name=unknownomptarget --> Entry 9: Base=0x00007fffd0f562e0, Begin=0x00007fffd0f56308, Size=4, Type=0x1000000000003, Name=unknownFor code: #include <cstdlib> #include <cstdio> struct S { int x; int y; int z; int *p1; int *p2; }; struct T : public S { int a; int b; int c; }; int main() { T v; v.p1 = (int*) calloc(8, sizeof(int)); v.p2 = (int*) calloc(8, sizeof(int)); #pragma omp target map(tofrom: v, v.x, v.y, v.z, v.p1[:8], v.a, v.b, v.c) { v.x++; v.y += 2; v.z += 3; v.p1[0] += 4; v.a += 7; v.b += 5; v.c += 6; } return 0; }d6b9747 toc30dcc3CompareSign up for freeto join this conversation on GitHub. Already have an account?Sign in to comment
Labels
clang:codegenIR generation bugs: mangling, exceptions, etc. clang:openmpOpenMP related changes to Clang clangClang issues not falling into any other category openmp
3 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.
The handling of overlapped structure mapping in CGOpenMPRuntime.cpp can lead to redundant zero-sized mapping nodes at runtime. This patch fixes it using a combination of approaches: trivially adjacent struct members won't have a mapping node created between them, and for more complicated cases (inheritance) the physical layout of the struct/class is used to make sure that elements aren't missed.
I've introduced a new class to track the state whilst iterating over the struct. This reduces a bit of redundancy in the code (accumulating CombinedInfo both during and after the loop), which I think is a bit neater.
Before:
After:
For code: