[Clang][OpenMP] Fixed an issue of segment fault when using target nowait

[tianshilei1992]

Summary

The implementation of target nowait just wraps the target region into a task. The essential four parameters (base ptr, ptr, size, mapper) are taken as firstprivate such that they will be copied to the private location. When there is no user-defined mapper, the mapper variable will be nullptr. However, it will be still copied to the corresponding place. Therefore, a memcpy will be generated and the source pointer will be nullptr, causing a segmentation fault. The root cause is when calling emitOffloadingArraysArgument, the last argument Options has a field about whether it requires a task. It only takes depend clause into account. In this patch, the nowait clause is also included.

There're two things that will be done in another patches:

1. target data nowait has not been supported yet. D90099 added the support.
2. When there is no mapper, the mapper array can be nullptr no matter whether it requires outer task or not. It can avoid an unnecessary data copy. This is an optimization that is covered in D90101.

[jdoerfert]

Generally OK. Unclear if we want to keep a flag so we don't need to test the pointer, but I guess it is fine too.

[ye-luo]

Getting this even when compiling without offload. You can use the reproducer from the original bug report.

clang++: /home/yeluo/opt/llvm-clang/llvm-project/llvm/include/llvm/ADT/APInt.h:1151: bool llvm::APInt::operator==(const llvm::APInt &) const: Assertion `BitWidth == RHS.BitWidth && "Comparison requires equal bit widths"' failed.
PLEASE submit a bug report to https://bugs.llvm.org/ and include the crash backtrace, preprocessed source, and associated run script.
Stack dump:
0.	Program arguments: /home/packages/llvm/master-patched/bin/clang++ -DADD_ -DH5_USE_16_API -DHAVE_CONFIG_H -Drestrict=__restrict__ -I/home/yeluo/opt/miniqmc/src -I/home/yeluo/opt/miniqmc/build_clang_offlaod_nowait/src -fopenmp -fomit-frame-pointer -fstrict-aliasing -D__forceinline=inline -march=native -O3 -DNDEBUG -ffast-math -std=c++11 -o CMakeFiles/qmcutil.dir/Utilities/tinyxml/tinyxml2.cpp.o -c /home/yeluo/opt/miniqmc/src/Utilities/tinyxml/tinyxml2.cpp 
1.	<eof> parser at end of file
2.	Per-module optimization passes
3.	Running pass 'CallGraph Pass Manager' on module '/home/yeluo/opt/miniqmc/src/Utilities/tinyxml/tinyxml2.cpp'.
4.	Running pass 'Combine redundant instructions' on function '@_ZN8tinyxml27XMLUtil10IsNameCharEh'
 #0 0x0000000001ecc523 llvm::sys::PrintStackTrace(llvm::raw_ostream&, int) (/home/packages/llvm/master-patched/bin/clang+++0x1ecc523)
 #1 0x0000000001eca25e llvm::sys::RunSignalHandlers() (/home/packages/llvm/master-patched/bin/clang+++0x1eca25e)
 #2 0x0000000001ecb8cd llvm::sys::CleanupOnSignal(unsigned long) (/home/packages/llvm/master-patched/bin/clang+++0x1ecb8cd)
 #3 0x0000000001e513b3 (anonymous namespace)::CrashRecoveryContextImpl::HandleCrash(int, unsigned long) (/home/packages/llvm/master-patched/bin/clang+++0x1e513b3)
 #4 0x0000000001e514ee CrashRecoverySignalHandler(int) (/home/packages/llvm/master-patched/bin/clang+++0x1e514ee)
 #5 0x00007f18f56923c0 __restore_rt (/lib/x86_64-linux-gnu/libpthread.so.0+0x153c0)
 #6 0x00007f18f512718b raise /build/glibc-ZN95T4/glibc-2.31/signal/../sysdeps/unix/sysv/linux/raise.c:51:1
 #7 0x00007f18f5106859 abort /build/glibc-ZN95T4/glibc-2.31/stdlib/abort.c:81:7
 #8 0x00007f18f5106729 get_sysdep_segment_value /build/glibc-ZN95T4/glibc-2.31/intl/loadmsgcat.c:509:8
 #9 0x00007f18f5106729 _nl_load_domain /build/glibc-ZN95T4/glibc-2.31/intl/loadmsgcat.c:970:34
#10 0x00007f18f5117f36 (/lib/x86_64-linux-gnu/libc.so.6+0x36f36)
#11 0x00000000019f4c00 llvm::InstCombinerImpl::foldOrOfICmps(llvm::ICmpInst*, llvm::ICmpInst*, llvm::BinaryOperator&) (/home/packages/llvm/master-patched/bin/clang+++0x19f4c00)
#12 0x00000000019fb023 llvm::InstCombinerImpl::visitOr(llvm::BinaryOperator&) (/home/packages/llvm/master-patched/bin/clang+++0x19fb023)
#13 0x00000000019d354c llvm::InstCombinerImpl::run() (/home/packages/llvm/master-patched/bin/clang+++0x19d354c)
#14 0x00000000019d5788 combineInstructionsOverFunction(llvm::Function&, llvm::InstCombineWorklist&, llvm::AAResults*, llvm::AssumptionCache&, llvm::TargetLibraryInfo&, llvm::TargetTransformInfo&, llvm::DominatorTree&, llvm::OptimizationRemarkEmitter&, llvm::BlockFrequencyInfo*, llvm::ProfileSummaryInfo*, unsigned int, llvm::LoopInfo*) (/home/packages/llvm/master-patched/bin/clang+++0x19d5788)
#15 0x00000000019d70b1 llvm::InstructionCombiningPass::runOnFunction(llvm::Function&) (/home/packages/llvm/master-patched/bin/clang+++0x19d70b1)
#16 0x00000000017c7a68 llvm::FPPassManager::runOnFunction(llvm::Function&) (/home/packages/llvm/master-patched/bin/clang+++0x17c7a68)
#17 0x00000000010d0033 (anonymous namespace)::CGPassManager::runOnModule(llvm::Module&) (/home/packages/llvm/master-patched/bin/clang+++0x10d0033)
#18 0x00000000017c8117 llvm::legacy::PassManagerImpl::run(llvm::Module&) (/home/packages/llvm/master-patched/bin/clang+++0x17c8117)
#19 0x00000000020fed4a clang::EmitBackendOutput(clang::DiagnosticsEngine&, clang::HeaderSearchOptions const&, clang::CodeGenOptions const&, clang::TargetOptions const&, clang::LangOptions const&, llvm::DataLayout const&, llvm::Module*, clang::BackendAction, std::unique_ptr<llvm::raw_pwrite_stream, std::default_delete<llvm::raw_pwrite_stream> >) (/home/packages/llvm/master-patched/bin/clang+++0x20fed4a)
#20 0x0000000002d29c9c clang::BackendConsumer::HandleTranslationUnit(clang::ASTContext&) (/home/packages/llvm/master-patched/bin/clang+++0x2d29c9c)
#21 0x00000000037e77e3 clang::ParseAST(clang::Sema&, bool, bool) (/home/packages/llvm/master-patched/bin/clang+++0x37e77e3)
#22 0x00000000026dc383 clang::FrontendAction::Execute() (/home/packages/llvm/master-patched/bin/clang+++0x26dc383)
#23 0x000000000266e4f2 clang::CompilerInstance::ExecuteAction(clang::FrontendAction&) (/home/packages/llvm/master-patched/bin/clang+++0x266e4f2)
#24 0x0000000002789bb2 clang::ExecuteCompilerInvocation(clang::CompilerInstance*) (/home/packages/llvm/master-patched/bin/clang+++0x2789bb2)
#25 0x0000000000a4568c cc1_main(llvm::ArrayRef<char const*>, char const*, void*) (/home/packages/llvm/master-patched/bin/clang+++0xa4568c)
#26 0x0000000000a437ec ExecuteCC1Tool(llvm::SmallVectorImpl<char const*>&) (/home/packages/llvm/master-patched/bin/clang+++0xa437ec)
#27 0x0000000002523de2 void llvm::function_ref<void ()>::callback_fn<clang::driver::CC1Command::Execute(llvm::ArrayRef<llvm::Optional<llvm::StringRef> >, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> >*, bool*) const::$_1>(long) (/home/packages/llvm/master-patched/bin/clang+++0x2523de2)
#28 0x0000000001e512c7 llvm::CrashRecoveryContext::RunSafely(llvm::function_ref<void ()>) (/home/packages/llvm/master-patched/bin/clang+++0x1e512c7)
#29 0x00000000025234f7 clang::driver::CC1Command::Execute(llvm::ArrayRef<llvm::Optional<llvm::StringRef> >, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> >*, bool*) const (/home/packages/llvm/master-patched/bin/clang+++0x25234f7)
#30 0x00000000024efd28 clang::driver::Compilation::ExecuteCommand(clang::driver::Command const&, clang::driver::Command const*&) const (/home/packages/llvm/master-patched/bin/clang+++0x24efd28)
#31 0x00000000024f0247 clang::driver::Compilation::ExecuteJobs(clang::driver::JobList const&, llvm::SmallVectorImpl<std::pair<int, clang::driver::Command const*> >&) const (/home/packages/llvm/master-patched/bin/clang+++0x24f0247)
#32 0x0000000002509758 clang::driver::Driver::ExecuteCompilation(clang::driver::Compilation&, llvm::SmallVectorImpl<std::pair<int, clang::driver::Command const*> >&) (/home/packages/llvm/master-patched/bin/clang+++0x2509758)
#33 0x0000000000a43158 main (/home/packages/llvm/master-patched/bin/clang+++0xa43158)
#34 0x00007f18f51080b3 __libc_start_main /build/glibc-ZN95T4/glibc-2.31/csu/../csu/libc-start.c:342:3
#35 0x0000000000a404de _start (/home/packages/llvm/master-patched/bin/clang+++0xa404de)
clang-12: error: clang frontend command failed with exit code 134 (use -v to see invocation)
clang version 12.0.0 (https://github.com/llvm/llvm-project.git ca73dcd8a9ed9cc3ca1c1cc97ab893747791a681)
Target: x86_64-unknown-linux-gnu
Thread model: posix
InstalledDir: /home/packages/llvm/master-patched/bin
clang-12: note: diagnostic msg: 
********************

PLEASE ATTACH THE FOLLOWING FILES TO THE BUG REPORT:
Preprocessed source(s) and associated run script(s) are located at:
clang-12: note: diagnostic msg: /tmp/tinyxml2-c2195c.cpp
clang-12: note: diagnostic msg: /tmp/tinyxml2-c2195c.sh
clang-12: note: diagnostic msg: 

********************

[tianshilei1992]

In D89844#2345337, @ye-luo wrote:

Getting this even when compiling without offload. You can use the reproducer from the original bug report.

I noticed one issue and I'm working on it now.

[tianshilei1992]

Fixed an issue that when there is mapper, the variable in the generated wrapped
task still uses the one in the encountering function

[jdoerfert]

can we add the reproducer to the tests please.

[tianshilei1992]

Tried to fix some cases

[tianshilei1992]

Removed useless code and fixed some cases

[tianshilei1992]

Fixed the issue using a way that no need to change test cases

[tianshilei1992]

Reverted changes for target data. It will be in another patch.

[tianshilei1992]

I don't think the failed case has anything to do with this patch. It must be a random issue in the bolt.

[tianshilei1992]

Added a test case

[jdoerfert]

LGTM

[mikerice]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

9670

It seems this change is no longer in the compiler, as it was reverted by the next commit to this file. Was that intentional? See: https://reviews.llvm.org/D90101.

[tianshilei1992]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

9670

Yeah, it is not needed anymore.

clang/lib/CodeGen/CGOpenMPRuntime.cpp

auto &&TargetThenGen = [this, &ThenGen, &D, &InputInfo, &MapTypesArray,
     // Map any list items in a map clause that were not captures because they
     // weren't referenced within the construct.
     MEHandler.generateAllInfo(CombinedInfo, /*NotTargetParams=*/true,
                               MappedVarSet);
 
     TargetDataInfo Info;
     // Fill up the arrays and create the arguments.
     emitOffloadingArrays(CGF, CombinedInfo, Info);
-    bool HasDependClauses = D.hasClausesOfKind<OMPDependClause>();
     emitOffloadingArraysArgument(CGF, Info.BasePointersArray,
                                  Info.PointersArray, Info.SizesArray,
                                  Info.MapTypesArray, Info.MappersArray, Info,
-                                 {/*ForEndTask=*/false, HasDependClauses});
+                                 {/*ForEndTask=*/false, RequiresOuterTask});
     InputInfo.NumberOfTargetItems = Info.NumberOfPtrs;
     InputInfo.BasePointersArray =
         Address(Info.BasePointersArray, CGM.getPointerAlign());
     InputInfo.PointersArray =
         Address(Info.PointersArray, CGM.getPointerAlign());
     InputInfo.SizesArray = Address(Info.SizesArray, CGM.getPointerAlign());
     InputInfo.MappersArray = Address(Info.MappersArray, CGM.getPointerAlign());
     MapTypesArray = Info.MapTypesArray;

openmp/libomptarget/test/offloading/bug47654.cpp

+// RUN: %libomptarget-compilexx-run-and-check-aarch64-unknown-linux-gnu
+// RUN: %libomptarget-compilexx-run-and-check-powerpc64-ibm-linux-gnu
+// RUN: %libomptarget-compilexx-run-and-check-powerpc64le-ibm-linux-gnu
+// RUN: %libomptarget-compilexx-run-and-check-x86_64-pc-linux-gnu
+// RUN: %libomptarget-compilexx-run-and-check-nvptx64-nvidia-cuda
+
+#include <cassert>
+#include <iostream>
+
+int main(int argc, char *argv[]) {
+  int i = 0, j = 0;
+
+#pragma omp target map(tofrom : i, j) nowait
+  {
+    i = 1;
+    j = 2;
+  }
+
+#pragma omp taskwait
+
+  assert(i == 1);
+  assert(j == 2);
+
+  std::cout << "PASS\n";
+
+  return 0;
+}
+
+// CHECK: PASS