[clang] [CIR][CUDA] Read fat binary in CIRGen and store its bytes on the module (PR #225971)
Konstantinos Parasyris via cfe-commits
cfe-commits at lists.llvm.org
Wed Sep 23 16:10:50 PDT 2026
https://github.com/koparasy created https://github.com/llvm/llvm-project/pull/225971
Two problems with the current implementation of GPU CIR:
1. `LoweringPrepare` read the fat binary from disk, via
`astCtx->getSourceManager().getFileManager().getVirtualFileSystem()`. A
transform pass should not do I/O, and this is one of the `ASTContext`
dependencies that keeps the post-CIRGen pipeline from being IR-to-IR.
Some of the related discussion on this has been done [here](https://discourse.llvm.org/t/rfc-clangir-making-cir-pipeline-boundaries-first-class-driver-artifacts/90998) and PR: #219048
2. `#cir.cu.binary_handle` stored the *file path*. That's a build input, not a
property of the program: it bakes one machine's directory layout into a
`.cir`.
The PR removes the file path and adds an attribute
CIRGen now reads the file in and records the contents
as `#cir.cu.device_binary`, a `StringAttr` of raw bytes. The
attribute is deleted once LoweringPrepare "consumes" it.
A couple of notes:
1. The GPU dialect has a similar representation. We may want to move
towards that direction slowly, but I don't see any value doing this now.
It will confuse the direction of this PR.
2. The test `clang/test/CIR/Diagnostics/mlir-error-routing.cu` has completely
changed, because it relied on an error occurring when the file attribute was missing.
Since we deleted that attribute completely, the test was failing.
>From 8756aca66e42312d005e3cdd10773b8adb5d317c Mon Sep 17 00:00:00 2001
From: Konstantinos Parasyris <konstantinos.parasyris at intel.com>
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 <noreply at anthropic.com>
---
.../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() {}
More information about the cfe-commits
mailing list