[clang] [clang-repl] Fix PTX emission for CUDA inputs with no device code (PR #226977)

[via cfe-commits]

llvmorg-github-actions[bot] wrote:


<!--LLVM PR SUMMARY COMMENT-->

@llvm/pr-subscribers-clang

Author: Aaron Jomy (aaronj0)

<details>
<summary>Changes</summary>

Previously, GeneratePTX() returned an error when PassManager::run() returned false. That value reports whether any pass changed the module, not whether emission succeeded. IIUC, addPassesToEmitFile() should be the only failure point, which is already checked.

The check was harmless until e8b75c172810 ("[NVPTX] Add NewPM boilerplate to NVPTXAssignValidGlobalNames"). Before it, that pass returned true unconditionally, so run() always reported a change. Now a device module without a function definition, produced in the case of host-only inputs, reports no change and fails the check, falsely erroring out with `Failed to emit PTX code.` Due to this, every host-only input to `clang-repl --cuda` is now rejected on main. The host-supports-cuda lit probe consists mostly of such inputs, so clang/test/Interpreter/CUDA reports UNSUPPORTED instead of failing.

I've also added a test with host-only inputs so a regression like this could be caught.

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


2 Files Affected:

- (modified) clang/lib/Interpreter/DeviceOffload.cpp (+1-3) 
- (added) clang/test/Interpreter/CUDA/empty-device-module.cu (+18) 


``````````diff
diff --git a/clang/lib/Interpreter/DeviceOffload.cpp b/clang/lib/Interpreter/DeviceOffload.cpp
index bf7653c518c30..d876e4cf10c3d 100644
--- a/clang/lib/Interpreter/DeviceOffload.cpp
+++ b/clang/lib/Interpreter/DeviceOffload.cpp
@@ -68,9 +68,7 @@ llvm::Expected<llvm::StringRef> IncrementalCUDADeviceParser::GeneratePTX() {
         llvm::inconvertibleErrorCode());
   }
 
-  if (!PM.run(*PTU.TheModule))
-    return llvm::make_error<llvm::StringError>("Failed to emit PTX code.",
-                                               llvm::inconvertibleErrorCode());
+  PM.run(*PTU.TheModule);
 
   PTXCode += '\0';
   while (PTXCode.size() % 8)
diff --git a/clang/test/Interpreter/CUDA/empty-device-module.cu b/clang/test/Interpreter/CUDA/empty-device-module.cu
new file mode 100644
index 0000000000000..fd9ba760e5a69
--- /dev/null
+++ b/clang/test/Interpreter/CUDA/empty-device-module.cu
@@ -0,0 +1,18 @@
+// Tests host-only inputs. They produce an empty device module, and emitting
+// PTX for it must not be reported as a failure just because no pass changed
+// the module.
+// RUN: cat %s | clang-repl --cuda | FileCheck %s
+
+extern "C" int printf(const char*, ...);
+
+int host_only = 42;
+printf("host_only: %d\n", host_only);
+// CHECK: host_only: 42
+
+__global__ void kernel() {}
+
+kernel<<<1,1>>>();
+printf("CUDA Error: %d\n", cudaGetLastError());
+// CHECK-NEXT: CUDA Error: 0
+
+%quit

``````````

</details>


https://github.com/llvm/llvm-project/pull/226977