https://github.com/aaronj0 updated https://github.com/llvm/llvm-project/pull/226975
>From 4b15d1272c6b2addea97bcba1c426efc097047a9 Mon Sep 17 00:00:00 2001 From: Aaron Jomy <[email protected]> Date: Thu, 24 Sep 2026 16:00:08 +0200 Subject: [PATCH 1/3] [clang-repl] Flush the CUDA device bootstrap module before the first PTU The host path sets the bootstrap module aside with CacheCodeGenModule() once the initial action has run, so the first PTU starts from a fresh module. The CUDA device path skipped this step. The device module that HandleTranslationUnit had already finalized during bootstrap stayed current, and the first device PTU finalized it a second time, resulting in CodeGen adding every module flag twice. The IR verifier then rejects the module. To reproduce, on a build with assertions, or with `-fverify-intermediate-code` (hidden on release builds since the driver disables the verifier there), the first input to `clang-repl --cuda` fails: ``` module flag identifiers must be unique (or of 'require' type) !"nvvm-reflect-ftz" module flag identifiers must be unique (or of 'require' type) !"PIC Level" module flag identifiers must be unique (or of 'require' type) !"frame-pointer" fatal error: error in backend: Broken module found, compilation aborted! ``` Tested on 22.x and main. The `host-supports-cuda` lit probe also hits this: clang/test/Interpreter/CUDA reports UNSUPPORTED instead of failing. After, the first device module carries each flag once and the verifier accepts it. The probe passes and the CUDA tests run on an assertions build. --- clang/lib/Interpreter/Interpreter.cpp | 3 +++ 1 file changed, 3 insertions(+) diff --git a/clang/lib/Interpreter/Interpreter.cpp b/clang/lib/Interpreter/Interpreter.cpp index 655db32477a0c..5d5629997e94f 100644 --- a/clang/lib/Interpreter/Interpreter.cpp +++ b/clang/lib/Interpreter/Interpreter.cpp @@ -514,6 +514,9 @@ Interpreter::createWithDevice(OffloadType Type, if (llvm::Error E = ExecuteIncrementalAction(*DCI, *Interp->DeviceAct)) return std::move(E); + // Set the finalized initial device module aside, as the host path does. + Interp->DeviceAct->CacheCodeGenModule(); + Interp->DeviceCI = std::move(DCI); if (Type == OffloadType::HIP) { >From c7fae79cfc403643a7ba01f9e643022e52158378 Mon Sep 17 00:00:00 2001 From: Aaron Jomy <[email protected]> Date: Tue, 29 Sep 2026 18:38:08 +0200 Subject: [PATCH 2/3] [clang-repl] Add a verifier lit test for the CUDA bootstrap module The test turns the IR verifier on, which release builds leave off, and checks that the first device module passes it. --- .../Interpreter/CUDA/device-module-verifier.cu | 14 ++++++++++++++ 1 file changed, 14 insertions(+) create mode 100644 clang/test/Interpreter/CUDA/device-module-verifier.cu diff --git a/clang/test/Interpreter/CUDA/device-module-verifier.cu b/clang/test/Interpreter/CUDA/device-module-verifier.cu new file mode 100644 index 0000000000000..5e9c6c1862cce --- /dev/null +++ b/clang/test/Interpreter/CUDA/device-module-verifier.cu @@ -0,0 +1,14 @@ +// Release builds skip the IR verifier. Turn it on and check that the first +// device module passes it. +// RUN: cat %s | clang-repl --cuda -Xcc -fverify-intermediate-code 2>&1 \ +// RUN: | FileCheck %s + +extern "C" int printf(const char*, ...); + +__global__ void kernel() {} +printf("kernel: %d\n", 0); +// CHECK-NOT: module flag identifiers must be unique +// CHECK-NOT: Broken module found +// CHECK: kernel: 0 + +%quit >From 3239be8d12cea9fa83767ab35577e986d90d5ffd Mon Sep 17 00:00:00 2001 From: Aaron Jomy <[email protected]> Date: Wed, 30 Sep 2026 10:04:07 +0200 Subject: [PATCH 3/3] [clang-repl] Test CUDA interpreter creation with the IR verifier enabled Add a unittest that builds a CUDA interpreter with -fverify-intermediate-code and parses one host-only input. A regression leads to a broken module (report_fatal_error), so run the interpreter in a child process as a death test. The test only requires the NVPTX backend and does not need a CUDA toolkit or a GPU present. Adds DeviceOffloadTest.cpp for the device side of the interpreter's unit tests. --- clang/unittests/Interpreter/CMakeLists.txt | 1 + .../Interpreter/DeviceOffloadTest.cpp | 76 +++++++++++++++++++ 2 files changed, 77 insertions(+) create mode 100644 clang/unittests/Interpreter/DeviceOffloadTest.cpp diff --git a/clang/unittests/Interpreter/CMakeLists.txt b/clang/unittests/Interpreter/CMakeLists.txt index 66b396b53cb55..6d999104054c3 100644 --- a/clang/unittests/Interpreter/CMakeLists.txt +++ b/clang/unittests/Interpreter/CMakeLists.txt @@ -35,6 +35,7 @@ set(CLANG_REPL_TEST_SOURCES InterpreterTest.cpp InterpreterExtensionsTest.cpp CodeCompletionTest.cpp + DeviceOffloadTest.cpp ) if(TARGET compiler-rt AND LLVM_ON_UNIX) diff --git a/clang/unittests/Interpreter/DeviceOffloadTest.cpp b/clang/unittests/Interpreter/DeviceOffloadTest.cpp new file mode 100644 index 0000000000000..9fa1e0b3e9e60 --- /dev/null +++ b/clang/unittests/Interpreter/DeviceOffloadTest.cpp @@ -0,0 +1,76 @@ +//===- unittests/Interpreter/DeviceOffloadTest.cpp - CUDA device tests ----===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// +// Unit tests for the Incremental CUDA compilation infrastructure under Clang's +// Interpreter library. The tests in this file only require the NVPTX backend +// and run without a CUDA toolkit/GPU. +// +//===----------------------------------------------------------------------===// + +#include "InterpreterTestFixture.h" + +#include "clang/Frontend/CompilerInstance.h" +#include "clang/Interpreter/Interpreter.h" + +#include "llvm/MC/TargetRegistry.h" +#include "llvm/Support/Error.h" +#include "llvm/Support/TargetSelect.h" +#include "llvm/TargetParser/Triple.h" + +#include "gtest/gtest.h" + +#include <cstdlib> + +using namespace clang; + +namespace { + +class DeviceOffloadTest : public InterpreterTestBase { +protected: + static void SetUpTestSuite() { + InterpreterTestBase::SetUpTestSuite(); + llvm::InitializeAllTargets(); + llvm::InitializeAllTargetMCs(); + llvm::InitializeAllAsmPrinters(); + } + + void SetUp() override { + InterpreterTestBase::SetUp(); + if (IsSkipped()) + return; + std::string Err; + if (!llvm::TargetRegistry::lookupTarget(llvm::Triple("nvptx64-nvidia-cuda"), + Err)) + GTEST_SKIP() << Err; + } +}; + +TEST_F(DeviceOffloadTest, FirstDeviceModuleVerifies) { +#if GTEST_HAS_DEATH_TEST + // Release builds leave the IR verifier off. A broken module ends in + // report_fatal_error, so the first device PTU is built in a child process. + EXPECT_EXIT( + { + // Without the runtime headers and libdevice no CUDA toolkit is needed. + IncrementalCompilerBuilder CB; + CB.SetCompilerArgs( + {"-nocudainc", "-nocudalib", "-fverify-intermediate-code"}); + auto DeviceCI = llvm::cantFail(CB.CreateDevice(OffloadType::CUDA)); + auto HostCI = llvm::cantFail(CB.CreateHost(OffloadType::CUDA)); + auto Interp = llvm::cantFail(Interpreter::createWithDevice( + OffloadType::CUDA, std::move(HostCI), std::move(DeviceCI))); + llvm::cantFail(Interp->Parse("int i = 0;")); + exit(0); + }, + ::testing::ExitedWithCode(0), ""); +#else + GTEST_SKIP() << "no death tests on this platform"; +#endif +} + +} // end anonymous namespace _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
