llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-offload Author: Amit Tiwari (amitamd7) <details> <summary>Changes</summary> --- Patch is 99.53 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/176699.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 (+43-452) - (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..e6f5f00a86922 100644 --- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp +++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp @@ -8034,11 +8034,27 @@ class MappableExprsHandler { 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..418c6142380eb 100644 --- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp +++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp @@ -21,7 +21,6 @@ #include "llvm/Analysis/CodeMetrics.h" #include "llvm/Analysis/LoopInfo.h" #include "llvm/Analysis/OptimizationRemarkEmitter.h" -#include "llvm/Analysis/PostDominators.h" #include "llvm/Analysis/ScalarEvolution.h" #include "llvm/Analysis/TargetLibraryInfo.h" #include "llvm/Bitcode/BitcodeReader.h" @@ -50,7 +49,6 @@ #include "llvm/IR/Value.h" #include "llvm/MC/TargetRegistry.h" #include "llvm/Support/CommandLine.h" -#include "llvm/Support/Error.h" #include "llvm/Support/ErrorHandling.h" #include "llvm/Support/FileSystem.h" #include "llvm/Support/VirtualFileSystem.h" @@ -404,19 +402,18 @@ Value *createFakeIntVal(IRBuilderBase &Builder, OpenMPIRBuilder::InsertPointTy OuterAllocaIP, llvm::SmallVectorImpl<Instruction *> &ToBeDeleted, OpenMPIRBuilder::InsertPointTy InnerAllocaIP, - const Twine &Name = "", bool AsPtr = true, - bool Is64Bit = false) { + const Twine &Name = "", bool AsPtr = true) { Builder.restoreIP(OuterAllocaIP); - IntegerType *IntTy = Is64Bit ? Builder.getInt64Ty() : Builder.getInt32Ty(); Instruction *FakeVal; AllocaInst *FakeValAddr = - Builder.CreateAlloca(IntTy, nullptr, Name + ".addr"); + Builder.CreateAlloca(Builder.getInt32Ty(), nullptr, Name + ".addr"); ToBeDeleted.push_back(FakeValAddr); if (AsPtr) { FakeVal = FakeValAddr; } else { - FakeVal = Builder.CreateLoad(IntTy, FakeValAddr, Name + ".val"); + FakeVal = + Builder.CreateLoad(Builder.getInt32Ty(), FakeValAddr, Name + ".val"); ToBeDeleted.push_back(FakeVal); } @@ -424,10 +421,11 @@ Value *createFakeIntVal(IRBuilderBase &Builder, Builder.restoreIP(InnerAllocaIP); Instruction *UseFakeVal; if (AsPtr) { - UseFakeVal = Builder.CreateLoad(IntTy, FakeVal, Name + ".use"); + UseFakeVal = + Builder.CreateLoad(Builder.getInt32Ty(), FakeVal, Name + ".use"); } else { - UseFakeVal = cast<BinaryOperator>(Builder.CreateAdd( - FakeVal, Is64Bit ? Builder.getInt64(10) : Builder.getInt32(10))); + UseFakeVal = + cast<BinaryOperator>(Builder.CreateAdd(FakeVal, Builder.getInt32(10))); } ToBeDeleted.push_back(UseFakeVal); return FakeVal; @@ -765,28 +763,6 @@ static void raiseUserConstantDataAllocasToEntryBlock(IRBuilderBase &Builder, } } -static void hoistNonEntryAllocasToEntryBlock(llvm::BasicBlock &Block) { - llvm::SmallVector<llvm::Instruction *> AllocasToMove; - - auto ShouldHoistAlloca = [](const llvm::AllocaInst &AllocaInst) { - // TODO: For now, we support simple static allocations, we might need to - // move non-static ones as well. However, this will need further analysis to - // move the lenght arguments as well. - return !AllocaInst.isArrayAllocation(); - }; - - for (llvm::Instruction &Inst : Block) - if (auto *AllocaInst = llvm::dyn_cast<llvm::AllocaInst>(&Inst)) - if (ShouldHoistAlloca(*AllocaInst)) - AllocasToMove.push_back(AllocaInst); - - auto InsertPoint = - Block.getParent()->getEntryBlock().getTerminator()->getIterator(); - - for (llvm::Instruction *AllocaInst : AllocasToMove) - AllocaInst->moveBefore(InsertPoint); -} - void OpenMPIRBuilder::finalize(Function *Fn) { SmallPtrSet<BasicBlock *, 32> ParallelRegionBlockSet; Sma... [truncated] `````````` </details> https://github.com/llvm/llvm-project/pull/176699 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
