https://github.com/vgvassilev updated https://github.com/llvm/llvm-project/pull/226975
>From 3208e175c7f87e905ac34036784bd317f2de53f7 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 cef7ef6aa65d9289289b4c6d3043bc460f43c481 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 37822db7da134491bbb656bac6259e19c4977500 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 device function definition. 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. The input keeps the device module non-empty, so the test does not depend on the separate fix for PTX emission from an empty device module. 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..964ee6c1dc572 --- /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("__attribute__((device)) void f() {}")); + 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
