llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-offload @llvm/pr-subscribers-clang-codegen Author: Amit Tiwari (amitamd7) <details> <summary>Changes</summary> ### Issue 1: Dimension override missing with variable count expressions When variable count expressions were used with stride, the constant subsection path computed size first. This marked `ArgSizes` with byte size semantics. Variable expression logic later triggered, but reused `ArgSizes` assuming "bytes" semantics **Result**: `ArgSizes` wasn't overwritten with dimension count, breaking non-contiguous mapping. ### Issue 2: Variable stride not recognized as non-contiguous `CGOpenMPRuntime.cpp` failed to detect `DeclRefExpr`, `MemberExpr`, `ArraySubscriptExpr` as non-contiguous. Issue 3: Missing expression semantics in `OMPIRBuilder` `OMPIRBuilder.cpp` didn't handle dimension count for `OMP_MAP_NON_CONTIG` flag **Fixes**: `clang/lib/CodeGen/CGOpenMPRuntime.cpp` - Variable stride detection + dimension count logic `llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp` - Expression semantics for non-contiguous Detect variable stride expressions (`DeclRefExpr/MemberExpr/ArraySubscriptExpr`) as non-contiguous Set `OMP_MAP_NON_CONTIG` flag (`0x100000000000`) for variable stride/count. Generate 3D descriptor structures with runtime dimensions. Fix dimension override to use dimension count instead of byte size. Added testcases to cover stack arrays, heap pointers, struct members, etc., for expression semantics in non-contiguous update. --- Patch is 73.51 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/176708.diff 22 Files Affected: - (modified) clang/lib/CodeGen/CGOpenMPRuntime.cpp (+16) - (added) clang/test/OpenMP/target_update_strided_ptr_variable_count_and_stride_messages.c (+62) - (added) clang/test/OpenMP/target_update_strided_ptr_variable_count_messages.c (+57) - (added) clang/test/OpenMP/target_update_strided_ptr_variable_stride_messages.c (+64) - (added) clang/test/OpenMP/target_update_strided_struct_variable_count_and_stride_messages.c (+72) - (added) clang/test/OpenMP/target_update_variable_count_and_stride_messages.c (+85) - (modified) llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp (+21-8) - (added) offload/test/offloading/strided_update_count_expression.c (+133) - (added) offload/test/offloading/strided_update_count_expression_complex.c (+289) - (added) offload/test/offloading/strided_update_count_expression_misc.c (+99) - (added) offload/test/offloading/strided_update_multiple_arrays_count_expression.c (+161) - (added) offload/test/offloading/strided_update_multiple_arrays_variable_stride.c (+145) - (added) offload/test/offloading/strided_update_variable_count_and_stride.c (+136) - (added) offload/test/offloading/strided_update_variable_stride.c (+135) - (added) offload/test/offloading/strided_update_variable_stride_complex.c (+293) - (added) offload/test/offloading/strided_update_variable_stride_misc.c (+94) - (added) offload/test/offloading/target_update_ptr_count_expression.c (+99) - (added) offload/test/offloading/target_update_ptr_variable_count_and_stride.c (+94) - (added) offload/test/offloading/target_update_ptr_variable_stride.c (+95) - (added) offload/test/offloading/target_update_strided_struct_count_expression.c (+97) - (added) offload/test/offloading/target_update_strided_struct_variable_count_and_stride.c (+96) - (added) offload/test/offloading/target_update_strided_struct_variable_stride.c (+95) ``````````diff diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp index 8981a0de6d0e4..027d1fb26bc97 100644 --- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp +++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp @@ -8033,12 +8033,28 @@ class MappableExprsHandler { const Expr *StrideExpr = OASE->getStride(); if (!StrideExpr) return false; + + assert(StrideExpr->getType()->isIntegerType() && + "Stride expression must be of integer type"); + + // If the stride is a variable (not a constant), it's non-contiguous. + const Expr *S = StrideExpr->IgnoreParenImpCasts(); + if (const auto *DRE = dyn_cast<DeclRefExpr>(S)) { + if (isa<VarDecl>(DRE->getDecl()) || + isa<ParmVarDecl>(DRE->getDecl())) + return true; + } + if (isa<MemberExpr>(S) || isa<ArraySubscriptExpr>(S)) + return true; + // If stride is not evaluatable as a constant, treat as + // non-contiguous. const auto Constant = StrideExpr->getIntegerConstantExpr(CGF.getContext()); if (!Constant) return false; + // Treat non-unitary strides as non-contiguous. return !Constant->isOne(); }); diff --git a/clang/test/OpenMP/target_update_strided_ptr_variable_count_and_stride_messages.c b/clang/test/OpenMP/target_update_strided_ptr_variable_count_and_stride_messages.c new file mode 100644 index 0000000000000..932cd6b1c97bb --- /dev/null +++ b/clang/test/OpenMP/target_update_strided_ptr_variable_count_and_stride_messages.c @@ -0,0 +1,62 @@ +// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s -Wuninitialized +// RUN: %clang_cc1 -verify -fopenmp-simd -ferror-limit 100 %s -Wuninitialized + +int main(int argc, char **argv) { + int len = 16; + int count = 8; + int stride = 2; + int stride_large = 5; + double *data; + + // Valid strided array sections with both variable count and variable stride (FROM) + #pragma omp target update from(data[0:count:stride]) // OK - both variable + {} + + #pragma omp target update from(data[0:len/2:stride]) // OK - count expression, variable stride + {} + + #pragma omp target update from(data[0:count:stride_large]) // OK - variable count, different stride + {} + + #pragma omp target update from(data[1:len-2:stride]) // OK - with offset, count expression + {} + + #pragma omp target update from(data[0:count/2:stride*2]) // OK - both expressions + {} + + #pragma omp target update from(data[0:(len+1)/2:stride+1]) // OK - complex expressions + {} + + #pragma omp target update from(data[2:count-2:len/4]) // OK - all expressions + {} + + // Edge cases + int stride_one = 1; + #pragma omp target update from(data[0:count:stride_one]) // OK - variable count, stride=1 + {} + + #pragma omp target update from(data[0:len/stride:stride]) // OK - count depends on stride + {} + + // Invalid compile-time constant strides with variable count + #pragma omp target update from(data[0:count:0]) // expected-error {{section stride is evaluated to a non-positive value 0}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} + + #pragma omp target update from(data[0:len/2:-1]) // expected-error {{section stride is evaluated to a non-positive value -1}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} + + #pragma omp target update from(data[1:count:-2]) // expected-error {{section stride is evaluated to a non-positive value -2}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} + + // Valid strided array sections with variable count and stride (TO) + #pragma omp target update to(data[0:count:stride]) // OK + {} + + #pragma omp target update to(data[0:len/2:stride]) // OK + {} + + #pragma omp target update to(data[0:count:stride*2]) // OK + {} + + // Invalid stride with TO + #pragma omp target update to(data[0:count:0]) // expected-error {{section stride is evaluated to a non-positive value 0}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} + + return 0; +} diff --git a/clang/test/OpenMP/target_update_strided_ptr_variable_count_messages.c b/clang/test/OpenMP/target_update_strided_ptr_variable_count_messages.c new file mode 100644 index 0000000000000..23fba9c8bc84f --- /dev/null +++ b/clang/test/OpenMP/target_update_strided_ptr_variable_count_messages.c @@ -0,0 +1,57 @@ +// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s -Wuninitialized +// RUN: %clang_cc1 -verify -fopenmp-simd -ferror-limit 100 %s -Wuninitialized + +int main(int argc, char **argv) { + int len = 16; + int count = 8; + int divisor = 2; + double *data; + + // Valid strided array sections with variable count expressions (FROM) + #pragma omp target update from(data[0:count:2]) // OK - variable count + {} + + #pragma omp target update from(data[0:len/2:2]) // OK - count expression + {} + + #pragma omp target update from(data[0:len-4:3]) // OK - count with subtraction + {} + + #pragma omp target update from(data[1:(len+1)/2:2]) // OK - complex count expression + {} + + #pragma omp target update from(data[0:count*2:3]) // OK - count multiplication + {} + + #pragma omp target update from(data[2:len%divisor:2]) // OK - count with modulo + {} + + // Variable count with stride = 1 (contiguous) + #pragma omp target update from(data[0:count]) // OK - variable count, implicit stride + {} + + #pragma omp target update from(data[0:len/divisor]) // OK - expression count, implicit stride + {} + + // Invalid stride expressions with variable count + #pragma omp target update from(data[0:count:0]) // expected-error {{section stride is evaluated to a non-positive value 0}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} + + #pragma omp target update from(data[0:len/2:-1]) // expected-error {{section stride is evaluated to a non-positive value -1}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} + + #pragma omp target update from(data[1:count:-2]) // expected-error {{section stride is evaluated to a non-positive value -2}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} + + // Valid strided array sections with variable count expressions (TO) + #pragma omp target update to(data[0:count:2]) // OK + {} + + #pragma omp target update to(data[0:len/2:2]) // OK + {} + + #pragma omp target update to(data[0:len-4:3]) // OK + {} + + // Invalid stride with TO + #pragma omp target update to(data[0:count:0]) // expected-error {{section stride is evaluated to a non-positive value 0}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} + + return 0; +} diff --git a/clang/test/OpenMP/target_update_strided_ptr_variable_stride_messages.c b/clang/test/OpenMP/target_update_strided_ptr_variable_stride_messages.c new file mode 100644 index 0000000000000..3f85ed0c48d66 --- /dev/null +++ b/clang/test/OpenMP/target_update_strided_ptr_variable_stride_messages.c @@ -0,0 +1,64 @@ +// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s -Wuninitialized +// RUN: %clang_cc1 -verify -fopenmp-simd -ferror-limit 100 %s -Wuninitialized + +int main(int argc, char **argv) { + int len = 16; + int stride = 2; + int stride_large = 5; + double *data; + + // Valid strided array sections with variable stride (FROM) + #pragma omp target update from(data[0:8:stride]) // OK - variable stride + {} + + #pragma omp target update from(data[0:4:stride_large]) // OK - different variable stride + {} + + #pragma omp target update from(data[1:6:stride]) // OK - with offset + {} + + #pragma omp target update from(data[0:5:stride+1]) // OK - stride expression + {} + + #pragma omp target update from(data[0:4:stride*2]) // OK - stride multiplication + {} + + #pragma omp target update from(data[2:3:len/4]) // OK - stride from expression + {} + + // Edge case: stride = 1 (should be contiguous, not non-contiguous) + int stride_one = 1; + #pragma omp target update from(data[0:8:stride_one]) // OK - stride=1 is contiguous + {} + + // Invalid variable stride expressions + int zero_stride = 0; + int neg_stride = -1; + + // Note: These are runtime checks, so no compile-time error + #pragma omp target update from(data[0:8:zero_stride]) // OK at compile-time (runtime will fail) + {} + + #pragma omp target update from(data[0:4:neg_stride]) // OK at compile-time (runtime will fail) + {} + + // Compile-time constant invalid strides + #pragma omp target update from(data[0:4:0]) // expected-error {{section stride is evaluated to a non-positive value 0}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} + + #pragma omp target update from(data[0:4:-1]) // expected-error {{section stride is evaluated to a non-positive value -1}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} + + // Valid strided array sections with variable stride (TO) + #pragma omp target update to(data[0:8:stride]) // OK + {} + + #pragma omp target update to(data[0:5:stride+1]) // OK + {} + + #pragma omp target update to(data[0:4:stride*2]) // OK + {} + + // Invalid stride with TO + #pragma omp target update to(data[0:4:0]) // expected-error {{section stride is evaluated to a non-positive value 0}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} + + return 0; +} diff --git a/clang/test/OpenMP/target_update_strided_struct_variable_count_and_stride_messages.c b/clang/test/OpenMP/target_update_strided_struct_variable_count_and_stride_messages.c new file mode 100644 index 0000000000000..70775d5c8322c --- /dev/null +++ b/clang/test/OpenMP/target_update_strided_struct_variable_count_and_stride_messages.c @@ -0,0 +1,72 @@ +// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s -Wuninitialized +// RUN: %clang_cc1 -verify -fopenmp-simd -ferror-limit 100 %s -Wuninitialized + +#define N 20 +typedef struct { + double data[N]; + int len; + int stride; +} T; + +int main(int argc, char **argv) { + T s; + s.len = 16; + s.stride = 2; + int count = 8; + int ext_stride = 3; + + // Valid strided struct member array sections with variable count/stride (FROM) + #pragma omp target update from(s.data[0:s.len/2:2]) // OK - member count expression + {} + + #pragma omp target update from(s.data[0:count:s.stride]) // OK - external count, member stride + {} + + #pragma omp target update from(s.data[0:s.len:ext_stride]) // OK - member count, external stride + {} + + #pragma omp target update from(s.data[0:count:ext_stride]) // OK - both external + {} + + #pragma omp target update from(s.data[0:s.len/2:s.stride]) // OK - both from struct + {} + + #pragma omp target update from(s.data[1:(s.len-2)/2:s.stride]) // OK - complex count expression + {} + + #pragma omp target update from(s.data[0:count*2:s.stride+1]) // OK - expressions for both + {} + + // Edge cases + int stride_one = 1; + #pragma omp target update from(s.data[0:s.len:stride_one]) // OK - stride=1 + {} + + #pragma omp target update from(s.data[0:s.len/s.stride:s.stride]) // OK - count depends on stride + {} + + // Invalid compile-time constant strides with variable count + #pragma omp target update from(s.data[0:s.len:0]) // expected-error {{section stride is evaluated to a non-positive value 0}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} + + #pragma omp target update from(s.data[0:count:-1]) // expected-error {{section stride is evaluated to a non-positive value -1}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} + + #pragma omp target update from(s.data[1:s.len/2:-2]) // expected-error {{section stride is evaluated to a non-positive value -2}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} + + // Valid strided struct member array sections with variable count and stride (TO) + #pragma omp target update to(s.data[0:s.len/2:2]) // OK + {} + + #pragma omp target update to(s.data[0:count:s.stride]) // OK + {} + + #pragma omp target update to(s.data[0:s.len:ext_stride]) // OK + {} + + #pragma omp target update to(s.data[0:count*2:s.stride+1]) // OK + {} + + // Invalid stride with TO + #pragma omp target update to(s.data[0:s.len:0]) // expected-error {{section stride is evaluated to a non-positive value 0}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} + + return 0; +} diff --git a/clang/test/OpenMP/target_update_variable_count_and_stride_messages.c b/clang/test/OpenMP/target_update_variable_count_and_stride_messages.c new file mode 100644 index 0000000000000..0082539538a32 --- /dev/null +++ b/clang/test/OpenMP/target_update_variable_count_and_stride_messages.c @@ -0,0 +1,85 @@ +// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s -Wuninitialized +// RUN: %clang_cc1 -verify -fopenmp-simd -ferror-limit 100 %s -Wuninitialized + +int main(int argc, char **argv) { + int len = 16; + int count = 8; + int stride = 2; + int divisor = 2; + double data[100]; + + // Valid strided array sections with variable count expressions (FROM) + #pragma omp target update from(data[0:count:2]) // OK - variable count + {} + + #pragma omp target update from(data[0:len/2:2]) // OK - count expression + {} + + #pragma omp target update from(data[0:len-4:3]) // OK - count with subtraction + {} + + #pragma omp target update from(data[1:(len+1)/2:2]) // OK - complex count expression + {} + + #pragma omp target update from(data[0:count*2:3]) // OK - count multiplication + {} + + #pragma omp target update from(data[2:len%divisor:2]) // OK - count with modulo + {} + + // Variable stride with constant/variable count + #pragma omp target update from(data[0:10:stride]) // OK - constant count, variable stride + {} + + #pragma omp target update from(data[0:count:stride]) // OK - both variable + {} + + #pragma omp target update from(data[0:len/2:stride]) // OK - count expression, variable stride + {} + + #pragma omp target update from(data[0:count:stride*2]) // OK - variable count, stride expression + {} + + #pragma omp target update from(data[0:len/divisor:stride+1]) // OK - both expressions + {} + + // Variable count with stride = 1 (contiguous) + #pragma omp target update from(data[0:count]) // OK - variable count, implicit stride + {} + + #pragma omp target update from(data[0:len/divisor]) // OK - expression count, implicit stride + {} + + // Edge cases + int stride_one = 1; + #pragma omp target update from(data[0:len:stride_one]) // OK - stride=1 variable + {} + + #pragma omp target update from(data[0:len/stride:stride]) // OK - count depends on stride + {} + + // Invalid stride expressions with variable count + #pragma omp target update from(data[0:count:0]) // expected-error {{section stride is evaluated to a non-positive value 0}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} + + #pragma omp target update from(data[0:len/2:-1]) // expected-error {{section stride is evaluated to a non-positive value -1}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} + + #pragma omp target update from(data[1:count:-2]) // expected-error {{section stride is evaluated to a non-positive value -2}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} + + // Valid strided array sections with variable count expressions (TO) + #pragma omp target update to(data[0:count:2]) // OK + {} + + #pragma omp target update to(data[0:len/2:stride]) // OK + {} + + #pragma omp target update to(data[0:count:stride]) // OK + {} + + #pragma omp target update to(data[0:len/divisor:stride+1]) // OK + {} + + // Invalid stride with TO + #pragma omp target update to(data[0:count:0]) // expected-error {{section stride is evaluated to a non-positive value 0}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} + + return 0; +} diff --git a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp index 8d7a207a91f5a..578e63351ca5f 100644 --- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp +++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp @@ -9773,16 +9773,29 @@ Error OpenMPIRBuilder::emitOffloadingArrays( ConstantInt::get(Int64Ty, 0)); SmallBitVector RuntimeSizes(CombinedInfo.Sizes.size()); for (unsigned I = 0, E = CombinedInfo.Sizes.size(); I < E; ++I) { + bool IsNonContigEntry = + IsNonContiguous && + (static_cast<std::underlying_type_t<OpenMPOffloadMappingFlags>>( + CombinedInfo.Types[I] & + OpenMPOffloadMappingFlags::OMP_MAP_NON_CONTIG) != 0); + // For NON_CONTIG entries ArgSizes must carry the dimension count + // (number of descriptor_dim records) – NOT the byte size expression. + // Variable subsection forms (e.g. 0:s.len/2:2) previously produced a + // non-constant size so we marked them runtime and stored the byte size, + // leading the runtime to treat it as DimSize and overrun descriptors. + if (IsNonContigEntry) { + // Dims must be long enough and positive. + assert(I < CombinedInfo.NonContigInfo.Dims.size() && + "Induction variable is in-bounds with the NON_CONTIG Dims array"); + const uint64_t DimCount = CombinedInfo.NonContigInfo.Dims[I]; + assert(DimCount > 0 && "NON_CONTIG DimCount must be > 0"); + ConstSizes[I] = + ConstantInt::get(Int64Ty, CombinedInfo.NonContigInfo.Dims[I]); + continue; + } if (auto *CI = dyn_cast<Constant>(CombinedInfo.Sizes[I])) { if (!isa<ConstantExpr>(CI) && !isa<GlobalValue>(CI)) { - if (IsNonContiguous && - static_cast<std::underlying_type_t<OpenMPOffloadMappingFlags>>( - CombinedInfo.Types[I] & - OpenMPOffloadMappingFlags::OMP_MAP_NON_CONTIG)) - ConstSizes[I] = - ConstantInt::get(Int64Ty, CombinedInfo.NonContigInfo.Dims[I]); - else - ConstSizes[I] = CI; + ConstSizes[I] = CI; continue; } } diff --git a/offload/test/offloading/strided_update_count_expression.c b/offload/test/offloading/strided_update_count_expression.c new file mode 100644 index 0000000000000..a87da289a9154 --- /dev/null +++ b/offload/test/offloading/strided_update_count_expression.c @@ -0,0 +1,133 @@ +// This test checks that "update from" and "update to" clauses in OpenMP are +// supported when elements are updated in a non-contiguous manner with variable +// count expression. Tests #pragma omp target update from/to(data[0:len/2:2]) +// where the count (len/2) is a variable expression, not a constant. + +// RUN: %libomptarget-compile-run-and-check-generic +#include <omp.h> +#include <stdio.h> + +int main() { + int len = 10; + double data[len]; + + // ==================================================================== + // TEST 1: Update FROM device (device -> host) + // ==================================================================== + +#pragma omp target map(tofrom : len, data[0 : len]) + { + for (int i = 0; i < len; i++) { + data[i] = i; + } + } + + printf("Test 1: Update FROM device\n"); + printf("original host array values:\n"); + for (int i = 0; i < len; i++) + printf("%f\n", data[i]); + +#pragma omp target data map(to : len, data[0 : len]) + { +#pragma omp target + for (int i = 0; i < len; i++) { + data[i] += i; + } + +#pragma omp target update from(data[0 : len / 2 : 2... [truncated] `````````` </details> https://github.com/llvm/llvm-project/pull/176708 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
