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 <[email protected]> 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 <[email protected]> --- .../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 <[email protected]> Date: Thu, 24 Sep 2026 16:41:04 -0700 Subject: [PATCH 2/2] Address feedback Co-Authored-By: Claude Opus 5.5 <[email protected]> --- 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@+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@+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@+1 {{expects !cir.int<u, 8> element type}} +module attributes {cir.cu.device_binary = "abc" : !cir.array<!u16i x 3>} { +} + +// ----- + +// expected-error@+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} { +} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
