https://github.com/aaronj0 updated https://github.com/llvm/llvm-project/pull/235754
>From 617300508cf1f8ab5949c6ad1da76fa68ad64d0d Mon Sep 17 00:00:00 2001 From: Aaron Jomy <[email protected]> Date: Mon, 26 Oct 2026 12:00:15 +0100 Subject: [PATCH 1/3] [clang-repl] Pass the target features to the device TargetMachine IncrementalCUDADeviceParser::GeneratePTX creates the NVPTX TargetMachine with an empty feature string. The +ptxNN feature that the driver derives from the CUDA version is dropped, and the backend emits the minimum PTX version for the target SM instead. A kernel using a newer instruction then fails to load with cudaErrorInvalidPtx, for example shfl.sync below sm_70. Join TargetOptions::Features into the feature string, as EmitAssemblyHelper::CreateTargetMachine does. Add a unittest that passes --cuda-feature=+ptx80 to the device compilation and checks the .version of the embedded PTX. --- clang/lib/Interpreter/DeviceOffload.cpp | 5 ++-- .../Interpreter/DeviceOffloadTest.cpp | 25 +++++++++++++++++++ 2 files changed, 28 insertions(+), 2 deletions(-) diff --git a/clang/lib/Interpreter/DeviceOffload.cpp b/clang/lib/Interpreter/DeviceOffload.cpp index 94fec23b1c0b2f..d7ae386466f40f 100644 --- a/clang/lib/Interpreter/DeviceOffload.cpp +++ b/clang/lib/Interpreter/DeviceOffload.cpp @@ -19,6 +19,7 @@ #include "clang/Frontend/CompilerInstance.h" #include "clang/Interpreter/PartialTranslationUnit.h" +#include "llvm/ADT/StringExtras.h" #include "llvm/IR/LegacyPassManager.h" #include "llvm/IR/Module.h" #include "llvm/MC/TargetRegistry.h" @@ -122,8 +123,8 @@ llvm::Expected<llvm::StringRef> IncrementalCUDADeviceParser::GeneratePTX() { std::error_code()); llvm::TargetOptions TO = llvm::TargetOptions(); std::unique_ptr<llvm::TargetMachine> TM(Target->createTargetMachine( - PTU.TheModule->getTargetTriple(), TargetOpts.CPU, "", TO, - llvm::Reloc::Model::PIC_)); + PTU.TheModule->getTargetTriple(), TargetOpts.CPU, + llvm::join(TargetOpts.Features, ","), TO, llvm::Reloc::Model::PIC_)); PTXCode.clear(); llvm::raw_svector_ostream dest(PTXCode); diff --git a/clang/unittests/Interpreter/DeviceOffloadTest.cpp b/clang/unittests/Interpreter/DeviceOffloadTest.cpp index fa624f23f435b7..e89f3c5e6c617a 100644 --- a/clang/unittests/Interpreter/DeviceOffloadTest.cpp +++ b/clang/unittests/Interpreter/DeviceOffloadTest.cpp @@ -130,6 +130,31 @@ TEST_F(DeviceOffloadTest, EmptyDeviceModule) { EXPECT_TRUE((*Buf)->getBuffer().contains(".target")); } +TEST_F(DeviceOffloadTest, PTXVersionFromTargetFeature) { + // The +ptxNN target feature decides the PTX ISA version. + IncrementalCompilerBuilder CB; + CB.SetCompilerArgs( + {"-nocudainc", "-nocudalib", "-Xarch_device", "--cuda-feature=+ptx80"}); + auto DeviceCI = CB.CreateDevice(OffloadType::CUDA); + ASSERT_THAT_EXPECTED(DeviceCI, llvm::Succeeded()); + auto HostCI = CB.CreateHost(OffloadType::CUDA); + ASSERT_THAT_EXPECTED(HostCI, llvm::Succeeded()); + auto Interp = Interpreter::createWithDevice( + std::move(*HostCI), std::move(*DeviceCI), OffloadType::CUDA); + ASSERT_THAT_EXPECTED(Interp, llvm::Succeeded()); + + auto PTU = (*Interp)->Parse("int i = 0;"); + ASSERT_THAT_EXPECTED(PTU, llvm::Succeeded()); + + const CompilerInstance *CI = (*Interp)->getCompilerInstance(); + auto Buf = CI->getVirtualFileSystem().getBufferForFile( + CI->getCodeGenOpts().OffloadBinaryToEmbedFile, /*FileSize=*/-1, + /*RequiresNullTerminator=*/false); + ASSERT_TRUE(static_cast<bool>(Buf)) << Buf.getError().message(); + llvm::StringRef PTX = (*Buf)->getBuffer(); + EXPECT_TRUE(PTX.contains(".version 8.0")) << PTX; +} + TEST(DeviceOffloadBuilderTest, MultipleOffloadArchsAreRejected) { // The driver creates one compile job per architecture. Only one can be used, // so this must be reported instead of silently using the first one. >From 43cead68feb2fb8f72daf1f12e0df8d6b00da291 Mon Sep 17 00:00:00 2001 From: Aaron Jomy <[email protected]> Date: Fri, 30 Oct 2026 21:51:52 +0100 Subject: [PATCH 2/3] [clang-repl] Check the PTX version of the device fatbin Add a lit test that pins --cuda-feature=+ptx80 on the device side and checks for .version 8.0 in the fatbin embedded in the -emit-llvm output. It fails without the target features passed to the device TargetMachine, with .version 3.2. --- clang/test/Interpreter/CUDA/ptx-version.cu | 14 ++++++++++++++ 1 file changed, 14 insertions(+) create mode 100644 clang/test/Interpreter/CUDA/ptx-version.cu diff --git a/clang/test/Interpreter/CUDA/ptx-version.cu b/clang/test/Interpreter/CUDA/ptx-version.cu new file mode 100644 index 00000000000000..536fd37de458ea --- /dev/null +++ b/clang/test/Interpreter/CUDA/ptx-version.cu @@ -0,0 +1,14 @@ +// The PTX version of the embedded fatbin must follow the pinned feature. +// RUN: rm -rf %t +// RUN: mkdir -p %t/lib +// RUN: %clang --target=%host-jit-triple -nostdlib -fPIC -shared \ +// RUN: %S/Inputs/cudart.c -o %t/lib/libcudart.so +// RUN: cat %s | clang-repl --cuda --cuda-path=%t -Xcc -nocudainc \ +// RUN: -Xcc -nocudalib -Xcc -include -Xcc %S/Inputs/cuda.h \ +// RUN: -Xcc -Xarch_device -Xcc --cuda-feature=+ptx80 \ +// RUN: -Xcc -Xclang -Xcc -emit-llvm | FileCheck %s + +__global__ void kernel() {} +// CHECK: private constant {{.*}}.version 8.0{{.*}}.target sm_ + +%quit >From 8b136610b767eecd0faab871982c2bd0d8133177 Mon Sep 17 00:00:00 2001 From: Aaron Jomy <[email protected]> Date: Sat, 31 Oct 2026 07:50:29 +0100 Subject: [PATCH 3/3] address review comments: drop unittest --- .../Interpreter/DeviceOffloadTest.cpp | 25 ------------------- 1 file changed, 25 deletions(-) diff --git a/clang/unittests/Interpreter/DeviceOffloadTest.cpp b/clang/unittests/Interpreter/DeviceOffloadTest.cpp index e89f3c5e6c617a..fa624f23f435b7 100644 --- a/clang/unittests/Interpreter/DeviceOffloadTest.cpp +++ b/clang/unittests/Interpreter/DeviceOffloadTest.cpp @@ -130,31 +130,6 @@ TEST_F(DeviceOffloadTest, EmptyDeviceModule) { EXPECT_TRUE((*Buf)->getBuffer().contains(".target")); } -TEST_F(DeviceOffloadTest, PTXVersionFromTargetFeature) { - // The +ptxNN target feature decides the PTX ISA version. - IncrementalCompilerBuilder CB; - CB.SetCompilerArgs( - {"-nocudainc", "-nocudalib", "-Xarch_device", "--cuda-feature=+ptx80"}); - auto DeviceCI = CB.CreateDevice(OffloadType::CUDA); - ASSERT_THAT_EXPECTED(DeviceCI, llvm::Succeeded()); - auto HostCI = CB.CreateHost(OffloadType::CUDA); - ASSERT_THAT_EXPECTED(HostCI, llvm::Succeeded()); - auto Interp = Interpreter::createWithDevice( - std::move(*HostCI), std::move(*DeviceCI), OffloadType::CUDA); - ASSERT_THAT_EXPECTED(Interp, llvm::Succeeded()); - - auto PTU = (*Interp)->Parse("int i = 0;"); - ASSERT_THAT_EXPECTED(PTU, llvm::Succeeded()); - - const CompilerInstance *CI = (*Interp)->getCompilerInstance(); - auto Buf = CI->getVirtualFileSystem().getBufferForFile( - CI->getCodeGenOpts().OffloadBinaryToEmbedFile, /*FileSize=*/-1, - /*RequiresNullTerminator=*/false); - ASSERT_TRUE(static_cast<bool>(Buf)) << Buf.getError().message(); - llvm::StringRef PTX = (*Buf)->getBuffer(); - EXPECT_TRUE(PTX.contains(".version 8.0")) << PTX; -} - TEST(DeviceOffloadBuilderTest, MultipleOffloadArchsAreRejected) { // The driver creates one compile job per architecture. Only one can be used, // so this must be reported instead of silently using the first one. _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
