Author: Aaron Jomy Date: 2026-10-31T07:37:10Z New Revision: 2323019b73220fd87aaf28175e830dca67fa4cdd
URL: https://github.com/llvm/llvm-project/commit/2323019b73220fd87aaf28175e830dca67fa4cdd DIFF: https://github.com/llvm/llvm-project/commit/2323019b73220fd87aaf28175e830dca67fa4cdd.diff LOG: [clang-repl] Pass the target features to the device TargetMachine (#235754) IncrementalCUDADeviceParser creates the NVPTX TargetMachine with an empty feature string. The +ptxNN feature that the driver derives from the CUDA version is dropped, causing the backend to emit 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 done in EmitAssemblyHelper::CreateTargetMachine. Add a unittest that passes --cuda-feature=+ptx80 to the device compilation and checks the .version of the embedded PTX. Added: clang/test/Interpreter/CUDA/ptx-version.cu Modified: clang/lib/Interpreter/DeviceOffload.cpp Removed: ################################################################################ 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/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 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
