https://github.com/koparasy updated https://github.com/llvm/llvm-project/pull/225971
>From bf096fab36f57915d1c7fd1109c1b4a3fab1d097 Mon Sep 17 00:00:00 2001 From: Konstantinos Parasyris <[email protected]> Date: Wed, 23 Sep 2026 16:00:22 -0700 Subject: [PATCH] [CIR][CUDA] Read fat binary in CIRGen and store its bytes on the module Co-Authored-By: Claude Opus 5 <[email protected]> --- .../clang/CIR/Dialect/IR/CIRCUDAAttrs.td | 18 ------ .../clang/CIR/Dialect/IR/CIRDialect.td | 4 +- clang/lib/CIR/CodeGen/CIRGenCUDANV.cpp | 56 ++++++++++++++++++- clang/lib/CIR/CodeGen/CIRGenModule.cpp | 11 ---- .../Dialect/Transforms/LoweringPrepare.cpp | 42 +++++--------- clang/test/CIR/CodeGenCUDA/device-stub.cu | 6 ++ .../test/CIR/CodeGenCUDA/missing-gpubinary.cu | 22 ++++++++ .../CIR/Diagnostics/mlir-error-routing.cpp | 21 +++++++ .../CIR/Diagnostics/mlir-error-routing.cu | 20 ------- 9 files changed, 122 insertions(+), 78 deletions(-) create mode 100644 clang/test/CIR/CodeGenCUDA/missing-gpubinary.cu create mode 100644 clang/test/CIR/Diagnostics/mlir-error-routing.cpp delete mode 100644 clang/test/CIR/Diagnostics/mlir-error-routing.cu diff --git a/clang/include/clang/CIR/Dialect/IR/CIRCUDAAttrs.td b/clang/include/clang/CIR/Dialect/IR/CIRCUDAAttrs.td index d993e1b2b11eb..ebe8eaf60cc43 100644 --- a/clang/include/clang/CIR/Dialect/IR/CIRCUDAAttrs.td +++ b/clang/include/clang/CIR/Dialect/IR/CIRCUDAAttrs.td @@ -50,24 +50,6 @@ def CIR_CUDAExternallyInitializedAttr : CIR_Attr<"CUDAExternallyInitialized", }]; let canHaveIllegalCXXABIType = 0; } -def CIR_CUDABinaryHandleAttr : CIR_Attr< - "CUDABinaryHandle", "cu.binary_handle" -> { - let summary = "Fat binary handle for device code."; - let description = - [{ - This attribute is attached to the ModuleOp and records the binary file - name passed to host. - - CUDA first compiles device-side code into a fat binary file. The file - name is then passed into host-side code, which is used to create a handle - and then generate various registration functions. - }]; - - let parameters = (ins "mlir::StringAttr":$name); - let assemblyFormat = "`<` $name `>`"; -} - // No wrapper attribute: the kind is only ever printed by // CIR_CUDAVarRegistrationInfoAttr's own assembly format. def CIR_CUDADeviceVarKind : CIR_I32Enum<"CUDADeviceVarKind", diff --git a/clang/include/clang/CIR/Dialect/IR/CIRDialect.td b/clang/include/clang/CIR/Dialect/IR/CIRDialect.td index 6f88d3a2b547f..2964d86e508e0 100644 --- a/clang/include/clang/CIR/Dialect/IR/CIRDialect.td +++ b/clang/include/clang/CIR/Dialect/IR/CIRDialect.td @@ -93,7 +93,9 @@ def CIR_Dialect : Dialect { static llvm::StringRef getTargetCPUAttrName() { return "cir.target-cpu"; } static llvm::StringRef getTuneCPUAttrName() { return "cir.tune-cpu"; } static llvm::StringRef getTargetFeaturesAttrName() { return "cir.target-features"; } - static llvm::StringRef getCUDABinaryHandleAttrName() { return "cir.cu.binary_handle"; } + // Raw bytes of the device-side fat binary, read by CIRGen so LoweringPrepare + // can build the runtime-registration globals without doing file I/O. + static llvm::StringRef getCUDADeviceBinaryAttrName() { return "cir.cu.device_binary"; } // Mangled symbol name of the C++20 named-module initializer function, // precomputed by CIRGen so later passes don't need a live ASTContext. static llvm::StringRef getCXXModuleInitFnNameAttrName() { return "cir.cxx_module_init_fn_name"; } diff --git a/clang/lib/CIR/CodeGen/CIRGenCUDANV.cpp b/clang/lib/CIR/CodeGen/CIRGenCUDANV.cpp index ab4baf336d379..fc95ab48eab94 100644 --- a/clang/lib/CIR/CodeGen/CIRGenCUDANV.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenCUDANV.cpp @@ -22,9 +22,14 @@ #include "clang/AST/GlobalDecl.h" #include "clang/Basic/AddressSpaces.h" #include "clang/Basic/Cuda.h" +#include "clang/Basic/DiagnosticFrontend.h" +#include "clang/Basic/FileManager.h" +#include "clang/Basic/SourceManager.h" #include "clang/CIR/Dialect/IR/CIRDialect.h" #include "clang/CIR/Dialect/IR/CIRTypes.h" #include "llvm/Support/Casting.h" +#include "llvm/Support/MemoryBuffer.h" +#include "llvm/Support/VirtualFileSystem.h" using namespace clang; using namespace clang::CIRGen; @@ -56,6 +61,7 @@ class CIRGenNVCUDARuntime : public CIRGenCUDARuntime { private: void emitDeviceStubBodyNew(CIRGenFunction &cgf, cir::FuncOp fn, FunctionArgList &args); + void recordDeviceBinary(); mlir::Value prepareKernelArgs(CIRGenFunction &cgf, mlir::Location loc, FunctionArgList &args); mlir::Operation *getKernelHandle(cir::FuncOp fn, GlobalDecl gd) override; @@ -517,9 +523,57 @@ void CIRGenNVCUDARuntime::handleGlobalReplace(cir::GlobalOp oldGV, } } +/// Whether this translation unit has anything for the CUDA runtime to register. +/// These are the same two attributes LoweringPrepare collects to decide whether +/// to build a module ctor, so both sides answer the question from one source. +static bool hasEntitiesToRegister(mlir::ModuleOp module) { + return llvm::any_of( + module.getBody()->getOperations(), [](mlir::Operation &op) { + return op.hasAttr(cir::CUDAKernelNameAttr::getMnemonic()) || + op.hasAttr(cir::CUDAVarRegistrationInfoAttr::getMnemonic()); + }); +} + +/// Read the device-side fat binary and record its contents on the module as +/// `cir.cu.device_binary`, for LoweringPrepare to build the fatbin global from. +/// +/// This mirrors the read in CGNVCUDARuntime::makeModuleCtorFunction, guards +/// included: nothing is read in a compilation that would build no module +/// constructor. +void CIRGenNVCUDARuntime::recordDeviceBinary() { + StringRef binaryName = cgm.getCodeGenOpts().OffloadBinaryToEmbedFile; + if (binaryName.empty()) + return; + + const LangOptions &langOpts = cgm.getLangOpts(); + if ((langOpts.HIP || !langOpts.GPURelocatableDeviceCode) && + !hasEntitiesToRegister(cgm.getModule())) + return; + + llvm::vfs::FileSystem &fs = cgm.getASTContext() + .getSourceManager() + .getFileManager() + .getVirtualFileSystem(); + llvm::ErrorOr<std::unique_ptr<llvm::MemoryBuffer>> binaryOrErr = + fs.getBufferForFile(binaryName, /*FileSize=*/-1, + /*RequiresNullTerminator=*/false); + if (std::error_code ec = binaryOrErr.getError()) { + cgm.getDiags().Report(diag::err_cannot_open_file) + << binaryName << ec.message(); + return; + } + + cgm.getModule()->setAttr( + cir::CIRDialect::getCUDADeviceBinaryAttrName(), + mlir::StringAttr::get(&cgm.getMLIRContext(), + binaryOrErr.get()->getBuffer())); +} + void CIRGenNVCUDARuntime::finalizeModule() { - if (!cgm.getLangOpts().CUDAIsDevice) + if (!cgm.getLangOpts().CUDAIsDevice) { + recordDeviceBinary(); return; + } // Mark ODR-used device variables as compiler used to prevent them from being // eliminated by optimization. This is necessary for device variables diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.cpp b/clang/lib/CIR/CodeGen/CIRGenModule.cpp index 61d5b03f393b0..dc6a365bf790b 100644 --- a/clang/lib/CIR/CodeGen/CIRGenModule.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenModule.cpp @@ -212,17 +212,6 @@ CIRGenModule::CIRGenModule(mlir::MLIRContext &mlirContext, /*line=*/0, /*column=*/0)); } - - // Set CUDA GPU binary handle. - if (langOpts.CUDA) { - llvm::StringRef cudaBinaryName = codeGenOpts.OffloadBinaryToEmbedFile; - if (!cudaBinaryName.empty()) { - theModule->setAttr(cir::CIRDialect::getCUDABinaryHandleAttrName(), - cir::CUDABinaryHandleAttr::get( - &mlirContext, mlir::StringAttr::get( - &mlirContext, cudaBinaryName))); - } - } } CIRGenModule::~CIRGenModule() = default; diff --git a/clang/lib/CIR/Dialect/Transforms/LoweringPrepare.cpp b/clang/lib/CIR/Dialect/Transforms/LoweringPrepare.cpp index 2abe5a82c49e6..20742ef2302f0 100644 --- a/clang/lib/CIR/Dialect/Transforms/LoweringPrepare.cpp +++ b/clang/lib/CIR/Dialect/Transforms/LoweringPrepare.cpp @@ -32,9 +32,7 @@ #include "llvm/ADT/TypeSwitch.h" #include "llvm/IR/Instructions.h" #include "llvm/Support/ErrorHandling.h" -#include "llvm/Support/MemoryBuffer.h" #include "llvm/Support/Path.h" -#include "llvm/Support/VirtualFileSystem.h" #include <map> #include <memory> @@ -2514,30 +2512,15 @@ void LoweringPreparePass::buildCUDAModuleCtor() { // There's no device-side binary, so no need to proceed for CUDA. // HIP has to create an external symbol in this case, which is NYI. - mlir::Attribute cudaBinaryHandleAttr = - mlirModule->getAttr(CIRDialect::getCUDABinaryHandleAttrName()); - if (!cudaBinaryHandleAttr) { + auto deviceBinaryAttr = mlirModule->getAttrOfType<mlir::StringAttr>( + CIRDialect::getCUDADeviceBinaryAttrName()); + if (!deviceBinaryAttr) { if (isHIP) assert(!cir::MissingFeatures::hipModuleCtor()); return; } - llvm::StringRef cudaGPUBinaryName = - mlir::cast<CUDABinaryHandleAttr>(cudaBinaryHandleAttr) - .getName() - .getValue(); - - llvm::vfs::FileSystem &vfs = - astCtx->getSourceManager().getFileManager().getVirtualFileSystem(); - llvm::ErrorOr<std::unique_ptr<llvm::MemoryBuffer>> gpuBinaryOrErr = - vfs.getBufferForFile(cudaGPUBinaryName); - if (std::error_code ec = gpuBinaryOrErr.getError()) { - mlirModule->emitError("cannot open GPU binary file: " + cudaGPUBinaryName + - ": " + ec.message()); - return; - } - std::unique_ptr<llvm::MemoryBuffer> gpuBinary = - std::move(gpuBinaryOrErr.get()); + llvm::StringRef deviceBinary = deviceBinaryAttr.getValue(); // Set up common types and builder. llvm::StringRef cudaPrefix = getCUDAPrefix(getLangOpts()); @@ -2563,8 +2546,7 @@ void LoweringPreparePass::buildCUDAModuleCtor() { getLangOpts().HIP ? ".hipFatBinSegment" : ".nvFatBinSegment"; // Create the fatbin string constant with GPU binary contents. - auto fatbinType = - ArrayType::get(&getContext(), charTy, gpuBinary->getBuffer().size()); + auto fatbinType = ArrayType::get(&getContext(), charTy, deviceBinary.size()); std::string fatbinStrName = addUnderscoredPrefix(cudaPrefix, "_fatbin_str"); GlobalOp fatbinStr = GlobalOp::create(builder, loc, fatbinStrName, fatbinType, /*isConstant=*/true, {}, @@ -2577,7 +2559,7 @@ void LoweringPreparePass::buildCUDAModuleCtor() { } fatbinStr.setInitialValueAttr(cir::ConstArrayAttr::get( - fatbinType, StringAttr::get(gpuBinary->getBuffer(), fatbinType))); + fatbinType, StringAttr::get(deviceBinary, fatbinType))); fatbinStr.setSection(fatbinConstName); fatbinStr.setPrivate(); @@ -2755,7 +2737,7 @@ void LoweringPreparePass::buildCUDAModuleCtor() { } std::optional<FuncOp> LoweringPreparePass::buildCUDAModuleDtor() { - if (!mlirModule->getAttr(CIRDialect::getCUDABinaryHandleAttrName())) + if (!mlirModule->getAttr(CIRDialect::getCUDADeviceBinaryAttrName())) return {}; llvm::StringRef prefix = getCUDAPrefix(getLangOpts()); @@ -2812,7 +2794,7 @@ std::optional<FuncOp> LoweringPreparePass::buildCUDAModuleDtor() { /// the dtor list would cause a double-free. It is meant to be registered via /// atexit() at the end of the module ctor. std::optional<FuncOp> LoweringPreparePass::buildHIPModuleDtor() { - if (!mlirModule->getAttr(CIRDialect::getCUDABinaryHandleAttrName())) + if (!mlirModule->getAttr(CIRDialect::getCUDADeviceBinaryAttrName())) return {}; llvm::StringRef prefix = getCUDAPrefix(getLangOpts()); @@ -3092,8 +3074,14 @@ void LoweringPreparePass::runOnOperation() { buildCXXGlobalInitFunc(); buildCXXGlobalTlsFunc(); - if (getLangOpts().CUDA && !getLangOpts().CUDAIsDevice) + if (getLangOpts().CUDA && !getLangOpts().CUDAIsDevice) { buildCUDAModuleCtor(); + // The bytes are in the fatbin global now; drop the attribute so a large fat + // binary isn't stored twice in an emitted .cir. This has to happen out here + // because the ctor and both dtor builders test the attribute to decide + // whether a device-side binary exists at all. + mlirModule->removeAttr(CIRDialect::getCUDADeviceBinaryAttrName()); + } buildGlobalCtorDtorList(); } diff --git a/clang/test/CIR/CodeGenCUDA/device-stub.cu b/clang/test/CIR/CodeGenCUDA/device-stub.cu index 1de4e5575ce1d..bc175500a45b2 100644 --- a/clang/test/CIR/CodeGenCUDA/device-stub.cu +++ b/clang/test/CIR/CodeGenCUDA/device-stub.cu @@ -123,6 +123,10 @@ __device__ _BitInt(36) c; // CIR: cir.global "private" constant cir_private @__cuda_fatbin_str = #cir.const_array<"GPU binary would be here." : !cir.array<!u8i x 25>> : !cir.array<!u8i x 25> {alignment = 8 : i64, section = ".nv_fatbin"} +// The bytes arrive as the #cir.cu.device_binary module attribute, which +// LoweringPrepare erases once they are in the global above. +// CIR-NOT: cir.cu.device_binary + // Check the fatbin wrapper struct: { magic, version, ptr to fatbin, null }, with section. // CIR: cir.global constant cir_private @__cuda_fatbin_wrapper = #cir.const_record<{ // CIR-SAME: #cir.int<1180844977> : !s32i, @@ -196,6 +200,7 @@ __device__ _BitInt(36) c; // LLVM: call i32 @atexit(ptr @__cuda_module_dtor) // No GPU binary — no registration infrastructure at all. +// NOGPUBIN-NOT: cir.cu.device_binary // NOGPUBIN-NOT: fatbin // NOGPUBIN-NOT: gpubin // NOGPUBIN-NOT: __cuda_register_globals @@ -337,6 +342,7 @@ __device__ _BitInt(36) c; // HIP-LLVM: ret void // No GPU binary: no fatbin, no handle, no registration scaffolding. +// HIP-NOGPUBIN-NOT: cir.cu.device_binary // HIP-NOGPUBIN-NOT: __hip_fatbin // HIP-NOGPUBIN-NOT: __hip_gpubin_handle // HIP-NOGPUBIN-NOT: __hip_register_globals diff --git a/clang/test/CIR/CodeGenCUDA/missing-gpubinary.cu b/clang/test/CIR/CodeGenCUDA/missing-gpubinary.cu new file mode 100644 index 0000000000000..964ec4a40bffd --- /dev/null +++ b/clang/test/CIR/CodeGenCUDA/missing-gpubinary.cu @@ -0,0 +1,22 @@ +// A missing fat binary must produce the same clang diagnostic from ClangIR as +// from classic codegen, not an MLIR pass error. + +// RUN: not %clang_cc1 -triple x86_64-linux-gnu -emit-cir %s -x cuda \ +// RUN: -target-sdk-version=12.3 -fcuda-include-gpubinary %t.nonexistent \ +// RUN: -o %t.cir 2>&1 | FileCheck %s + +// RUN: not %clang_cc1 -triple x86_64-linux-gnu -fclangir -emit-llvm %s -x cuda \ +// RUN: -target-sdk-version=12.3 -fcuda-include-gpubinary %t.nonexistent \ +// RUN: -o %t.ll 2>&1 | FileCheck %s + +// RUN: not %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s -x cuda \ +// RUN: -target-sdk-version=12.3 -fcuda-include-gpubinary %t.nonexistent \ +// RUN: -o %t-ogcg.ll 2>&1 | FileCheck %s + +#include "Inputs/cuda.h" + +// A kernel is needed: with nothing to register neither CIRGen nor classic +// codegen reads the fat binary at all. +__global__ void kernel() {} + +// CHECK: fatal error: cannot open file '{{.*}}.nonexistent': diff --git a/clang/test/CIR/Diagnostics/mlir-error-routing.cpp b/clang/test/CIR/Diagnostics/mlir-error-routing.cpp new file mode 100644 index 0000000000000..5ce9da67f1f25 --- /dev/null +++ b/clang/test/CIR/Diagnostics/mlir-error-routing.cpp @@ -0,0 +1,21 @@ +// RUN: not %clang_cc1 -triple x86_64-apple-macosx10.15 -fclangir -emit-llvm \ +// RUN: %s -o %t.ll 2>&1 | FileCheck %s + +// LoweringPrepare emits an MLIR-side error via mlir::Operation::emitError for +// a thread_local variable on a target whose thread wrapper is replaceable. +// CIRDiagnosticHandler must surface it in clang's `file:line:col: error: ...` +// format rather than MLIR's `loc("file":N:M): error: ...` default. + +// CHECK: mlir-error-routing.cpp:[[#@LINE+7]]:1: error: Unhandled thread wrapper attributes for CC and Nounwind +// CHECK-NOT: loc({{.*}}): error: Unhandled thread wrapper attributes + +struct S { + S(); + ~S(); +}; +thread_local S s; +S *use() { return &s; } + +// The generic CIR-to-CIR transform fatal error must not be reported on top of +// the specific one: CIRGenAction gates it on hasErrorOccurred(). +// CHECK-NOT: error: CIR-to-CIR transformation failed diff --git a/clang/test/CIR/Diagnostics/mlir-error-routing.cu b/clang/test/CIR/Diagnostics/mlir-error-routing.cu deleted file mode 100644 index 25759ace8b5bd..0000000000000 --- a/clang/test/CIR/Diagnostics/mlir-error-routing.cu +++ /dev/null @@ -1,20 +0,0 @@ -// RUN: not %clang_cc1 -triple x86_64-linux-gnu -fclangir -emit-llvm -x cuda \ -// RUN: -target-sdk-version=12.3 -fcuda-include-gpubinary %t.missing.bin \ -// RUN: %s -o %t.ll 2>&1 | FileCheck %s - -// LoweringPrepare emits an MLIR-side error via mlir::Operation::emitError when -// the requested CUDA gpubinary cannot be opened. With CIRDiagnosticHandler -// installed, that diagnostic surfaces through clang's DiagnosticsEngine in -// clang's standard format (`error: ...`) rather than MLIR's -// `loc("file":N:M): error: ...` default-handler format. - -// CHECK: error: cannot open GPU binary file: {{.*}}.missing.bin -// CHECK-NOT: loc({{.*}}): error: cannot open GPU binary file - -// The generic CIR-to-CIR transform fatal error must NOT be reported on top of -// the specific MLIR-relayed error. CIRGenAction gates the fallback diag on -// clang::DiagnosticsEngine::hasErrorOccurred() so users see one root cause, -// not two. -// CHECK-NOT: error: CIR-to-CIR transformation failed - -__attribute__((global)) void kernel() {} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
