[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
Fri Sep 25 08:08:06 PDT 2026
https://github.com/koparasy updated https://github.com/llvm/llvm-project/pull/225971
>From e7b55736ded9bc935e9ecb9976e579da2fade021 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 1/2] [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 | 61 ++++++++++++++++++-
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, 127 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 b2a2707a67b66..72e9471d78707 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..b4e5840f25100 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,62 @@ 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) {
+ // A walk, not a scan of the module body, so this stays in agreement with the
+ // recursive walk LoweringPrepare collects them with.
+ return module
+ ->walk([](mlir::Operation *op) {
+ if (op->hasAttr(cir::CUDAKernelNameAttr::getMnemonic()) ||
+ op->hasAttr(cir::CUDAVarRegistrationInfoAttr::getMnemonic()))
+ return mlir::WalkResult::interrupt();
+ return mlir::WalkResult::advance();
+ })
+ .wasInterrupted();
+}
+
+/// 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 adffa7dfe2969..007788d7e27c7 100644
--- a/clang/lib/CIR/CodeGen/CIRGenModule.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenModule.cpp
@@ -216,17 +216,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 dff7646c640b0..eb61b26bf3ab5 100644
--- a/clang/lib/CIR/Dialect/Transforms/LoweringPrepare.cpp
+++ b/clang/lib/CIR/Dialect/Transforms/LoweringPrepare.cpp
@@ -32,10 +32,8 @@
#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/VersionTuple.h"
-#include "llvm/Support/VirtualFileSystem.h"
#include <map>
#include <memory>
@@ -2538,30 +2536,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());
@@ -2587,8 +2570,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, {},
@@ -2601,7 +2583,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();
@@ -2781,7 +2763,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());
@@ -2838,7 +2820,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());
@@ -3118,8 +3100,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 768a013dc8319..a480cee5d8ebd 100644
--- a/clang/test/CIR/CodeGenCUDA/device-stub.cu
+++ b/clang/test/CIR/CodeGenCUDA/device-stub.cu
@@ -130,6 +130,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,
@@ -213,6 +217,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
@@ -354,6 +359,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() {}
>From 2bab06c8785d297ada31904ba2b8aaa5e83e70be Mon Sep 17 00:00:00 2001
From: Konstantinos Parasyris <konstantinos.parasyris at intel.com>
Date: Thu, 24 Sep 2026 16:41:04 -0700
Subject: [PATCH 2/2] Address feedback
Co-Authored-By: Claude Opus 5.5 <noreply at anthropic.com>
---
clang/lib/CIR/CodeGen/CIRGenCUDANV.cpp | 15 +++++---
clang/lib/CIR/Dialect/IR/CIRDialect.cpp | 15 ++++++++
.../Dialect/Transforms/LoweringPrepare.cpp | 20 +++++------
.../CIR/IR/invalid-cuda-device-binary.cir | 34 +++++++++++++++++++
4 files changed, 68 insertions(+), 16 deletions(-)
create mode 100644 clang/test/CIR/IR/invalid-cuda-device-binary.cir
diff --git a/clang/lib/CIR/CodeGen/CIRGenCUDANV.cpp b/clang/lib/CIR/CodeGen/CIRGenCUDANV.cpp
index b4e5840f25100..79153ca788756 100644
--- a/clang/lib/CIR/CodeGen/CIRGenCUDANV.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenCUDANV.cpp
@@ -568,10 +568,17 @@ void CIRGenNVCUDARuntime::recordDeviceBinary() {
return;
}
- cgm.getModule()->setAttr(
- cir::CIRDialect::getCUDADeviceBinaryAttrName(),
- mlir::StringAttr::get(&cgm.getMLIRContext(),
- binaryOrErr.get()->getBuffer()));
+ // Typed as the fatbin global's array type so LoweringPrepare can use this
+ // attribute as the initializer as-is: attributes are uniqued on {value, type}
+ // and never freed, so building a second, typed copy there would keep the fat
+ // binary in memory twice.
+ StringRef bytes = binaryOrErr.get()->getBuffer();
+ mlir::MLIRContext &ctx = cgm.getMLIRContext();
+ auto charTy = cir::IntType::get(&ctx, cgm.getTarget().getCharWidth(),
+ /*isSigned=*/false);
+ auto fatbinTy = cir::ArrayType::get(charTy, bytes.size());
+ cgm.getModule()->setAttr(cir::CIRDialect::getCUDADeviceBinaryAttrName(),
+ mlir::StringAttr::get(bytes, fatbinTy));
}
void CIRGenNVCUDARuntime::finalizeModule() {
diff --git a/clang/lib/CIR/Dialect/IR/CIRDialect.cpp b/clang/lib/CIR/Dialect/IR/CIRDialect.cpp
index ce406707f5942..d5a587ff6d81e 100644
--- a/clang/lib/CIR/Dialect/IR/CIRDialect.cpp
+++ b/clang/lib/CIR/Dialect/IR/CIRDialect.cpp
@@ -270,6 +270,21 @@ cir::CIRDialect::verifyOperationAttribute(mlir::Operation *op,
<< "' attribute to be attached to '"
<< mlir::ModuleOp::getOperationName() << "'";
+ // LoweringPrepare uses this attribute directly as the fatbin global's
+ // initializer, so it must be a valid #cir.const_array payload for its type.
+ if (attrName == getCUDADeviceBinaryAttrName()) {
+ auto bytes = mlir::dyn_cast<mlir::StringAttr>(attr.getValue());
+ auto arrayTy =
+ bytes ? mlir::dyn_cast<cir::ArrayType>(bytes.getType()) : nullptr;
+ if (!arrayTy || arrayTy.getSize() != bytes.size())
+ return op->emitOpError()
+ << "expects '" << getCUDADeviceBinaryAttrName()
+ << "' to be a string typed as an array of its length";
+ return cir::ConstArrayAttr::verify([&] { return op->emitOpError(); },
+ arrayTy, bytes,
+ /*trailingZerosNum=*/0);
+ }
+
return success();
}
diff --git a/clang/lib/CIR/Dialect/Transforms/LoweringPrepare.cpp b/clang/lib/CIR/Dialect/Transforms/LoweringPrepare.cpp
index eb61b26bf3ab5..9a3c9c9745eaa 100644
--- a/clang/lib/CIR/Dialect/Transforms/LoweringPrepare.cpp
+++ b/clang/lib/CIR/Dialect/Transforms/LoweringPrepare.cpp
@@ -2544,8 +2544,6 @@ void LoweringPreparePass::buildCUDAModuleCtor() {
return;
}
- llvm::StringRef deviceBinary = deviceBinaryAttr.getValue();
-
// Set up common types and builder.
llvm::StringRef cudaPrefix = getCUDAPrefix(getLangOpts());
mlir::Location loc = mlirModule->getLoc();
@@ -2556,9 +2554,6 @@ void LoweringPreparePass::buildCUDAModuleCtor() {
PointerType voidPtrTy = builder.getVoidPtrTy();
PointerType voidPtrPtrTy = builder.getPointerTo(voidPtrTy);
IntType intTy = builder.getSIntNTy(32);
- IntType charTy =
- cir::IntType::get(&getContext(), getTargetInfo().getCharWidth(),
- /*isSigned=*/false);
// --- Create fatbin globals ---
@@ -2570,7 +2565,8 @@ void LoweringPreparePass::buildCUDAModuleCtor() {
getLangOpts().HIP ? ".hipFatBinSegment" : ".nvFatBinSegment";
// Create the fatbin string constant with GPU binary contents.
- auto fatbinType = ArrayType::get(&getContext(), charTy, deviceBinary.size());
+ // The dialect verifier guarantees the attribute is typed as the array.
+ auto fatbinType = mlir::cast<ArrayType>(deviceBinaryAttr.getType());
std::string fatbinStrName = addUnderscoredPrefix(cudaPrefix, "_fatbin_str");
GlobalOp fatbinStr = GlobalOp::create(builder, loc, fatbinStrName, fatbinType,
/*isConstant=*/true, {},
@@ -2582,8 +2578,8 @@ void LoweringPreparePass::buildCUDAModuleCtor() {
fatbinStr.setAlignment(8);
}
- fatbinStr.setInitialValueAttr(cir::ConstArrayAttr::get(
- fatbinType, StringAttr::get(deviceBinary, fatbinType)));
+ fatbinStr.setInitialValueAttr(
+ cir::ConstArrayAttr::get(fatbinType, deviceBinaryAttr));
fatbinStr.setSection(fatbinConstName);
fatbinStr.setPrivate();
@@ -3102,10 +3098,10 @@ void LoweringPreparePass::runOnOperation() {
buildCXXGlobalTlsFunc();
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.
+ // The fatbin global now references the same attribute; drop the module's
+ // reference so an emitted .cir doesn't print the bytes twice. 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());
}
diff --git a/clang/test/CIR/IR/invalid-cuda-device-binary.cir b/clang/test/CIR/IR/invalid-cuda-device-binary.cir
new file mode 100644
index 0000000000000..ef4d2c6b22374
--- /dev/null
+++ b/clang/test/CIR/IR/invalid-cuda-device-binary.cir
@@ -0,0 +1,34 @@
+// RUN: cir-opt %s -verify-diagnostics -split-input-file
+
+!u8i = !cir.int<u, 8>
+
+module attributes {cir.cu.device_binary = "abc" : !cir.array<!u8i x 3>} {
+}
+
+// -----
+
+// expected-error at +1 {{expects 'cir.cu.device_binary' to be a string typed as an array of its length}}
+module attributes {cir.cu.device_binary = "abc"} {
+}
+
+// -----
+
+!u8i = !cir.int<u, 8>
+
+// expected-error at +1 {{expects 'cir.cu.device_binary' to be a string typed as an array of its length}}
+module attributes {cir.cu.device_binary = "abc" : !cir.array<!u8i x 4>} {
+}
+
+// -----
+
+!u16i = !cir.int<u, 16>
+
+// expected-error at +1 {{expects !cir.int<u, 8> element type}}
+module attributes {cir.cu.device_binary = "abc" : !cir.array<!u16i x 3>} {
+}
+
+// -----
+
+// expected-error at +1 {{expects 'cir.cu.device_binary' to be a string typed as an array of its length}}
+module attributes {cir.cu.device_binary = 42 : i32} {
+}
More information about the cfe-commits
mailing list