Skip to content

Conversation

anutosh491
Copy link
Member

clang-repl --cuda was previously crashing with a segmentation fault, instead of reporting a clean error

(base) anutosh491@Anutoshs-MacBook-Air bin % ./clang-repl --cuda
#0 0x0000000111da4fbc llvm::sys::PrintStackTrace(llvm::raw_ostream&, int) (/opt/local/libexec/llvm-20/lib/libLLVM.dylib+0x150fbc)
#1 0x0000000111da31dc llvm::sys::RunSignalHandlers() (/opt/local/libexec/llvm-20/lib/libLLVM.dylib+0x14f1dc)
#2 0x0000000111da5628 SignalHandler(int) (/opt/local/libexec/llvm-20/lib/libLLVM.dylib+0x151628)
#3 0x000000019b242de4 (/usr/lib/system/libsystem_platform.dylib+0x180482de4)
#4 0x0000000107f638d0 clang::IncrementalCUDADeviceParser::IncrementalCUDADeviceParser(std::__1::unique_ptr<clang::CompilerInstance, std::__1::default_delete<clang::CompilerInstance>>, clang::CompilerInstance&, llvm::IntrusiveRefCntPtr<llvm::vfs::InMemoryFileSystem>, llvm::Error&, std::__1::list<clang::PartialTranslationUnit, std::__1::allocator<clang::PartialTranslationUnit>> const&) (/opt/local/libexec/llvm-20/lib/libclang-cpp.dylib+0x216b8d0)
#5 0x0000000107f638d0 clang::IncrementalCUDADeviceParser::IncrementalCUDADeviceParser(std::__1::unique_ptr<clang::CompilerInstance, std::__1::default_delete<clang::CompilerInstance>>, clang::CompilerInstance&, llvm::IntrusiveRefCntPtr<llvm::vfs::InMemoryFileSystem>, llvm::Error&, std::__1::list<clang::PartialTranslationUnit, std::__1::allocator<clang::PartialTranslationUnit>> const&) (/opt/local/libexec/llvm-20/lib/libclang-cpp.dylib+0x216b8d0)
#6 0x0000000107f6bac8 clang::Interpreter::createWithCUDA(std::__1::unique_ptr<clang::CompilerInstance, std::__1::default_delete<clang::CompilerInstance>>, std::__1::unique_ptr<clang::CompilerInstance, std::__1::default_delete<clang::CompilerInstance>>) (/opt/local/libexec/llvm-20/lib/libclang-cpp.dylib+0x2173ac8)
#7 0x000000010206f8a8 main (/opt/local/libexec/llvm-20/bin/clang-repl+0x1000038a8)
#8 0x000000019ae8c274 
Segmentation fault: 11

The underlying issue was that the DeviceCompilerInstance (used for device-side CUDA compilation) was never initialized with a Sema, which is required before constructing the IncrementalCUDADeviceParser.

: IncrementalParser(*DeviceInstance, Err), PTUs(PTUs), VFS(FS),

Unlike the host-side CompilerInstance which runs ExecuteAction inside the Interpreter constructor (thereby setting up Sema), the device-side CI was passed into the parser uninitialized, leading to an assertion or crash when accessing its internals.

To fix this, I refactored the Interpreter::create method to include an optional DeviceCI parameter. If provided, we know we need to take care of this instance too. Only then do we construct the IncrementalCUDADeviceParser.

@llvmbot llvmbot added clang Clang issues not falling into any other category clang:frontend Language frontend issues, e.g. anything involving "Sema" labels Apr 19, 2025
@llvmbot
Copy link
Member

llvmbot commented Apr 19, 2025

@llvm/pr-subscribers-clang

Author: Anutosh Bhat (anutosh491)

Changes

clang-repl --cuda was previously crashing with a segmentation fault, instead of reporting a clean error

(base) anutosh491@<!-- -->Anutoshs-MacBook-Air bin % ./clang-repl --cuda
#<!-- -->0 0x0000000111da4fbc llvm::sys::PrintStackTrace(llvm::raw_ostream&amp;, int) (/opt/local/libexec/llvm-20/lib/libLLVM.dylib+0x150fbc)
#<!-- -->1 0x0000000111da31dc llvm::sys::RunSignalHandlers() (/opt/local/libexec/llvm-20/lib/libLLVM.dylib+0x14f1dc)
#<!-- -->2 0x0000000111da5628 SignalHandler(int) (/opt/local/libexec/llvm-20/lib/libLLVM.dylib+0x151628)
#<!-- -->3 0x000000019b242de4 (/usr/lib/system/libsystem_platform.dylib+0x180482de4)
#<!-- -->4 0x0000000107f638d0 clang::IncrementalCUDADeviceParser::IncrementalCUDADeviceParser(std::__1::unique_ptr&lt;clang::CompilerInstance, std::__1::default_delete&lt;clang::CompilerInstance&gt;&gt;, clang::CompilerInstance&amp;, llvm::IntrusiveRefCntPtr&lt;llvm::vfs::InMemoryFileSystem&gt;, llvm::Error&amp;, std::__1::list&lt;clang::PartialTranslationUnit, std::__1::allocator&lt;clang::PartialTranslationUnit&gt;&gt; const&amp;) (/opt/local/libexec/llvm-20/lib/libclang-cpp.dylib+0x216b8d0)
#<!-- -->5 0x0000000107f638d0 clang::IncrementalCUDADeviceParser::IncrementalCUDADeviceParser(std::__1::unique_ptr&lt;clang::CompilerInstance, std::__1::default_delete&lt;clang::CompilerInstance&gt;&gt;, clang::CompilerInstance&amp;, llvm::IntrusiveRefCntPtr&lt;llvm::vfs::InMemoryFileSystem&gt;, llvm::Error&amp;, std::__1::list&lt;clang::PartialTranslationUnit, std::__1::allocator&lt;clang::PartialTranslationUnit&gt;&gt; const&amp;) (/opt/local/libexec/llvm-20/lib/libclang-cpp.dylib+0x216b8d0)
#<!-- -->6 0x0000000107f6bac8 clang::Interpreter::createWithCUDA(std::__1::unique_ptr&lt;clang::CompilerInstance, std::__1::default_delete&lt;clang::CompilerInstance&gt;&gt;, std::__1::unique_ptr&lt;clang::CompilerInstance, std::__1::default_delete&lt;clang::CompilerInstance&gt;&gt;) (/opt/local/libexec/llvm-20/lib/libclang-cpp.dylib+0x2173ac8)
#<!-- -->7 0x000000010206f8a8 main (/opt/local/libexec/llvm-20/bin/clang-repl+0x1000038a8)
#<!-- -->8 0x000000019ae8c274 
Segmentation fault: 11

The underlying issue was that the DeviceCompilerInstance (used for device-side CUDA compilation) was never initialized with a Sema, which is required before constructing the IncrementalCUDADeviceParser.

: IncrementalParser(*DeviceInstance, Err), PTUs(PTUs), VFS(FS),

Unlike the host-side CompilerInstance which runs ExecuteAction inside the Interpreter constructor (thereby setting up Sema), the device-side CI was passed into the parser uninitialized, leading to an assertion or crash when accessing its internals.

To fix this, I refactored the Interpreter::create method to include an optional DeviceCI parameter. If provided, we know we need to take care of this instance too. Only then do we construct the IncrementalCUDADeviceParser.


Full diff: https://github.com/llvm/llvm-project/pull/136404.diff

3 Files Affected:

  • (modified) clang/include/clang/Interpreter/Interpreter.h (+2-1)
  • (modified) clang/lib/Interpreter/DeviceOffload.cpp (+2-1)
  • (modified) clang/lib/Interpreter/Interpreter.cpp (+33-15)
diff --git a/clang/include/clang/Interpreter/Interpreter.h b/clang/include/clang/Interpreter/Interpreter.h
index b1b63aedf86ab..20d22c4437a97 100644
--- a/clang/include/clang/Interpreter/Interpreter.h
+++ b/clang/include/clang/Interpreter/Interpreter.h
@@ -129,7 +129,8 @@ class Interpreter {
 public:
   virtual ~Interpreter();
   static llvm::Expected<std::unique_ptr<Interpreter>>
-  create(std::unique_ptr<CompilerInstance> CI);
+  create(std::unique_ptr<CompilerInstance> CI,
+    std::unique_ptr<CompilerInstance> DeviceCI = nullptr);
   static llvm::Expected<std::unique_ptr<Interpreter>>
   createWithCUDA(std::unique_ptr<CompilerInstance> CI,
                  std::unique_ptr<CompilerInstance> DCI);
diff --git a/clang/lib/Interpreter/DeviceOffload.cpp b/clang/lib/Interpreter/DeviceOffload.cpp
index 1999d63d1aa04..9a7be006250a0 100644
--- a/clang/lib/Interpreter/DeviceOffload.cpp
+++ b/clang/lib/Interpreter/DeviceOffload.cpp
@@ -34,14 +34,15 @@ IncrementalCUDADeviceParser::IncrementalCUDADeviceParser(
       TargetOpts(HostInstance.getTargetOpts()) {
   if (Err)
     return;
-  DeviceCI = std::move(DeviceInstance);
   StringRef Arch = TargetOpts.CPU;
   if (!Arch.starts_with("sm_") || Arch.substr(3).getAsInteger(10, SMVersion)) {
+    DeviceInstance.release();
     Err = llvm::joinErrors(std::move(Err), llvm::make_error<llvm::StringError>(
                                                "Invalid CUDA architecture",
                                                llvm::inconvertibleErrorCode()));
     return;
   }
+  DeviceCI = std::move(DeviceInstance);
 }
 
 llvm::Expected<TranslationUnitDecl *>
diff --git a/clang/lib/Interpreter/Interpreter.cpp b/clang/lib/Interpreter/Interpreter.cpp
index f8c8d0a425659..1af3f24774028 100644
--- a/clang/lib/Interpreter/Interpreter.cpp
+++ b/clang/lib/Interpreter/Interpreter.cpp
@@ -451,13 +451,44 @@ const char *const Runtimes = R"(
 )";
 
 llvm::Expected<std::unique_ptr<Interpreter>>
-Interpreter::create(std::unique_ptr<CompilerInstance> CI) {
+Interpreter::create(std::unique_ptr<CompilerInstance> CI,
+                    std::unique_ptr<CompilerInstance> DeviceCI) {
   llvm::Error Err = llvm::Error::success();
   auto Interp =
       std::unique_ptr<Interpreter>(new Interpreter(std::move(CI), Err));
   if (Err)
     return std::move(Err);
 
+  if (DeviceCI) {
+    // auto DeviceLLVMCtx = std::make_unique<llvm::LLVMContext>();
+    // auto DeviceTSCtx =
+    //     std::make_unique<llvm::orc::ThreadSafeContext>(std::move(DeviceLLVMCtx));
+
+    // llvm::Error DeviceErr = llvm::Error::success();
+    // llvm::ErrorAsOutParameter EAO(&DeviceErr);
+
+    // auto DeviceAct = std::make_unique<IncrementalAction>(
+    //     *DeviceCI, *DeviceTSCtx->getContext(), DeviceErr, *Interp);
+
+    // if (DeviceErr)
+    //   return std::move(DeviceErr);
+
+    //DeviceCI->ExecuteAction(*DeviceAct);
+    DeviceCI->ExecuteAction(*Interp->Act);
+
+    llvm::IntrusiveRefCntPtr<llvm::vfs::InMemoryFileSystem> IMVFS =
+        std::make_unique<llvm::vfs::InMemoryFileSystem>();
+
+    auto DeviceParser = std::make_unique<IncrementalCUDADeviceParser>(
+        std::move(DeviceCI), *Interp->getCompilerInstance(), IMVFS, Err,
+        Interp->PTUs);
+
+    if (Err)
+      return std::move(Err);
+
+    Interp->DeviceParser = std::move(DeviceParser);
+  }
+
   // Add runtime code and set a marker to hide it from user code. Undo will not
   // go through that.
   auto PTU = Interp->Parse(Runtimes);
@@ -481,20 +512,7 @@ Interpreter::createWithCUDA(std::unique_ptr<CompilerInstance> CI,
   OverlayVFS->pushOverlay(IMVFS);
   CI->createFileManager(OverlayVFS);
 
-  auto Interp = Interpreter::create(std::move(CI));
-  if (auto E = Interp.takeError())
-    return std::move(E);
-
-  llvm::Error Err = llvm::Error::success();
-  auto DeviceParser = std::make_unique<IncrementalCUDADeviceParser>(
-      std::move(DCI), *(*Interp)->getCompilerInstance(), IMVFS, Err,
-      (*Interp)->PTUs);
-  if (Err)
-    return std::move(Err);
-
-  (*Interp)->DeviceParser = std::move(DeviceParser);
-
-  return Interp;
+  return Interpreter::create(std::move(CI), std::move(DCI));
 }
 
 const CompilerInstance *Interpreter::getCompilerInstance() const {

Copy link

github-actions bot commented Apr 19, 2025

✅ With the latest revision this PR passed the C/C++ code formatter.

@anutosh491
Copy link
Member Author

anutosh491 commented Apr 19, 2025

cc @AaronBallman @vgvassilev

This fixes this discussion y'all had #107737 (comment)

Unfortunately I am on an Apple m1 so I can just justify seeing this on my side

(base) anutosh491@Anutoshs-MacBook-Air bin % ./clang-repl --cuda
clang-repl: Invalid CUDA architecture

But actually this error is more justified that what we are seeing with clang-repl-19

(base) anutosh491@Anutoshs-MacBook-Air bin % ./clang-repl-19 --cuda
clang-repl: dlopen(libcudart.so, 0x0009): tried: 'libcudart.so' (no such file), '/System/Volumes/Preboot/Cryptexes/OSlibcudart.so' (no such file), '/opt/local/libexec/llvm-19/lib/../lib/libcudart.so' (no such file), '/opt/local/libexec/llvm-19/bin/../lib/libcudart.so' (no such file), '/usr/lib/libcudart.so' (no such file, not in dyld cache), 'libcudart.so' (no such file)

So everything spawns from here

if (CudaEnabled) {
Interp = ExitOnErr(
clang::Interpreter::createWithCUDA(std::move(CI), std::move(DeviceCI)));
if (CudaPath.empty()) {
ExitOnErr(Interp->LoadDynamicLibrary("libcudart.so"));

But for
i) clang-repl-20: The call to createwithcuda would error out from here

The first one comes from here

StringRef Arch = TargetOpts.CPU;
if (!Arch.starts_with("sm_") || Arch.substr(3).getAsInteger(10, SMVersion)) {
Err = llvm::joinErrors(std::move(Err), llvm::make_error<llvm::StringError>(
"Invalid CUDA architecture",
llvm::inconvertibleErrorCode()));
return;

ii) clang-repl-19: Passes that, successfully creates the interpreter and ends up failing at the Loading step (although we should expect it to fail on the top itself isn't it ?)

@anutosh491 anutosh491 changed the title [clan-repl] : Fix clang-repl crash with --cuda flag [clang-repl] : Fix clang-repl crash with --cuda flag Apr 19, 2025
@anutosh491
Copy link
Member Author

Hey @vgvassilev @argentite

With these changes, I am positive we atleast have the design in place to run any of these Cuda tests (https://github.com/llvm/llvm-project/tree/main/clang/test/Interpreter/CUDA)

For example running the sanity.cu file

anutosh491@vv-nuc:/build/anutosh491/llvm-project/build/bin$ ./clang-repl --cuda
clang-repl> extern "C" int printf(const char*, ...);
clang-repl> __global__ void test_func() {}
clang-repl> test_func<<<1,1>>>();
clang-repl> printf("CUDA Error: %d", cudaGetLastError());
CUDA Error: 0
clang-repl> 

The only thing which I am kinda struggling to figure out is how to put the virtual filesystem to use. So to get the above tests working, I had to just comment out what we do through VFS and use the real filesystem

  // VFS->addFile(FatbinFileName, 0,
  //              llvm::MemoryBuffer::getMemBuffer(
  //                  llvm::StringRef(FatbinContent.data(), FatbinContent.size()),
  //                  "", false));
  // Write directly to disk
  std::error_code EC;
  llvm::raw_fd_ostream OS(FatbinFileName, EC);
  if (EC) {
    return llvm::make_error<llvm::StringError>(
        "Failed to write fatbin: " + EC.message(), llvm::inconvertibleErrorCode());
  }
  OS.write(FatbinContent.data(), FatbinContent.size());
  OS.close();

So hopefully the VFS used here

VFS->addFile(FatbinFileName, 0,

and here

if (!CudaGpuBinaryFileName.empty()) {
auto VFS = CGM.getFileSystem();

Should be the same, but my logs tell me otherwise. I added some simple logs to verify the same and I see this

[DeviceOffload] Current working directory: /
[DeviceOffload] Fatbin file successfully found in VFS: /incr_module_3.fatbin
[CGCUDANV] CudaGpuBinaryFileName: /incr_module_3.fatbin
[CGCUDANV] VFS CWD: /build/anutosh491/llvm-project/build/bin
[CGCUDANV] status() failed: No such file or directory

So if you see both VFS point to different locations.
I tried some workarounds but haven't been able to figure this out and possibly might need some help.

So the current changes would just fail with

anutosh491@vv-nuc:/build/anutosh491/llvm-project/build/bin$ ./clang-repl --cuda
clang-repl> extern "C" int printf(const char*, ...);
clang-repl> __global__ void test_func() {}
fatal error: cannot open file '/incr_module_3.fatbin': No such file or directory
Stack dump without symbol names (ensure you have llvm-symbolizer in your PATH or set the environment var `LLVM_SYMBOLIZER_PATH` to point to it):
0  clang-repl 0x0000555557fde34f llvm::sys::PrintStackTrace(llvm::raw_ostream&, int) + 63
1  clang-repl 0x0000555557fdc264 llvm::sys::RunSignalHandlers() + 52
2  clang-repl 0x0000555557fdc5b9
3  libc.so.6  0x00007ffff795c520
4  clang-repl 0x00005555581608e7 clang::Interpreter::GenModule(clang::IncrementalAction*) + 519
5  clang-repl 0x0000555558160dee clang::Interpreter::RegisterPTU(clang::TranslationUnitDecl*, std::unique_ptr<llvm::Module, std::default_delete<llvm::Module>>, clang::IncrementalAction*) + 206
6  clang-repl 0x0000555558161027 clang::Interpreter::Parse(llvm::StringRef) + 471
7  clang-repl 0x0000555558162360 clang::Interpreter::ParseAndExecute(llvm::StringRef, clang::Value*) + 48
8  clang-repl 0x0000555556263075 main + 4085
9  libc.so.6  0x00007ffff7943d90
10 libc.so.6  0x00007ffff7943e40 __libc_start_main + 128
11 clang-repl 0x00005555563da585 _start + 37
Segmentation fault (core dumped)

Once we start fetching these files from the correct location, this should be ready !

@anutosh491
Copy link
Member Author

Was able to address the issue here.

Making a commit soon !

@anutosh491 anutosh491 force-pushed the cuda branch 2 times, most recently from c4c00cb to 6c64e64 Compare April 24, 2025 12:55
@anutosh491
Copy link
Member Author

Making a note of the major changes below !

@anutosh491 anutosh491 requested a review from vgvassilev April 24, 2025 16:26
@anutosh491 anutosh491 force-pushed the cuda branch 3 times, most recently from bbf41f4 to d10dab4 Compare April 25, 2025 19:28
Copy link
Contributor

@vgvassilev vgvassilev left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

LGTM!

@vgvassilev vgvassilev merged commit 21fb19f into llvm:main Apr 26, 2025
11 checks passed
@llvm-ci
Copy link
Collaborator

llvm-ci commented Apr 26, 2025

LLVM Buildbot has detected a new failure on builder lldb-arm-ubuntu running on linaro-lldb-arm-ubuntu while building clang at step 6 "test".

Full details are available at: https://lab.llvm.org/buildbot/#/builders/18/builds/15088

Here is the relevant piece of the build log for the reference
Step 6 (test) failure: build (failure)
...
PASS: lldb-api :: symbol_ondemand/shared_library/TestSharedLibOnDemand.py (1146 of 3001)
PASS: lldb-api :: terminal/TestSTTYBeforeAndAfter.py (1147 of 3001)
PASS: lldb-api :: test_utils/TestDecorators.py (1148 of 3001)
PASS: lldb-api :: test_utils/TestInlineTest.py (1149 of 3001)
PASS: lldb-api :: test_utils/TestPExpectTest.py (1150 of 3001)
PASS: lldb-api :: test_utils/base/TestBaseTest.py (1151 of 3001)
PASS: lldb-api :: python_api/watchpoint/watchlocation/TestTargetWatchAddress.py (1152 of 3001)
PASS: lldb-api :: tools/lldb-dap/attach/TestDAP_attachByPortNum.py (1153 of 3001)
UNSUPPORTED: lldb-api :: tools/lldb-dap/breakpoint-events/TestDAP_breakpointEvents.py (1154 of 3001)
PASS: lldb-api :: terminal/TestEditline.py (1155 of 3001)
FAIL: lldb-api :: tools/lldb-dap/attach/TestDAP_attach.py (1156 of 3001)
******************** TEST 'lldb-api :: tools/lldb-dap/attach/TestDAP_attach.py' FAILED ********************
Script:
--
/usr/bin/python3.10 /home/tcwg-buildbot/worker/lldb-arm-ubuntu/llvm-project/lldb/test/API/dotest.py -u CXXFLAGS -u CFLAGS --env LLVM_LIBS_DIR=/home/tcwg-buildbot/worker/lldb-arm-ubuntu/build/./lib --env LLVM_INCLUDE_DIR=/home/tcwg-buildbot/worker/lldb-arm-ubuntu/build/include --env LLVM_TOOLS_DIR=/home/tcwg-buildbot/worker/lldb-arm-ubuntu/build/./bin --arch armv8l --build-dir /home/tcwg-buildbot/worker/lldb-arm-ubuntu/build/lldb-test-build.noindex --lldb-module-cache-dir /home/tcwg-buildbot/worker/lldb-arm-ubuntu/build/lldb-test-build.noindex/module-cache-lldb/lldb-api --clang-module-cache-dir /home/tcwg-buildbot/worker/lldb-arm-ubuntu/build/lldb-test-build.noindex/module-cache-clang/lldb-api --executable /home/tcwg-buildbot/worker/lldb-arm-ubuntu/build/./bin/lldb --compiler /home/tcwg-buildbot/worker/lldb-arm-ubuntu/build/./bin/clang --dsymutil /home/tcwg-buildbot/worker/lldb-arm-ubuntu/build/./bin/dsymutil --make /usr/bin/gmake --llvm-tools-dir /home/tcwg-buildbot/worker/lldb-arm-ubuntu/build/./bin --lldb-obj-root /home/tcwg-buildbot/worker/lldb-arm-ubuntu/build/tools/lldb --lldb-libs-dir /home/tcwg-buildbot/worker/lldb-arm-ubuntu/build/./lib /home/tcwg-buildbot/worker/lldb-arm-ubuntu/llvm-project/lldb/test/API/tools/lldb-dap/attach -p TestDAP_attach.py
--
Exit Code: 1

Command Output (stdout):
--
lldb version 21.0.0git (https://github.com/llvm/llvm-project.git revision 21fb19f3b5d572f608e959af895d781b9b24fbbd)
  clang revision 21fb19f3b5d572f608e959af895d781b9b24fbbd
  llvm revision 21fb19f3b5d572f608e959af895d781b9b24fbbd
Skipping the following test categories: ['libc++', 'dsym', 'gmodules', 'debugserver', 'objc']

--
Command Output (stderr):
--
========= DEBUG ADAPTER PROTOCOL LOGS =========
1745651006.068298101 --> (stdin/stdout) {"command":"initialize","type":"request","arguments":{"adapterID":"lldb-native","clientID":"vscode","columnsStartAt1":true,"linesStartAt1":true,"locale":"en-us","pathFormat":"path","supportsRunInTerminalRequest":true,"supportsVariablePaging":true,"supportsVariableType":true,"supportsStartDebuggingRequest":true,"supportsProgressReporting":true,"$__lldb_sourceInitFile":false},"seq":1}
1745651006.072639704 <-- (stdin/stdout) {"body":{"$__lldb_version":"lldb version 21.0.0git (https://github.com/llvm/llvm-project.git revision 21fb19f3b5d572f608e959af895d781b9b24fbbd)\n  clang revision 21fb19f3b5d572f608e959af895d781b9b24fbbd\n  llvm revision 21fb19f3b5d572f608e959af895d781b9b24fbbd","completionTriggerCharacters":["."," ","\t"],"exceptionBreakpointFilters":[{"default":false,"filter":"cpp_catch","label":"C++ Catch"},{"default":false,"filter":"cpp_throw","label":"C++ Throw"},{"default":false,"filter":"objc_catch","label":"Objective-C Catch"},{"default":false,"filter":"objc_throw","label":"Objective-C Throw"}],"supportTerminateDebuggee":true,"supportsBreakpointLocationsRequest":true,"supportsCancelRequest":true,"supportsCompletionsRequest":true,"supportsConditionalBreakpoints":true,"supportsConfigurationDoneRequest":true,"supportsDataBreakpoints":true,"supportsDelayedStackTraceLoading":true,"supportsDisassembleRequest":true,"supportsEvaluateForHovers":true,"supportsExceptionInfoRequest":true,"supportsExceptionOptions":true,"supportsFunctionBreakpoints":true,"supportsHitConditionalBreakpoints":true,"supportsInstructionBreakpoints":true,"supportsLogPoints":true,"supportsModulesRequest":true,"supportsReadMemoryRequest":true,"supportsRestartRequest":true,"supportsSetVariable":true,"supportsStepInTargetsRequest":true,"supportsSteppingGranularity":true,"supportsValueFormattingOptions":true},"command":"initialize","request_seq":1,"seq":0,"success":true,"type":"response"}
1745651006.073540688 --> (stdin/stdout) {"command":"attach","type":"request","arguments":{"program":"/tmp/lit-tmp-lz9mwtzv/tmpv73_0lcd","initCommands":["settings clear -all","settings set symbols.enable-external-lookup false","settings set target.inherit-tcc true","settings set target.disable-aslr false","settings set target.detach-on-error false","settings set target.auto-apply-fixits false","settings set plugin.process.gdb-remote.packet-timeout 60","settings set symbols.clang-modules-cache-path \"/home/tcwg-buildbot/worker/lldb-arm-ubuntu/build/lldb-test-build.noindex/module-cache-lldb/lldb-api\"","settings set use-color false","settings set show-statusline false"]},"seq":2}
1745651006.074194908 <-- (stdin/stdout) {"body":{"category":"console","output":"Running initCommands:\n"},"event":"output","seq":0,"type":"event"}
1745651006.074261189 <-- (stdin/stdout) {"body":{"category":"console","output":"(lldb) settings clear -all\n"},"event":"output","seq":0,"type":"event"}
1745651006.074277401 <-- (stdin/stdout) {"body":{"category":"console","output":"(lldb) settings set symbols.enable-external-lookup false\n"},"event":"output","seq":0,"type":"event"}
1745651006.074291468 <-- (stdin/stdout) {"body":{"category":"console","output":"(lldb) settings set target.inherit-tcc true\n"},"event":"output","seq":0,"type":"event"}
1745651006.074304342 <-- (stdin/stdout) {"body":{"category":"console","output":"(lldb) settings set target.disable-aslr false\n"},"event":"output","seq":0,"type":"event"}
1745651006.074316502 <-- (stdin/stdout) {"body":{"category":"console","output":"(lldb) settings set target.detach-on-error false\n"},"event":"output","seq":0,"type":"event"}
1745651006.074328423 <-- (stdin/stdout) {"body":{"category":"console","output":"(lldb) settings set target.auto-apply-fixits false\n"},"event":"output","seq":0,"type":"event"}
1745651006.074341059 <-- (stdin/stdout) {"body":{"category":"console","output":"(lldb) settings set plugin.process.gdb-remote.packet-timeout 60\n"},"event":"output","seq":0,"type":"event"}
1745651006.074354887 <-- (stdin/stdout) {"body":{"category":"console","output":"(lldb) settings set symbols.clang-modules-cache-path \"/home/tcwg-buildbot/worker/lldb-arm-ubuntu/build/lldb-test-build.noindex/module-cache-lldb/lldb-api\"\n"},"event":"output","seq":0,"type":"event"}
1745651006.074389696 <-- (stdin/stdout) {"body":{"category":"console","output":"(lldb) settings set use-color false\n"},"event":"output","seq":0,"type":"event"}
1745651006.074403286 <-- (stdin/stdout) {"body":{"category":"console","output":"(lldb) settings set show-statusline false\n"},"event":"output","seq":0,"type":"event"}
1745651006.251921415 <-- (stdin/stdout) {"command":"attach","request_seq":2,"seq":0,"success":true,"type":"response"}
1745651006.252081156 <-- (stdin/stdout) {"body":{"isLocalProcess":true,"name":"/tmp/lit-tmp-lz9mwtzv/tmpv73_0lcd","startMethod":"attach","systemProcessId":2701527},"event":"process","seq":0,"type":"event"}
1745651006.252115726 <-- (stdin/stdout) {"event":"initialized","seq":0,"type":"event"}
1745651006.253348351 --> (stdin/stdout) {"command":"setBreakpoints","type":"request","arguments":{"source":{"name":"main.c","path":"main.c"},"sourceModified":false,"lines":[28],"breakpoints":[{"line":28}]},"seq":3}
1745651006.265577793 <-- (stdin/stdout) {"body":{"breakpoints":[{"column":3,"id":1,"instructionReference":"0x4607D8","line":28,"source":{"name":"main.c","path":"main.c"},"verified":true}]},"command":"setBreakpoints","request_seq":3,"seq":0,"success":true,"type":"response"}
1745651006.266095400 <-- (stdin/stdout) {"body":{"breakpoint":{"column":3,"id":1,"instructionReference":"0x4607D8","line":28,"verified":true},"reason":"changed"},"event":"breakpoint","seq":0,"type":"event"}

@anutosh491
Copy link
Member Author

/cherry-pick 21fb19f

@llvmbot
Copy link
Member

llvmbot commented Apr 28, 2025

Failed to cherry-pick: 21fb19f

https://github.com/llvm/llvm-project/actions/runs/14705725537

Please manually backport the fix and push it to your github fork. Once this is done, please create a pull request

@anutosh491
Copy link
Member Author

/cherry-pick 21fb19f

@llvmbot
Copy link
Member

llvmbot commented Apr 28, 2025

/pull-request #137615

@llvmbot llvmbot moved this from Needs Triage to Done in LLVM Release Status Apr 28, 2025
swift-ci pushed a commit to swiftlang/llvm-project that referenced this pull request Apr 29, 2025
`clang-repl --cuda` was previously crashing with a segmentation fault,
instead of reporting a clean error
```
(base) anutosh491@Anutoshs-MacBook-Air bin % ./clang-repl --cuda
#0 0x0000000111da4fbc llvm::sys::PrintStackTrace(llvm::raw_ostream&, int) (/opt/local/libexec/llvm-20/lib/libLLVM.dylib+0x150fbc)
#1 0x0000000111da31dc llvm::sys::RunSignalHandlers() (/opt/local/libexec/llvm-20/lib/libLLVM.dylib+0x14f1dc)
#2 0x0000000111da5628 SignalHandler(int) (/opt/local/libexec/llvm-20/lib/libLLVM.dylib+0x151628)
#3 0x000000019b242de4 (/usr/lib/system/libsystem_platform.dylib+0x180482de4)
#4 0x0000000107f638d0 clang::IncrementalCUDADeviceParser::IncrementalCUDADeviceParser(std::__1::unique_ptr<clang::CompilerInstance, std::__1::default_delete<clang::CompilerInstance>>, clang::CompilerInstance&, llvm::IntrusiveRefCntPtr<llvm::vfs::InMemoryFileSystem>, llvm::Error&, std::__1::list<clang::PartialTranslationUnit, std::__1::allocator<clang::PartialTranslationUnit>> const&) (/opt/local/libexec/llvm-20/lib/libclang-cpp.dylib+0x216b8d0)
#5 0x0000000107f638d0 clang::IncrementalCUDADeviceParser::IncrementalCUDADeviceParser(std::__1::unique_ptr<clang::CompilerInstance, std::__1::default_delete<clang::CompilerInstance>>, clang::CompilerInstance&, llvm::IntrusiveRefCntPtr<llvm::vfs::InMemoryFileSystem>, llvm::Error&, std::__1::list<clang::PartialTranslationUnit, std::__1::allocator<clang::PartialTranslationUnit>> const&) (/opt/local/libexec/llvm-20/lib/libclang-cpp.dylib+0x216b8d0)
#6 0x0000000107f6bac8 clang::Interpreter::createWithCUDA(std::__1::unique_ptr<clang::CompilerInstance, std::__1::default_delete<clang::CompilerInstance>>, std::__1::unique_ptr<clang::CompilerInstance, std::__1::default_delete<clang::CompilerInstance>>) (/opt/local/libexec/llvm-20/lib/libclang-cpp.dylib+0x2173ac8)
#7 0x000000010206f8a8 main (/opt/local/libexec/llvm-20/bin/clang-repl+0x1000038a8)
#8 0x000000019ae8c274
Segmentation fault: 11
```

The underlying issue was that the `DeviceCompilerInstance` (used for
device-side CUDA compilation) was never initialized with a `Sema`, which
is required before constructing the `IncrementalCUDADeviceParser`.

https://github.com/llvm/llvm-project/blob/89687e6f383b742a3c6542dc673a84d9f82d02de/clang/lib/Interpreter/DeviceOffload.cpp#L32

https://github.com/llvm/llvm-project/blob/89687e6f383b742a3c6542dc673a84d9f82d02de/clang/lib/Interpreter/IncrementalParser.cpp#L31

Unlike the host-side `CompilerInstance` which runs `ExecuteAction`
inside the Interpreter constructor (thereby setting up Sema), the
device-side CI was passed into the parser uninitialized, leading to an
assertion or crash when accessing its internals.

To fix this, I refactored the `Interpreter::create` method to include an
optional `DeviceCI` parameter. If provided, we know we need to take care
of this instance too. Only then do we construct the
`IncrementalCUDADeviceParser`.

(cherry picked from commit 21fb19f)
IanWood1 pushed a commit to IanWood1/llvm-project that referenced this pull request May 6, 2025
`clang-repl --cuda` was previously crashing with a segmentation fault,
instead of reporting a clean error
```
(base) anutosh491@Anutoshs-MacBook-Air bin % ./clang-repl --cuda
#0 0x0000000111da4fbc llvm::sys::PrintStackTrace(llvm::raw_ostream&, int) (/opt/local/libexec/llvm-20/lib/libLLVM.dylib+0x150fbc)
llvm#1 0x0000000111da31dc llvm::sys::RunSignalHandlers() (/opt/local/libexec/llvm-20/lib/libLLVM.dylib+0x14f1dc)
llvm#2 0x0000000111da5628 SignalHandler(int) (/opt/local/libexec/llvm-20/lib/libLLVM.dylib+0x151628)
llvm#3 0x000000019b242de4 (/usr/lib/system/libsystem_platform.dylib+0x180482de4)
llvm#4 0x0000000107f638d0 clang::IncrementalCUDADeviceParser::IncrementalCUDADeviceParser(std::__1::unique_ptr<clang::CompilerInstance, std::__1::default_delete<clang::CompilerInstance>>, clang::CompilerInstance&, llvm::IntrusiveRefCntPtr<llvm::vfs::InMemoryFileSystem>, llvm::Error&, std::__1::list<clang::PartialTranslationUnit, std::__1::allocator<clang::PartialTranslationUnit>> const&) (/opt/local/libexec/llvm-20/lib/libclang-cpp.dylib+0x216b8d0)
llvm#5 0x0000000107f638d0 clang::IncrementalCUDADeviceParser::IncrementalCUDADeviceParser(std::__1::unique_ptr<clang::CompilerInstance, std::__1::default_delete<clang::CompilerInstance>>, clang::CompilerInstance&, llvm::IntrusiveRefCntPtr<llvm::vfs::InMemoryFileSystem>, llvm::Error&, std::__1::list<clang::PartialTranslationUnit, std::__1::allocator<clang::PartialTranslationUnit>> const&) (/opt/local/libexec/llvm-20/lib/libclang-cpp.dylib+0x216b8d0)
llvm#6 0x0000000107f6bac8 clang::Interpreter::createWithCUDA(std::__1::unique_ptr<clang::CompilerInstance, std::__1::default_delete<clang::CompilerInstance>>, std::__1::unique_ptr<clang::CompilerInstance, std::__1::default_delete<clang::CompilerInstance>>) (/opt/local/libexec/llvm-20/lib/libclang-cpp.dylib+0x2173ac8)
llvm#7 0x000000010206f8a8 main (/opt/local/libexec/llvm-20/bin/clang-repl+0x1000038a8)
llvm#8 0x000000019ae8c274 
Segmentation fault: 11
```


The underlying issue was that the `DeviceCompilerInstance` (used for
device-side CUDA compilation) was never initialized with a `Sema`, which
is required before constructing the `IncrementalCUDADeviceParser`.


https://github.com/llvm/llvm-project/blob/89687e6f383b742a3c6542dc673a84d9f82d02de/clang/lib/Interpreter/DeviceOffload.cpp#L32


https://github.com/llvm/llvm-project/blob/89687e6f383b742a3c6542dc673a84d9f82d02de/clang/lib/Interpreter/IncrementalParser.cpp#L31

Unlike the host-side `CompilerInstance` which runs `ExecuteAction`
inside the Interpreter constructor (thereby setting up Sema), the
device-side CI was passed into the parser uninitialized, leading to an
assertion or crash when accessing its internals.

To fix this, I refactored the `Interpreter::create` method to include an
optional `DeviceCI` parameter. If provided, we know we need to take care
of this instance too. Only then do we construct the
`IncrementalCUDADeviceParser`.
anutosh491 added a commit that referenced this pull request Aug 26, 2025
…alParser (#137458)

Read discussion : #136404 (comment)
and the following comments for context

Motivation
1) `IncrementalAction` is designed to keep Frontend statealive across
inputs. As per the docstring: “IncrementalAction ensures it keeps its
underlying action's objects alive as long as the IncrementalParser needs
them.”
2) To align responsibilities with that contract, the parser layer (host:
`IncrementalParser`, device: `IncrementalCUDADeviceParser`) should
manage PTU registration and module generation, while the interpreter
orchestrates at a higher level.

What this PR does
1) Moves CodeGen surfaces behind IncrementalAction:
GenModule(), getCodeGen(), and the cached “first CodeGen module” now
live in IncrementalAction.

2) Moves PTU ownership to the parser layer:
Adds IncrementalParser::RegisterPTU(…) (and device counterpart)

3) Add device-side registration in IncrementalCUDADeviceParser.

4) Remove Interpreter::{getCodeGen, GenModule, RegisterPTU}.
llvm-sync bot pushed a commit to arm/arm-toolchain that referenced this pull request Aug 26, 2025
…o IncrementalParser (#137458)

Read discussion : llvm/llvm-project#136404 (comment)
and the following comments for context

Motivation
1) `IncrementalAction` is designed to keep Frontend statealive across
inputs. As per the docstring: “IncrementalAction ensures it keeps its
underlying action's objects alive as long as the IncrementalParser needs
them.”
2) To align responsibilities with that contract, the parser layer (host:
`IncrementalParser`, device: `IncrementalCUDADeviceParser`) should
manage PTU registration and module generation, while the interpreter
orchestrates at a higher level.

What this PR does
1) Moves CodeGen surfaces behind IncrementalAction:
GenModule(), getCodeGen(), and the cached “first CodeGen module” now
live in IncrementalAction.

2) Moves PTU ownership to the parser layer:
Adds IncrementalParser::RegisterPTU(…) (and device counterpart)

3) Add device-side registration in IncrementalCUDADeviceParser.

4) Remove Interpreter::{getCodeGen, GenModule, RegisterPTU}.
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
clang:frontend Language frontend issues, e.g. anything involving "Sema" clang Clang issues not falling into any other category
Projects
Development

Successfully merging this pull request may close these issues.

4 participants