[clang] 19420c0 - [OpenMP] Fix non-contiguous array omp target update (#156889)
via cfe-commits
cfe-commits at lists.llvm.org
Thu Mar 26 07:55:38 PDT 2026
Author: Ivan R. Ivanov
Date: 2026-03-26T15:55:31+01:00
New Revision: 19420c0e7753db58651c7aa6b98c5d35a06830e0
URL: https://github.com/llvm/llvm-project/commit/19420c0e7753db58651c7aa6b98c5d35a06830e0
DIFF: https://github.com/llvm/llvm-project/commit/19420c0e7753db58651c7aa6b98c5d35a06830e0.diff
LOG: [OpenMP] Fix non-contiguous array omp target update (#156889)
The existing implementation has three issues which this patch addresses.
1. The last dimension which represents the bytes in the type, has the
wrong stride and count. For example, for a 4 byte int, count=1 and
stride=4. The correct representation here is count=4 and stride=1
because there are 4 bytes (count=4) that we need to copy and we do not
skip any bytes (stride=1).
2. The size of the data copy was computed using the last dimension.
However, this is incorrect in cases where some of the final dimensions
get merged into one. In this case we need to take the combined size of
the merged dimensions, which is (Count * Stride) of the first merged
dimension.
3. The Offset into a dimension was computed as a multiple of its Stride.
However, this Stride which is in bytes, already includes the stride
multiplier given by the user. This means that when the user specified
1:3:2, i.e. elements 1, 3, 5, the runtime incorrectly copied elements 2,
4, 6. Fix this by precomputing at compile time the Offset to be in bytes
by correctly multiplying the offset by the stride of the dimension
without the user-specified multiplier.
Added:
offload/test/offloading/strided_offset_multidim_update.c
Modified:
clang/lib/CodeGen/CGOpenMPRuntime.cpp
clang/test/OpenMP/target_update_codegen.cpp
clang/test/OpenMP/target_update_count_expression_codegen.c
clang/test/OpenMP/target_update_variable_stride_codegen.c
offload/libomptarget/omptarget.cpp
offload/test/offloading/non_contiguous_update.cpp
offload/test/offloading/strided_update_count_expression_complex.c
offload/test/offloading/strided_update_variable_stride_misc.c
Removed:
################################################################################
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index 3a57c06e40329..332b439c87472 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -8458,8 +8458,8 @@ class MappableExprsHandler {
// For supporting stride in array section, we need to initialize the first
// dimension size as 1, first offset as 0, and first count as 1
MapValuesArrayTy CurOffsets = {llvm::ConstantInt::get(CGF.CGM.Int64Ty, 0)};
- MapValuesArrayTy CurCounts = {llvm::ConstantInt::get(CGF.CGM.Int64Ty, 1)};
- MapValuesArrayTy CurStrides;
+ MapValuesArrayTy CurCounts;
+ MapValuesArrayTy CurStrides = {llvm::ConstantInt::get(CGF.CGM.Int64Ty, 1)};
MapValuesArrayTy DimSizes{llvm::ConstantInt::get(CGF.CGM.Int64Ty, 1)};
uint64_t ElementTypeSize;
@@ -8483,8 +8483,8 @@ class MappableExprsHandler {
"Should be either ConstantArray or VariableArray if not the "
"first Component");
- // Get element size if CurStrides is empty.
- if (CurStrides.empty()) {
+ // Get element size if CurCounts is empty.
+ if (CurCounts.empty()) {
const Type *ElementType = nullptr;
if (CAT)
ElementType = CAT->getElementType().getTypePtr();
@@ -8516,7 +8516,7 @@ class MappableExprsHandler {
ElementType = ElementType->getPointeeOrArrayElementType();
ElementTypeSize =
Context.getTypeSizeInChars(ElementType).getQuantity();
- CurStrides.push_back(
+ CurCounts.push_back(
llvm::ConstantInt::get(CGF.Int64Ty, ElementTypeSize));
}
}
@@ -8576,7 +8576,6 @@ class MappableExprsHandler {
CGF.Int64Ty,
/*isSigned=*/false);
}
- CurOffsets.push_back(Offset);
// Count
const Expr *CountExpr = OASE->getLength();
@@ -8613,11 +8612,12 @@ class MappableExprsHandler {
CurCounts.push_back(Count);
// Stride_n' = Stride_n * (D_0 * D_1 ... * D_n-1) * Unit size
+ // Offset_n' = Offset_n * (D_0 * D_1 ... * D_n-1) * Unit size
// Take `int arr[5][5][5]` and `arr[0:2:2][1:2:1][0:2:2]` as an example:
// Offset Count Stride
- // D0 0 1 4 (int) <- dummy dimension
+ // D0 0 4 1 (int) <- dummy dimension
// D1 0 2 8 (2 * (1) * 4)
- // D2 1 2 20 (1 * (1 * 5) * 4)
+ // D2 100 2 20 (1 * (1 * 5) * 4)
// D3 0 2 200 (2 * (1 * 5 * 4) * 4)
const Expr *StrideExpr = OASE->getStride();
llvm::Value *Stride =
@@ -8630,6 +8630,10 @@ class MappableExprsHandler {
CurStrides.push_back(CGF.Builder.CreateNUWMul(DimProd, Stride));
else
CurStrides.push_back(DimProd);
+
+ Offset = CGF.Builder.CreateNUWMul(DimProd, Offset);
+ CurOffsets.push_back(Offset);
+
if (DI != DimSizes.end())
++DI;
}
diff --git a/clang/test/OpenMP/target_update_codegen.cpp b/clang/test/OpenMP/target_update_codegen.cpp
index dc51e8969518b..1925d90651721 100644
--- a/clang/test/OpenMP/target_update_codegen.cpp
+++ b/clang/test/OpenMP/target_update_codegen.cpp
@@ -1144,7 +1144,7 @@ void foo(int arg) {
// CK20: store i64 {{32|64}}, ptr [[STRIDE]],
// CK20: [[DIM_2:%.+]] = getelementptr inbounds [3 x [[STRUCT_DESCRIPTOR]]], ptr [[DIMS]], {{.+}} 0, {{.+}} 1
// CK20: [[OFFSET_2:%.+]] = getelementptr inbounds nuw [[STRUCT_DESCRIPTOR]], ptr [[DIM_2]], {{.+}} 0, {{.+}} 0
- // CK20: store i64 1, ptr [[OFFSET_2]],
+ // CK20: store i64 {{8|16}}, ptr [[OFFSET_2]],
// CK20: [[COUNT_2:%.+]] = getelementptr inbounds nuw [[STRUCT_DESCRIPTOR]], ptr [[DIM_2]], {{.+}} 0, {{.+}} 1
// CK20: store i64 4, ptr [[COUNT_2]],
// CK20: [[STRIDE_2:%.+]] = getelementptr inbounds nuw [[STRUCT_DESCRIPTOR]], ptr [[DIM_2]], {{.+}} 0, {{.+}} 2
@@ -1153,9 +1153,9 @@ void foo(int arg) {
// CK20: [[OFFSET_3:%.+]] = getelementptr inbounds nuw [[STRUCT_DESCRIPTOR]], ptr [[DIM_3]], {{.+}} 0, {{.+}} 0
// CK20: store i64 0, ptr [[OFFSET_3]],
// CK20: [[COUNT_3:%.+]] = getelementptr inbounds nuw [[STRUCT_DESCRIPTOR]], ptr [[DIM_3]], {{.+}} 0, {{.+}} 1
- // CK20: store i64 1, ptr [[COUNT_3]],
+ // CK20: store i64 {{8|16}}, ptr [[COUNT_3]],
// CK20: [[STRIDE_3:%.+]] = getelementptr inbounds nuw [[STRUCT_DESCRIPTOR]], ptr [[DIM_3]], {{.+}} 0, {{.+}} 2
- // CK20: store i64 {{8|16}}, ptr [[STRIDE_3]],
+ // CK20: store i64 1, ptr [[STRIDE_3]],
// CK20-DAG: call void @__tgt_target_data_update_mapper(ptr @{{.+}}, i64 -1, i32 1, ptr [[GEPBP:%.+]], ptr [[GEPP:%.+]], ptr [[MSIZE]], ptr [[MTYPE]]{{.+}})
// CK20-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP]]
// CK20-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
@@ -1212,7 +1212,7 @@ struct ST {
// CK21: store i64 {{400|800}}, ptr [[STRIDE_1]],
// CK21: [[DIM_2:%.+]] = getelementptr inbounds [4 x [[STRUCT_DESCRIPTOR]]], ptr [[DIMS]], {{.+}} 0, {{.+}} 1
// CK21: [[OFFSET_2:%.+]] = getelementptr inbounds nuw [[STRUCT_DESCRIPTOR]], ptr [[DIM_2]], {{.+}} 0, {{.+}} 0
- // CK21: store i64 1, ptr [[OFFSET_2]],
+ // CK21: store i64 {{40|80}}, ptr [[OFFSET_2]],
// CK21: [[COUNT_2:%.+]] = getelementptr inbounds nuw [[STRUCT_DESCRIPTOR]], ptr [[DIM_2]], {{.+}} 0, {{.+}} 1
// CK21: store i64 3, ptr [[COUNT_2]],
// CK21: [[STRIDE_2:%.+]] = getelementptr inbounds nuw [[STRUCT_DESCRIPTOR]], ptr [[DIM_2]], {{.+}} 0, {{.+}} 2
@@ -1228,9 +1228,9 @@ struct ST {
// CK21: [[OFFSET_4:%.+]] = getelementptr inbounds nuw [[STRUCT_DESCRIPTOR]], ptr [[DIM_4]], {{.+}} 0, {{.+}} 0
// CK21: store i64 0, ptr [[OFFSET_4]],
// CK21: [[COUNT_4:%.+]] = getelementptr inbounds nuw [[STRUCT_DESCRIPTOR]], ptr [[DIM_4]], {{.+}} 0, {{.+}} 1
- // CK21: store i64 1, ptr [[COUNT_4]],
+ // CK21: store i64 {{4|8}}, ptr [[COUNT_4]],
// CK21: [[STRIDE_4:%.+]] = getelementptr inbounds nuw [[STRUCT_DESCRIPTOR]], ptr [[DIM_4]], {{.+}} 0, {{.+}} 2
- // CK21: store i64 {{4|8}}, ptr [[STRIDE_4]],
+ // CK21: store i64 1, ptr [[STRIDE_4]],
// CK21-DAG: call void @__tgt_target_data_update_mapper(ptr @{{.+}}, i64 -1, i32 2, ptr [[GEPBP:%.+]], ptr [[GEPP:%.+]], ptr [[GEPSZ:%.+]], ptr [[MTYPE]]{{.+}})
// CK21-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP]]
// CK21-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
@@ -1287,7 +1287,7 @@ struct ST {
// CK22: store i64 200, ptr [[STRIDE]],
// CK22: [[DIM_2:%.+]] = getelementptr inbounds [4 x [[STRUCT_DESCRIPTOR]]], ptr [[DIMS]], {{.+}} 0, {{.+}} 1
// CK22: [[OFFSET:%.+]] = getelementptr inbounds nuw [[STRUCT_DESCRIPTOR]], ptr [[DIM_2]], {{.+}} 0, {{.+}} 0
- // CK22: store i64 1, ptr [[OFFSET]],
+ // CK22: store i64 40, ptr [[OFFSET]],
// CK22: [[COUNT:%.+]] = getelementptr inbounds nuw [[STRUCT_DESCRIPTOR]], ptr [[DIM_2]], {{.+}} 0, {{.+}} 1
// CK22: store i64 3, ptr [[COUNT]],
// CK22: [[STRIDE:%.+]] = getelementptr inbounds nuw [[STRUCT_DESCRIPTOR]], ptr [[DIM_2]], {{.+}} 0, {{.+}} 2
@@ -1303,9 +1303,9 @@ struct ST {
// CK22: [[OFFSET:%.+]] = getelementptr inbounds nuw [[STRUCT_DESCRIPTOR]], ptr [[DIM_4]], {{.+}} 0, {{.+}} 0
// CK22: store i64 0, ptr [[OFFSET]],
// CK22: [[COUNT:%.+]] = getelementptr inbounds nuw [[STRUCT_DESCRIPTOR]], ptr [[DIM_4]], {{.+}} 0, {{.+}} 1
- // CK22: store i64 1, ptr [[COUNT]],
+ // CK22: store i64 4, ptr [[COUNT]],
// CK22: [[STRIDE:%.+]] = getelementptr inbounds nuw [[STRUCT_DESCRIPTOR]], ptr [[DIM_4]], {{.+}} 0, {{.+}} 2
- // CK22: store i64 4, ptr [[STRIDE]],
+ // CK22: store i64 1, ptr [[STRIDE]],
// CK22-DAG: call void @__tgt_target_data_update_mapper(ptr @{{.+}}, i64 -1, i32 1, ptr [[GEPBP:%.+]], ptr [[GEPP:%.+]], ptr [[MSIZE]], ptr [[MTYPE]]{{.+}})
// CK22-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP]]
// CK22-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
@@ -1368,7 +1368,7 @@ void foo(int arg) {
// CK23: store i64 200, ptr [[STRIDE]],
// CK23: [[DIM_2:%.+]] = getelementptr inbounds [4 x [[STRUCT_DESCRIPTOR]]], ptr [[DIMS]], {{.+}} 0, {{.+}} 1
// CK23: [[OFFSET_2:%.+]] = getelementptr inbounds nuw [[STRUCT_DESCRIPTOR]], ptr [[DIM_2]], {{.+}} 0, {{.+}} 0
- // CK23: store i64 1, ptr [[OFFSET_2]],
+ // CK23: store i64 20, ptr [[OFFSET_2]],
// CK23: [[COUNT_2:%.+]] = getelementptr inbounds nuw [[STRUCT_DESCRIPTOR]], ptr [[DIM_2]], {{.+}} 0, {{.+}} 1
// CK23: store i64 2, ptr [[COUNT_2]],
// CK23: [[STRIDE_2:%.+]] = getelementptr inbounds nuw [[STRUCT_DESCRIPTOR]], ptr [[DIM_2]], {{.+}} 0, {{.+}} 2
@@ -1384,9 +1384,9 @@ void foo(int arg) {
// CK23: [[OFFSET_4:%.+]] = getelementptr inbounds nuw [[STRUCT_DESCRIPTOR]], ptr [[DIM_4]], {{.+}} 0, {{.+}} 0
// CK23: store i64 0, ptr [[OFFSET_4]],
// CK23: [[COUNT_4:%.+]] = getelementptr inbounds nuw [[STRUCT_DESCRIPTOR]], ptr [[DIM_4]], {{.+}} 0, {{.+}} 1
- // CK23: store i64 1, ptr [[COUNT_4]],
+ // CK23: store i64 4, ptr [[COUNT_4]],
// CK23: [[STRIDE_4:%.+]] = getelementptr inbounds nuw [[STRUCT_DESCRIPTOR]], ptr [[DIM_4]], {{.+}} 0, {{.+}} 2
- // CK23: store i64 4, ptr [[STRIDE_4]],
+ // CK23: store i64 1, ptr [[STRIDE_4]],
// CK23-DAG: call void @__tgt_target_data_update_mapper(ptr @{{.+}}, i64 -1, i32 1, ptr [[GEPBP:%.+]], ptr [[GEPP:%.+]], ptr [[MSIZE]], ptr [[MTYPE]]{{.+}})
// CK23-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP]]
// CK23-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
@@ -1430,6 +1430,7 @@ void foo(int arg) {
// CK24: [[MUL:%.+]] = mul nuw i64 8,
// CK24: [[SUB:%.+]] = sub nuw i64 4, [[ARG:%.+]]
// CK24: [[LEN:%.+]] = udiv {{.+}} [[SUB]], 1
+ // CK24: [[MUL_ARG:%.+]] = mul nuw i64 40, [[ARG]]
// CK24: [[BP0:%.+]] = getelementptr inbounds [1 x ptr], ptr [[BP:%.+]], {{.+}} 0, {{.+}} 0
// CK24: store ptr [[ARR]], ptr [[BP0]],
// CK24: [[P0:%.+]] = getelementptr inbounds [1 x ptr], ptr [[P:%.+]], {{.+}} 0, {{.+}} 0
@@ -1443,7 +1444,7 @@ void foo(int arg) {
// CK24: store i64 320, ptr [[STRIDE]],
// CK24: [[DIM_2:%.+]] = getelementptr inbounds [4 x [[STRUCT_DESCRIPTOR]]], ptr [[DIMS]], {{.+}} 0, {{.+}} 1
// CK24: [[OFFSET_2:%.+]] = getelementptr inbounds nuw [[STRUCT_DESCRIPTOR]], ptr [[DIM_2]], {{.+}} 0, {{.+}} 0
- // CK24: store i64 [[ARG]], ptr [[OFFSET_2]],
+ // CK24: store i64 [[MUL_ARG]], ptr [[OFFSET_2]],
// CK24: [[COUNT_2:%.+]] = getelementptr inbounds nuw [[STRUCT_DESCRIPTOR]], ptr [[DIM_2]], {{.+}} 0, {{.+}} 1
// CK24: store i64 [[LEN]], ptr [[COUNT_2]],
// CK24: [[STRIDE_2:%.+]] = getelementptr inbounds nuw [[STRUCT_DESCRIPTOR]], ptr [[DIM_2]], {{.+}} 0, {{.+}} 2
@@ -1459,9 +1460,9 @@ void foo(int arg) {
// CK24: [[OFFSET_4:%.+]] = getelementptr inbounds nuw [[STRUCT_DESCRIPTOR]], ptr [[DIM_4]], {{.+}} 0, {{.+}} 0
// CK24: store i64 0, ptr [[OFFSET_4]],
// CK24: [[COUNT_4:%.+]] = getelementptr inbounds nuw [[STRUCT_DESCRIPTOR]], ptr [[DIM_4]], {{.+}} 0, {{.+}} 1
- // CK24: store i64 1, ptr [[COUNT_4]],
+ // CK24: store i64 8, ptr [[COUNT_4]],
// CK24: [[STRIDE_4:%.+]] = getelementptr inbounds nuw [[STRUCT_DESCRIPTOR]], ptr [[DIM_4]], {{.+}} 0, {{.+}} 2
- // CK24: store i64 8, ptr [[STRIDE_4]],
+ // CK24: store i64 1, ptr [[STRIDE_4]],
// CK24-DAG: call void @__tgt_target_data_update_mapper(ptr @{{.+}}, i64 -1, i32 1, ptr [[GEPBP:%.+]], ptr [[GEPP:%.+]], ptr [[MSIZE]], ptr [[MTYPE]]{{.+}})
// CK24-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP]]
// CK24-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
@@ -1526,7 +1527,7 @@ void foo(int arg) {
// CK25: store i64 20, ptr [[STRIDE_2]],
// CK25: [[DIM_3:%.+]] = getelementptr inbounds [4 x [[STRUCT_DESCRIPTOR]]], ptr [[DIMS]], {{.+}} 0, {{.+}} 2
// CK25: [[OFFSET_3:%.+]] = getelementptr inbounds nuw [[STRUCT_DESCRIPTOR]], ptr [[DIM_3]], {{.+}} 0, {{.+}} 0
- // CK25: store i64 1, ptr [[OFFSET_3]],
+ // CK25: store i64 4, ptr [[OFFSET_3]],
// CK25: [[COUNT_3:%.+]] = getelementptr inbounds nuw [[STRUCT_DESCRIPTOR]], ptr [[DIM_3]], {{.+}} 0, {{.+}} 1
// CK25: store i64 4, ptr [[COUNT_3]],
// CK25: [[STRIDE_3:%.+]] = getelementptr inbounds nuw [[STRUCT_DESCRIPTOR]], ptr [[DIM_3]], {{.+}} 0, {{.+}} 2
@@ -1535,9 +1536,9 @@ void foo(int arg) {
// CK25: [[OFFSET_4:%.+]] = getelementptr inbounds nuw [[STRUCT_DESCRIPTOR]], ptr [[DIM_4]], {{.+}} 0, {{.+}} 0
// CK25: store i64 0, ptr [[OFFSET_4]],
// CK25: [[COUNT_4:%.+]] = getelementptr inbounds nuw [[STRUCT_DESCRIPTOR]], ptr [[DIM_4]], {{.+}} 0, {{.+}} 1
- // CK25: store i64 1, ptr [[COUNT_4]],
+ // CK25: store i64 4, ptr [[COUNT_4]],
// CK25: [[STRIDE_4:%.+]] = getelementptr inbounds nuw [[STRUCT_DESCRIPTOR]], ptr [[DIM_4]], {{.+}} 0, {{.+}} 2
- // CK25: store i64 4, ptr [[STRIDE_4]],
+ // CK25: store i64 1, ptr [[STRIDE_4]],
// CK25: [[PTRS:%.+]] = getelementptr inbounds [3 x ptr], ptr %.offload_ptrs, i32 0, i32 0
// CK25: store ptr [[DIMS]], ptr [[PTRS]],
// CK25: [[DIM_5:%.+]] = getelementptr inbounds [3 x [[STRUCT_DESCRIPTOR]]], ptr [[DIMS_2]], {{.+}} 0, {{.+}} 0
@@ -1549,7 +1550,7 @@ void foo(int arg) {
// CK25: store i64 12, ptr [[STRIDE_2_1]],
// CK25: [[DIM_6:%.+]] = getelementptr inbounds [3 x [[STRUCT_DESCRIPTOR]]], ptr [[DIMS_2]], {{.+}} 0, {{.+}} 1
// CK25: [[OFFSET_2_2:%.+]] = getelementptr inbounds nuw [[STRUCT_DESCRIPTOR]], ptr [[DIM_6]], {{.+}} 0, {{.+}} 0
- // CK25: store i64 1, ptr [[OFFSET_2_2]],
+ // CK25: store i64 4, ptr [[OFFSET_2_2]],
// CK25: [[COUNT_2_2:%.+]] = getelementptr inbounds nuw [[STRUCT_DESCRIPTOR]], ptr [[DIM_6]], {{.+}} 0, {{.+}} 1
// CK25: store i64 2, ptr [[COUNT_2_2]],
// CK25: [[STRIDE_2_2:%.+]] = getelementptr inbounds nuw [[STRUCT_DESCRIPTOR]], ptr [[DIM_6]], {{.+}} 0, {{.+}} 2
@@ -1558,9 +1559,9 @@ void foo(int arg) {
// CK25: [[OFFSET_2_3:%.+]] = getelementptr inbounds nuw [[STRUCT_DESCRIPTOR]], ptr [[DIM_7]], {{.+}} 0, {{.+}} 0
// CK25: store i64 0, ptr [[OFFSET_2_3]],
// CK25: [[COUNT_2_3:%.+]] = getelementptr inbounds nuw [[STRUCT_DESCRIPTOR]], ptr [[DIM_7]], {{.+}} 0, {{.+}} 1
- // CK25: store i64 1, ptr [[COUNT_2_3]],
+ // CK25: store i64 4, ptr [[COUNT_2_3]],
// CK25: [[STRIDE_2_3:%.+]] = getelementptr inbounds nuw [[STRUCT_DESCRIPTOR]], ptr [[DIM_7]], {{.+}} 0, {{.+}} 2
- // CK25: store i64 4, ptr [[STRIDE_2_3]],
+ // CK25: store i64 1, ptr [[STRIDE_2_3]],
// CK25: [[PTRS_2:%.+]] = getelementptr inbounds [3 x ptr], ptr %.offload_ptrs, i32 0, i32 2
// CK25: store ptr [[DIMS_2]], ptr [[PTRS_2]],
// CK25-DAG: call void @__tgt_target_data_update_mapper(ptr @{{.+}}, i64 -1, i32 3, ptr [[GEPBP:%.+]], ptr [[GEPP:%.+]], ptr [[MSIZE]], ptr [[MTYPE]]{{.+}})
diff --git a/clang/test/OpenMP/target_update_count_expression_codegen.c b/clang/test/OpenMP/target_update_count_expression_codegen.c
index dc1a45a788846..e11f043f3602f 100644
--- a/clang/test/OpenMP/target_update_count_expression_codegen.c
+++ b/clang/test/OpenMP/target_update_count_expression_codegen.c
@@ -59,9 +59,9 @@ void test_contig_byte_size() {
// CHECK-NEXT: [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT_DESCRIPTOR_DIM]], ptr [[TMP7]], i32 0, i32 0
// CHECK-NEXT: store i64 0, ptr [[TMP8]], align 8
// CHECK-NEXT: [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT_DESCRIPTOR_DIM]], ptr [[TMP7]], i32 0, i32 1
-// CHECK-NEXT: store i64 1, ptr [[TMP9]], align 8
+// CHECK-NEXT: store i64 4, ptr [[TMP9]], align 8
// CHECK-NEXT: [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT_DESCRIPTOR_DIM]], ptr [[TMP7]], i32 0, i32 2
-// CHECK-NEXT: store i64 4, ptr [[TMP10]], align 8
+// CHECK-NEXT: store i64 1, ptr [[TMP10]], align 8
// CHECK-NEXT: [[TMP11:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
// CHECK-NEXT: store ptr [[DIMS]], ptr [[TMP11]], align 8
// CHECK-NEXT: [[TMP12:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
diff --git a/clang/test/OpenMP/target_update_variable_stride_codegen.c b/clang/test/OpenMP/target_update_variable_stride_codegen.c
index afc5459787c2e..5d4a582bf2090 100644
--- a/clang/test/OpenMP/target_update_variable_stride_codegen.c
+++ b/clang/test/OpenMP/target_update_variable_stride_codegen.c
@@ -67,9 +67,9 @@ void test_constant_stride_one() {
// CHECK-NEXT: [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT_DESCRIPTOR_DIM]], ptr [[TMP10]], i32 0, i32 0
// CHECK-NEXT: store i64 0, ptr [[TMP11]], align 8
// CHECK-NEXT: [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT_DESCRIPTOR_DIM]], ptr [[TMP10]], i32 0, i32 1
-// CHECK-NEXT: store i64 1, ptr [[TMP12]], align 8
+// CHECK-NEXT: store i64 4, ptr [[TMP12]], align 8
// CHECK-NEXT: [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT_DESCRIPTOR_DIM]], ptr [[TMP10]], i32 0, i32 2
-// CHECK-NEXT: store i64 4, ptr [[TMP13]], align 8
+// CHECK-NEXT: store i64 1, ptr [[TMP13]], align 8
// CHECK-NEXT: [[TMP14:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
// CHECK-NEXT: store ptr [[DIMS]], ptr [[TMP14]], align 8
// CHECK-NEXT: [[TMP15:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
@@ -105,9 +105,9 @@ void test_constant_stride_one() {
// CHECK-NEXT: [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT_DESCRIPTOR_DIM_0]], ptr [[TMP10]], i32 0, i32 0
// CHECK-NEXT: store i64 0, ptr [[TMP11]], align 8
// CHECK-NEXT: [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT_DESCRIPTOR_DIM_0]], ptr [[TMP10]], i32 0, i32 1
-// CHECK-NEXT: store i64 1, ptr [[TMP12]], align 8
+// CHECK-NEXT: store i64 4, ptr [[TMP12]], align 8
// CHECK-NEXT: [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT_DESCRIPTOR_DIM_0]], ptr [[TMP10]], i32 0, i32 2
-// CHECK-NEXT: store i64 4, ptr [[TMP13]], align 8
+// CHECK-NEXT: store i64 1, ptr [[TMP13]], align 8
// CHECK-NEXT: [[TMP14:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
// CHECK-NEXT: store ptr [[DIMS]], ptr [[TMP14]], align 8
// CHECK-NEXT: [[TMP15:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
diff --git a/offload/libomptarget/omptarget.cpp b/offload/libomptarget/omptarget.cpp
index 344c388e794af..5fe8dd705b5c5 100644
--- a/offload/libomptarget/omptarget.cpp
+++ b/offload/libomptarget/omptarget.cpp
@@ -1507,7 +1507,7 @@ static int targetDataNonContiguous(ident_t *Loc, DeviceTy &Device,
if (CurrentDim < DimSize) {
for (unsigned int I = 0; I < NonContig[CurrentDim].Count; ++I) {
uint64_t CurOffset =
- (NonContig[CurrentDim].Offset + I) * NonContig[CurrentDim].Stride;
+ NonContig[CurrentDim].Offset + I * NonContig[CurrentDim].Stride;
// we only need to transfer the first element for the last dimension
// since we've already got a contiguous piece.
if (CurrentDim != DimSize - 1 || I == 0) {
@@ -1578,9 +1578,18 @@ int targetDataUpdate(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
if (ArgTypes[I] & OMP_TGT_MAPTYPE_NON_CONTIG) {
__tgt_target_non_contig *NonContig = (__tgt_target_non_contig *)Args[I];
int32_t DimSize = ArgSizes[I];
- uint64_t Size =
- NonContig[DimSize - 1].Count * NonContig[DimSize - 1].Stride;
+ ODBG(ODT_DataTransfer) << "Non contig descriptor:";
+ for (int I = 0; I < DimSize; I++)
+ ODBG(ODT_DataTransfer)
+ << " Dim " << I << ": Offset " << NonContig[I].Offset << " Count "
+ << NonContig[I].Count << " Stride " << NonContig[I].Stride;
int32_t MergedDim = getNonContigMergedDimension(NonContig, DimSize);
+ ODBG(ODT_DataTransfer) << "Merged " << MergedDim << " dimensions";
+ __tgt_target_non_contig &FirstMergedDim =
+ NonContig[DimSize - MergedDim - 1];
+ uint64_t Size = FirstMergedDim.Count * FirstMergedDim.Stride;
+ ODBG(ODT_DataTransfer) << "Transfer size " << Size;
+ ODBG(ODT_DataTransfer) << "Base Ptr " << ArgsBase[I];
Ret = targetDataNonContiguous(
Loc, Device, ArgsBase[I], NonContig, Size, ArgTypes[I],
/*current_dim=*/0, DimSize - MergedDim, /*offset=*/0, AsyncInfo);
diff --git a/offload/test/offloading/non_contiguous_update.cpp b/offload/test/offloading/non_contiguous_update.cpp
index 609f0f967fb17..011e25a4a3b59 100644
--- a/offload/test/offloading/non_contiguous_update.cpp
+++ b/offload/test/offloading/non_contiguous_update.cpp
@@ -7,9 +7,9 @@
// Data structure definitions copied from OpenMP RTL.
struct __tgt_target_non_contig {
- int64_t offset;
- int64_t width;
- int64_t stride;
+ int64_t Offset;
+ int64_t Count;
+ int64_t Stride;
};
enum tgt_map_type { OMP_TGT_MAPTYPE_NON_CONTIG = 0x100000000000 };
@@ -26,68 +26,109 @@ void __tgt_target_data_update(int64_t device_id, int32_t arg_num,
#endif
int main() {
- // case 1
- // int arr[3][4][5][6];
- // #pragma omp target update to(arr[0:2][1:3][1:2][:])
- // set up descriptor
- __tgt_target_non_contig non_contig[5] = {
- {0, 2, 480}, {1, 3, 120}, {1, 2, 24}, {0, 6, 4}, {0, 1, 4}};
- int64_t size = 4, type = OMP_TGT_MAPTYPE_NON_CONTIG;
+ {
+ // case 1
+ // int32_t arr[3][4][5][6];
+ // #pragma omp target update to(arr[0:2][1:3][1:2][:])
+ // set up descriptor
+ __tgt_target_non_contig non_contig[5] = {{0, 2, 4 * 5 * 6 * 4},
+ {1 * 5 * 6 * 4, 3, 5 * 6 * 4},
+ {6 * 4, 2, 6 * 4},
+ {0, 6, 4},
+ {0, 4, 1}};
+ int64_t size = sizeof(non_contig) / sizeof(non_contig[0]),
+ type = OMP_TGT_MAPTYPE_NON_CONTIG;
- void *base;
- void *begin = &non_contig;
- int64_t *sizes = &size;
- int64_t *types = &type;
+ void *base;
+ void *begin = &non_contig;
+ int64_t *sizes = &size;
+ int64_t *types = &type;
- // The below diagram is the visualization of the non-contiguous transfer after
- // optimization. Note that each element represent the innermost dimension
- // (unit size = 24) since the stride * count of last dimension is equal to the
- // stride of second last dimension.
- //
- // OOOOO OOOOO OOOOO
- // OXXOO OXXOO OOOOO
- // OXXOO OXXOO OOOOO
- // OXXOO OXXOO OOOOO
- __tgt_target_data_update(/*device_id*/ -1, /*arg_num*/ 1, &base, &begin,
- sizes, types);
- // DEBUG: offset 144
- // DEBUG: offset 264
- // DEBUG: offset 384
- // DEBUG: offset 624
- // DEBUG: offset 744
- // DEBUG: offset 864
+ // The below diagram is the visualization of the non-contiguous transfer
+ // after optimization. Note that each element represent the merged innermost
+ // dimension (unit size = 24) since the stride * count of last dimension is
+ // equal to the stride of second last dimension.
+ //
+ // OOOOO OOOOO OOOOO
+ // OXXOO OXXOO OOOOO
+ // OXXOO OXXOO OOOOO
+ // OXXOO OXXOO OOOOO
+ __tgt_target_data_update(/*device_id*/ -1, /*arg_num*/ 1, &base, &begin,
+ sizes, types);
+ // DEBUG: offset 144 len 48
+ // DEBUG: offset 264 len 48
+ // DEBUG: offset 384 len 48
+ // DEBUG: offset 624 len 48
+ // DEBUG: offset 744 len 48
+ // DEBUG: offset 864 len 48
+ }
- // case 2
- // double darr[3][4][5];
- // #pragma omp target update to(darr[0:2:2][2:2][:2:2])
- // set up descriptor
- __tgt_target_non_contig non_contig_2[4] = {
- {0, 2, 320}, {2, 2, 40}, {0, 2, 16}, {0, 1, 8}};
- int64_t size_2 = 4, type_2 = OMP_TGT_MAPTYPE_NON_CONTIG;
+ {
+ // case 2
+ // int64_t darr[3][4][5];
+ // #pragma omp target update to(darr[0:2:2][2:2][:2:2])
+ // set up descriptor
+ __tgt_target_non_contig non_contig[4] = {
+ {0, 2, 2 * 4 * 5 * 8}, {2 * 5 * 8, 2, 5 * 8}, {0, 2, 2 * 8}, {0, 8, 1}};
+ int64_t size = sizeof(non_contig) / sizeof(non_contig[0]),
+ type = OMP_TGT_MAPTYPE_NON_CONTIG;
- void *base_2;
- void *begin_2 = &non_contig_2;
- int64_t *sizes_2 = &size_2;
- int64_t *types_2 = &type_2;
+ void *base;
+ void *begin = &non_contig;
+ int64_t *sizes = &size;
+ int64_t *types = &type;
+
+ // The below diagram is the visualization of the non-contiguous transfer
+ // after optimization. Note that each element represent the innermost
+ // dimension (unit size = 8).
+ //
+ // OOOOO OOOOO OOOOO
+ // OOOOO OOOOO OOOOO
+ // XOXOO OOOOO XOXOO
+ // XOXOO OOOOO XOXOO
+ __tgt_target_data_update(/*device_id*/ -1, /*arg_num*/ 1, &base, &begin,
+ sizes, types);
+ // DEBUG: offset 80 len 8
+ // DEBUG: offset 96 len 8
+ // DEBUG: offset 120 len 8
+ // DEBUG: offset 136 len 8
+ // DEBUG: offset 400 len 8
+ // DEBUG: offset 416 len 8
+ // DEBUG: offset 440 len 8
+ // DEBUG: offset 456 len 8
+ }
+
+ {
+ // case 3
+ // int32_t darr[6][6];
+ // #pragma omp target update to(darr[1:2:2][2:3])
+ // set up descriptor
+ __tgt_target_non_contig non_contig[3] = {{1 * 6 * 4 * 1, 2, 2 * 6 * 4 * 1},
+ {2 * 4 * 1, 3, 1 * 4 * 1},
+ {0, 4, 1}};
+ int64_t size = sizeof(non_contig) / sizeof(non_contig[0]),
+ type = OMP_TGT_MAPTYPE_NON_CONTIG;
+
+ void *base;
+ void *begin = &non_contig;
+ int64_t *sizes = &size;
+ int64_t *types = &type;
+
+ // The below diagram is the visualization of the non-contiguous transfer
+ // after optimization. Note that each element represent the merged innermost
+ // dimension (unit size = 4).
+ //
+ // OOOOOO
+ // OOXXXO
+ // OOOOOO
+ // OOXXXO
+ // OOOOOO
+ // OOOOOO
+ __tgt_target_data_update(/*device_id*/ -1, /*arg_num*/ 1, &base, &begin,
+ sizes, types);
+ // DEBUG: offset 32 len 12
+ // DEBUG: offset 80 len 12
+ }
- // The below diagram is the visualization of the non-contiguous transfer after
- // optimization. Note that each element represent the innermost dimension
- // (unit size = 24) since the stride * count of last dimension is equal to the
- // stride of second last dimension.
- //
- // OOOOO OOOOO OOOOO
- // OOOOO OOOOO OOOOO
- // XOXOO OOOOO XOXOO
- // XOXOO OOOOO XOXOO
- __tgt_target_data_update(/*device_id*/ -1, /*arg_num*/ 1, &base_2, &begin_2,
- sizes_2, types_2);
- // DEBUG: offset 80
- // DEBUG: offset 96
- // DEBUG: offset 120
- // DEBUG: offset 136
- // DEBUG: offset 400
- // DEBUG: offset 416
- // DEBUG: offset 440
- // DEBUG: offset 456
return 0;
}
diff --git a/offload/test/offloading/strided_offset_multidim_update.c b/offload/test/offloading/strided_offset_multidim_update.c
new file mode 100644
index 0000000000000..bc0b45cc5699b
--- /dev/null
+++ b/offload/test/offloading/strided_offset_multidim_update.c
@@ -0,0 +1,95 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+// XFAIL: intelgpu
+
+// Make sure multi-dimensional strided offset update works correctly.
+
+#include <stdio.h>
+
+#define N 6
+#define SLICE1 1 : 2 : 2
+#define SLICE2 2 : 3
+int main() {
+ int darr[N][N];
+
+ for (int i = 0; i < N; i++)
+ for (int j = 0; j < N; j++)
+ darr[i][j] = 100 + 10 * i + j;
+
+ printf("Full array\n");
+ for (int i = 0; i < N; i++) {
+ for (int j = 0; j < N; j++) {
+ printf("\t%d", darr[i][j]);
+ }
+ printf("\n");
+ }
+
+#pragma omp target enter data map(alloc : darr[0 : N][0 : N])
+
+ // Zero out target array.
+#pragma omp target
+ for (int i = 0; i < N; i++)
+ for (int j = 0; j < N; j++)
+ darr[i][j] = 0;
+
+ // Only copy over the slice to the device.
+#pragma omp target update to(darr[SLICE1][SLICE2])
+ // Then copy over the entire array to the host.
+#pragma omp target exit data map(from : darr[0 : N][0 : N])
+
+ printf("Only slice (to)\n");
+ for (int i = 0; i < N; i++) {
+ for (int j = 0; j < N; j++) {
+ printf("\t%d", darr[i][j]);
+ }
+ printf("\n");
+ }
+
+#pragma omp target enter data map(alloc : darr[0 : N][0 : N])
+
+ // Initialize on the device
+#pragma omp target
+ for (int i = 0; i < N; i++)
+ for (int j = 0; j < N; j++)
+ darr[i][j] = 100 + 10 * i + j;
+
+ // Zero out host array.
+ for (int i = 0; i < N; i++)
+ for (int j = 0; j < N; j++)
+ darr[i][j] = 0;
+
+ // Copy over only the slice to the host
+#pragma omp target update from(darr[SLICE1][SLICE2])
+#pragma omp target exit data map(delete : darr[0 : N][0 : N])
+
+ printf("Only slice (from)\n");
+ for (int i = 0; i < N; i++) {
+ for (int j = 0; j < N; j++) {
+ printf("\t%d", darr[i][j]);
+ }
+ printf("\n");
+ }
+
+ return 0;
+}
+
+// CHECK: Full array
+// CHECK-NEXT: 100 101 102 103 104 105
+// CHECK-NEXT: 110 111 112 113 114 115
+// CHECK-NEXT: 120 121 122 123 124 125
+// CHECK-NEXT: 130 131 132 133 134 135
+// CHECK-NEXT: 140 141 142 143 144 145
+// CHECK-NEXT: 150 151 152 153 154 155
+// CHECK-NEXT: Only slice (to)
+// CHECK-NEXT: 0 0 0 0 0 0
+// CHECK-NEXT: 0 0 112 113 114 0
+// CHECK-NEXT: 0 0 0 0 0 0
+// CHECK-NEXT: 0 0 132 133 134 0
+// CHECK-NEXT: 0 0 0 0 0 0
+// CHECK-NEXT: 0 0 0 0 0 0
+// CHECK-NEXT: Only slice (from)
+// CHECK-NEXT: 0 0 0 0 0 0
+// CHECK-NEXT: 0 0 112 113 114 0
+// CHECK-NEXT: 0 0 0 0 0 0
+// CHECK-NEXT: 0 0 132 133 134 0
+// CHECK-NEXT: 0 0 0 0 0 0
+// CHECK-NEXT: 0 0 0 0 0 0
diff --git a/offload/test/offloading/strided_update_count_expression_complex.c b/offload/test/offloading/strided_update_count_expression_complex.c
index a2ebdcd1f510d..92fb9c7da0c46 100644
--- a/offload/test/offloading/strided_update_count_expression_complex.c
+++ b/offload/test/offloading/strided_update_count_expression_complex.c
@@ -1,5 +1,5 @@
// RUN: %libomptarget-compile-run-and-check-generic
-// XFAIL: intelgpu
+// XFAIL: *
// Tests non-contiguous array sections with complex expression-based count
// scenarios including multiple struct arrays and non-zero offset.
@@ -130,6 +130,8 @@ void test_2_complex_count_with_offset() {
}
// Count: (len-offset)/2 with stride 2
+ // s1.arr[2:4:2]
+ // s2.arr[1:4:2]
#pragma omp target update from( \
s1.arr[s1.offset : (s1.len - s1.offset) / 2 : 2], \
s2.arr[s2.offset : (s2.len - s2.offset) / 2 : 2])
@@ -238,14 +240,14 @@ void test_2_complex_count_with_offset() {
// CHECK: s1 results:
// CHECK-NEXT: 0.000000
// CHECK-NEXT: 1.000000
-// CHECK-NEXT: 2.000000
-// CHECK-NEXT: 6.000000
// CHECK-NEXT: 4.000000
-// CHECK-NEXT: 10.000000
-// CHECK-NEXT: 6.000000
-// CHECK-NEXT: 14.000000
+// CHECK-NEXT: 3.000000
// CHECK-NEXT: 8.000000
-// CHECK-NEXT: 18.000000
+// CHECK-NEXT: 5.000000
+// CHECK-NEXT: 12.000000
+// CHECK-NEXT: 7.000000
+// CHECK-NEXT: 16.000000
+// CHECK-NEXT: 9.000000
// CHECK: s2 results:
// CHECK-NEXT: 0.000000
// CHECK-NEXT: 20.000000
diff --git a/offload/test/offloading/strided_update_variable_stride_misc.c b/offload/test/offloading/strided_update_variable_stride_misc.c
index 1efabb20db29b..d5579fe48887e 100644
--- a/offload/test/offloading/strided_update_variable_stride_misc.c
+++ b/offload/test/offloading/strided_update_variable_stride_misc.c
@@ -12,7 +12,7 @@ void test_1_variable_stride_one() {
// Initialize data on host
for (int i = 0; i < 10; i++) {
- data1[i] = i;
+ data1[i] = i + 1;
}
#pragma omp target data map(to : stride_one, data1[0 : 10])
@@ -20,7 +20,7 @@ void test_1_variable_stride_one() {
#pragma omp target
{
for (int i = 0; i < 10; i++) {
- data1[i] += i;
+ data1[i] += i + 1;
}
}
@@ -38,7 +38,7 @@ void test_2_variable_stride_large() {
// Initialize data on host
for (int i = 0; i < 10; i++) {
- data2[i] = i;
+ data2[i] = i + 1;
}
#pragma omp target data map(to : stride_large, data2[0 : 10])
@@ -46,7 +46,7 @@ void test_2_variable_stride_large() {
#pragma omp target
{
for (int i = 0; i < 10; i++) {
- data2[i] += i;
+ data2[i] += i + 1;
}
}
@@ -65,25 +65,25 @@ int main() {
}
// CHECK: Test 1: Variable stride = 1
-// CHECK-NEXT: 0.000000
-// CHECK-NEXT: 1.000000
// CHECK-NEXT: 2.000000
-// CHECK-NEXT: 3.000000
// CHECK-NEXT: 4.000000
-// CHECK-NEXT: 5.000000
// CHECK-NEXT: 6.000000
-// CHECK-NEXT: 7.000000
// CHECK-NEXT: 8.000000
-// CHECK-NEXT: 9.000000
+// CHECK-NEXT: 10.000000
+// CHECK-NEXT: 12.000000
+// CHECK-NEXT: 14.000000
+// CHECK-NEXT: 16.000000
+// CHECK-NEXT: 18.000000
+// CHECK-NEXT: 20.000000
// CHECK: Test 2: Variable stride = 5
-// CHECK-NEXT: 0.000000
-// CHECK-NEXT: 1.000000
+// CHECK-NEXT: 2.000000
// CHECK-NEXT: 2.000000
// CHECK-NEXT: 3.000000
// CHECK-NEXT: 4.000000
-// CHECK-NEXT: 10.000000
-// CHECK-NEXT: 6.000000
+// CHECK-NEXT: 5.000000
+// CHECK-NEXT: 12.000000
// CHECK-NEXT: 7.000000
// CHECK-NEXT: 8.000000
// CHECK-NEXT: 9.000000
+// CHECK-NEXT: 10.000000
More information about the cfe-commits
mailing list