[flang-commits] [flang] d925f39 - [flang][cuda] register type information pages after all units register (#231619)
via flang-commits
flang-commits at lists.llvm.org
Tue Oct 13 20:59:45 PDT 2026
Author: Valentin Clement (バレンタイン クレメン)
Date: 2026-10-13T20:59:40-07:00
New Revision: d925f39805d70c32c7ca3643326d3fd629d8f3c1
URL: https://github.com/llvm/llvm-project/commit/d925f39805d70c32c7ca3643326d3fd629d8f3c1
DIFF: https://github.com/llvm/llvm-project/commit/d925f39805d70c32c7ca3643326d3fd629d8f3c1.diff
LOG: [flang][cuda] register type information pages after all units register (#231619)
Registering the pages of the shared type information calls the CUDA
runtime, which initializes the module. With relocatable device code, the
managed variables registered by later units were then never populated.
Register the pages in the initialization constructor, which runs after
the registration constructors of all the units.
Added:
flang/test/Fir/CUDA/cuda-shared-type-info-no-managed.mlir
flang/test/Fir/CUDA/cuda-shared-type-info-used.mlir
Modified:
flang-rt/lib/cuda/registration.cpp
flang/include/flang/Optimizer/Builder/CUFCommon.h
flang/include/flang/Optimizer/Transforms/Passes.td
flang/include/flang/Runtime/CUDA/registration.h
flang/lib/Optimizer/Transforms/CUDA/CUFAddConstructor.cpp
flang/lib/Optimizer/Transforms/CUDA/CUFSharedTypeInfo.cpp
flang/test/Fir/CUDA/cuda-shared-type-info.mlir
Removed:
################################################################################
diff --git a/flang-rt/lib/cuda/registration.cpp b/flang-rt/lib/cuda/registration.cpp
index ec93a061cf716..4477c04b74767 100644
--- a/flang-rt/lib/cuda/registration.cpp
+++ b/flang-rt/lib/cuda/registration.cpp
@@ -60,10 +60,32 @@ void RTDEF(CUFInitModule)(void **module) { __cudaInitModule(module); }
void RTDEF(CUFRegisterHostMemoryRange)(void *begin, void *end) {
if (!begin || end <= begin)
return;
+ // With HMM or ATS, the device reads pageable host memory directly.
+ int device{0};
+ int pageableAccess{0};
+ if (cudaGetDevice(&device) != cudaSuccess ||
+ cudaDeviceGetAttribute(&pageableAccess, cudaDevAttrPageableMemoryAccess,
+ device) != cudaSuccess)
+ (void)cudaGetLastError();
+ else if (pageableAccess)
+ return;
+ // Registering a page that holds other data would break copies of that data
+ // that extend past the page. The compiler aligns the start of the range and
+ // the global ending it, at end - 1, to a page boundary, so that only the
+ // pages from the start up to that global need to be registered. An edge
+ // that is not aligned, because the pages are larger than that alignment or
+ // the objects predate it, can share its page with other data. That page is
+ // registered anyway: the device faults on type information that is not
+ // registered, whereas only copies extending past the page fail otherwise.
const auto pageSize{static_cast<std::uintptr_t>(sysconf(_SC_PAGESIZE))};
- const auto first{reinterpret_cast<std::uintptr_t>(begin) & ~(pageSize - 1)};
- const auto last{
- (reinterpret_cast<std::uintptr_t>(end) + pageSize - 1) & ~(pageSize - 1)};
+ const auto beginAddr{reinterpret_cast<std::uintptr_t>(begin)};
+ const auto endAddr{reinterpret_cast<std::uintptr_t>(end)};
+ const std::uintptr_t first{beginAddr & ~(pageSize - 1)};
+ const std::uintptr_t last{(endAddr - 1) % pageSize == 0
+ ? endAddr - 1
+ : (endAddr + pageSize - 1) & ~(pageSize - 1)};
+ if (last <= first)
+ return;
cudaError_t err{cudaHostRegister(reinterpret_cast<void *>(first),
last - first, cudaHostRegisterPortable | cudaHostRegisterMapped)};
if (err == cudaErrorHostMemoryAlreadyRegistered) {
diff --git a/flang/include/flang/Optimizer/Builder/CUFCommon.h b/flang/include/flang/Optimizer/Builder/CUFCommon.h
index 6c2afc4ec0e4d..f364018aceca4 100644
--- a/flang/include/flang/Optimizer/Builder/CUFCommon.h
+++ b/flang/include/flang/Optimizer/Builder/CUFCommon.h
@@ -15,6 +15,7 @@
#include "mlir/IR/BuiltinOps.h"
#include "llvm/ADT/DenseMap.h"
#include "llvm/ADT/SetVector.h"
+#include "llvm/TargetParser/Triple.h"
static constexpr llvm::StringRef cudaDeviceModuleName = "cuda_device_mod";
static constexpr llvm::StringRef cudaSharedMemSuffix = "__shared_mem__";
@@ -26,6 +27,18 @@ static constexpr llvm::StringRef cudaDeviceCopySuffix = ".device";
/// under -gpu=mem:managed. It must be a valid C identifier so the linker
/// defines the __start_ and __stop_ symbols bounding it.
static constexpr llvm::StringRef cudaSharedTypeInfoSection = "__nv_type_info";
+/// Name prefix of the page-aligned global emitted last in that section by each
+/// translation unit, so that the section starts on a page boundary and the
+/// pages before the last one hold only type information.
+static constexpr llvm::StringRef cudaSharedTypeInfoEndPrefix =
+ "__nv_type_info_end_";
+/// Alignment of that global: the largest page size of \p triple, so that the
+/// layout holds whatever page size the system uses.
+inline int64_t getCudaSharedTypeInfoAlignment(const llvm::Triple &triple) {
+ if (triple.isAArch64() || triple.isPPC64())
+ return 65536;
+ return 4096;
+}
/// Attribute on the GPU module mapping the name of each shared type
/// descriptor to the managed pointer holding its host address.
static constexpr llvm::StringRef cudaSharedTypeDescsAttrName =
diff --git a/flang/include/flang/Optimizer/Transforms/Passes.td b/flang/include/flang/Optimizer/Transforms/Passes.td
index a161af2826079..edc98e2f1ec14 100644
--- a/flang/include/flang/Optimizer/Transforms/Passes.td
+++ b/flang/include/flang/Optimizer/Transforms/Passes.td
@@ -664,10 +664,15 @@ def CUFSharedTypeInfo : Pass<"cuf-shared-type-info", "mlir::ModuleOp"> {
descriptor address from them in device code. The `acc.declare` attribute
is dropped so OpenACC does not copy the type information to the device.
+ Each translation unit defining such type information also ends the
+ section with a page-aligned global, so that the registered pages hold
+ nothing but type information.
+
Device code then sees the host binding tables, so type-bound procedures
cannot be called from device code.
}];
- let dependentDialects = ["mlir::gpu::GPUDialect", "cuf::CUFDialect"];
+ let dependentDialects = ["mlir::gpu::GPUDialect", "cuf::CUFDialect",
+ "mlir::LLVM::LLVMDialect"];
}
def CUFAddConstructor : Pass<"cuf-add-constructor", "mlir::ModuleOp"> {
diff --git a/flang/include/flang/Runtime/CUDA/registration.h b/flang/include/flang/Runtime/CUDA/registration.h
index 3286a277a8a2b..6c19f58e36528 100644
--- a/flang/include/flang/Runtime/CUDA/registration.h
+++ b/flang/include/flang/Runtime/CUDA/registration.h
@@ -37,9 +37,11 @@ void RTDECL(CUFRegisterManagedVariable)(
/// unified memory addresses.
void RTDECL(CUFInitModule)(void **module);
-/// Register the pages spanning [\p begin, \p end) with the CUDA runtime so
-/// device code can access them through their host address. Registering the
-/// same range more than once is allowed.
+/// Register the pages holding [\p begin, \p end) with the CUDA runtime so
+/// device code can access them through their host address. When \p end - 1 is
+/// on a page boundary, its page is left out.
+/// Nothing is registered when the device can access pageable host memory.
+/// Registering the same range more than once is allowed.
void RTDECL(CUFRegisterHostMemoryRange)(void *begin, void *end);
} // extern "C"
diff --git a/flang/lib/Optimizer/Transforms/CUDA/CUFAddConstructor.cpp b/flang/lib/Optimizer/Transforms/CUDA/CUFAddConstructor.cpp
index b4c9d324f84f7..4cab45fb42de0 100644
--- a/flang/lib/Optimizer/Transforms/CUDA/CUFAddConstructor.cpp
+++ b/flang/lib/Optimizer/Transforms/CUDA/CUFAddConstructor.cpp
@@ -407,6 +407,19 @@ struct CUFAddConstructor
auto entryBlock = func.addEntryBlock(builder);
builder.setInsertionPointToStart(entryBlock);
mlir::LLVM::LLVMFuncOp initCtor;
+ // Returns the block at the end of the initialization constructor, which
+ // is terminated once all of its code is emitted.
+ auto getInitCtorBlock = [&]() -> mlir::Block * {
+ if (!initCtor) {
+ mlir::OpBuilder::InsertionGuard guard(builder);
+ builder.setInsertionPointToEnd(mod.getBody());
+ initCtor = mlir::LLVM::LLVMFuncOp::create(
+ builder, loc, cudaFortranInitCtorName, funcTy);
+ initCtor.setLinkage(mlir::LLVM::Linkage::Internal);
+ initCtor.addEntryBlock(builder);
+ }
+ return &initCtor.getBody().back();
+ };
if (needAllocatorRegistration) {
llvm::StringRef allocatorRegistrationFunctionName =
@@ -424,44 +437,6 @@ struct CUFAddConstructor
mlir::LLVM::CallOp::create(builder, loc, funcTy, cufRegisterAllocatorRef);
}
- // Device code reads the runtime type information defined here through its
- // host address, so its pages must be registered with the CUDA runtime.
- // The section bounds are local to the executable or shared library being
- // linked, and every unit defining type information registers the same
- // range.
- if (cudaManagedTypeInfo && definesSharedTypeInfo(mod)) {
- auto llvmPtrTy = mlir::LLVM::LLVMPointerType::get(ctx);
- auto getSectionBound = [&](llvm::StringRef prefix) -> mlir::Value {
- std::string name = (prefix + cudaSharedTypeInfoSection).str();
- auto bound = mod.lookupSymbol<mlir::LLVM::GlobalOp>(name);
- if (!bound) {
- mlir::OpBuilder::InsertionGuard guard(builder);
- builder.setInsertionPointToEnd(mod.getBody());
- bound = mlir::LLVM::GlobalOp::create(
- builder, loc, mlir::IntegerType::get(ctx, 8),
- /*isConstant=*/false, mlir::LLVM::Linkage::ExternWeak, name,
- mlir::Attribute{});
- bound.setVisibility_(mlir::LLVM::Visibility::Hidden);
- }
- return mlir::LLVM::AddressOfOp::create(builder, loc, bound);
- };
- mlir::Value begin = getSectionBound("__start_");
- mlir::Value end = getSectionBound("__stop_");
- auto rangeFuncTy = mlir::LLVM::LLVMFunctionType::get(
- voidTy, {llvmPtrTy, llvmPtrTy}, /*isVarArg=*/false);
- llvm::StringRef rangeFuncName = RTNAME_STRING(CUFRegisterHostMemoryRange);
- if (!mod.lookupSymbol(rangeFuncName)) {
- mlir::OpBuilder::InsertionGuard guard(builder);
- builder.setInsertionPointToEnd(mod.getBody());
- auto rangeFunc = mlir::LLVM::LLVMFuncOp::create(
- builder, loc, rangeFuncName, rangeFuncTy);
- rangeFunc.setVisibility(mlir::SymbolTable::Visibility::Private);
- }
- mlir::LLVM::CallOp::create(builder, loc, rangeFuncTy,
- mlir::SymbolRefAttr::get(ctx, rangeFuncName),
- mlir::ValueRange{begin, end});
- }
-
auto gpuMod = symTab.lookup<mlir::gpu::GPUModuleOp>(cudaDeviceModuleName);
if (gpuMod) {
mlir::SymbolTable gpuSymTable(gpuMod);
@@ -600,11 +575,7 @@ struct CUFAddConstructor
mlir::LLVM::ZeroOp::create(builder, loc, llvmPtrTy);
mlir::LLVM::ReturnOp::create(builder, loc, nullHandle);
- builder.setInsertionPointToEnd(mod.getBody());
- initCtor = mlir::LLVM::LLVMFuncOp::create(
- builder, loc, cudaFortranInitCtorName, funcTy);
- initCtor.setLinkage(mlir::LLVM::Linkage::Internal);
- mlir::Block *initBlock = initCtor.addEntryBlock(builder);
+ mlir::Block *initBlock = getInitCtorBlock();
builder.setInsertionPointToEnd(entryBlock);
mlir::LLVM::StoreOp::create(
@@ -662,10 +633,56 @@ struct CUFAddConstructor
mlir::LLVM::BrOp::create(builder, loc, nextBlock);
builder.setInsertionPointToStart(nextBlock);
}
- mlir::LLVM::ReturnOp::create(builder, loc, mlir::ValueRange{});
}
}
}
+
+ // Device code reads the runtime type information defined here through its
+ // host address, so its pages must be registered with the CUDA runtime.
+ // The section bounds are local to the executable or shared library being
+ // linked, and every unit defining type information registers the same
+ // range. The registration calls the CUDA runtime, which initializes the
+ // module and ignores the managed variables registered afterwards, so it
+ // runs in the initialization constructor.
+ if (cudaManagedTypeInfo && definesSharedTypeInfo(mod)) {
+ mlir::OpBuilder::InsertionGuard guard(builder);
+ builder.setInsertionPointToEnd(getInitCtorBlock());
+ auto llvmPtrTy = mlir::LLVM::LLVMPointerType::get(ctx);
+ auto getSectionBound = [&](llvm::StringRef prefix) -> mlir::Value {
+ std::string name = (prefix + cudaSharedTypeInfoSection).str();
+ auto bound = mod.lookupSymbol<mlir::LLVM::GlobalOp>(name);
+ if (!bound) {
+ mlir::OpBuilder::InsertionGuard guard(builder);
+ builder.setInsertionPointToEnd(mod.getBody());
+ bound = mlir::LLVM::GlobalOp::create(
+ builder, loc, mlir::IntegerType::get(ctx, 8),
+ /*isConstant=*/false, mlir::LLVM::Linkage::ExternWeak, name,
+ mlir::Attribute{});
+ bound.setVisibility_(mlir::LLVM::Visibility::Hidden);
+ }
+ return mlir::LLVM::AddressOfOp::create(builder, loc, bound);
+ };
+ mlir::Value begin = getSectionBound("__start_");
+ mlir::Value end = getSectionBound("__stop_");
+ auto rangeFuncTy = mlir::LLVM::LLVMFunctionType::get(
+ voidTy, {llvmPtrTy, llvmPtrTy}, /*isVarArg=*/false);
+ llvm::StringRef rangeFuncName = RTNAME_STRING(CUFRegisterHostMemoryRange);
+ if (!mod.lookupSymbol(rangeFuncName)) {
+ mlir::OpBuilder::InsertionGuard guard(builder);
+ builder.setInsertionPointToEnd(mod.getBody());
+ auto rangeFunc = mlir::LLVM::LLVMFuncOp::create(
+ builder, loc, rangeFuncName, rangeFuncTy);
+ rangeFunc.setVisibility(mlir::SymbolTable::Visibility::Private);
+ }
+ mlir::LLVM::CallOp::create(builder, loc, rangeFuncTy,
+ mlir::SymbolRefAttr::get(ctx, rangeFuncName),
+ mlir::ValueRange{begin, end});
+ }
+ if (initCtor) {
+ mlir::OpBuilder::InsertionGuard guard(builder);
+ builder.setInsertionPointToEnd(getInitCtorBlock());
+ mlir::LLVM::ReturnOp::create(builder, loc, mlir::ValueRange{});
+ }
if (emitCudaCompiledMarker) {
// Keep the sentinel reference alive: an unused non-volatile load would
// be folded away before it reaches the object file.
diff --git a/flang/lib/Optimizer/Transforms/CUDA/CUFSharedTypeInfo.cpp b/flang/lib/Optimizer/Transforms/CUDA/CUFSharedTypeInfo.cpp
index a4cee32c27a29..7e2ce5fb8c7b0 100644
--- a/flang/lib/Optimizer/Transforms/CUDA/CUFSharedTypeInfo.cpp
+++ b/flang/lib/Optimizer/Transforms/CUDA/CUFSharedTypeInfo.cpp
@@ -11,9 +11,11 @@
#include "flang/Optimizer/Dialect/CUF/CUFOps.h"
#include "flang/Optimizer/Dialect/FIROps.h"
#include "flang/Optimizer/Dialect/FIRType.h"
+#include "flang/Optimizer/Dialect/Support/FIRContext.h"
#include "flang/Optimizer/Support/InternalNames.h"
#include "flang/Optimizer/Transforms/Passes.h"
#include "mlir/Dialect/GPU/IR/GPUDialect.h"
+#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
#include "mlir/Dialect/OpenACC/OpenACC.h"
#include "mlir/IR/Builders.h"
#include "mlir/IR/SymbolTable.h"
@@ -45,6 +47,61 @@ static std::string getUnitTag(mlir::ModuleOp mod) {
return llvm::utohexstr(llvm::xxh3_64bits(key), /*LowerCase=*/true);
}
+/// Add the LLVM global \p name to llvm.used, creating the array or appending
+/// to the one already in \p mod.
+static mlir::LogicalResult
+addToUsed(mlir::ModuleOp mod, mlir::SymbolTable &symTab, llvm::StringRef name) {
+ constexpr llvm::StringLiteral usedName{"llvm.used"};
+ mlir::MLIRContext *ctx = mod.getContext();
+ mlir::Location loc = mod.getLoc();
+ auto ptrTy = mlir::LLVM::LLVMPointerType::get(ctx);
+ mlir::OpBuilder builder(ctx);
+
+ auto used = symTab.lookup<mlir::LLVM::GlobalOp>(usedName);
+ mlir::LLVM::ReturnOp oldReturn;
+ uint64_t numOld = 0;
+ if (used) {
+ auto oldTy =
+ mlir::dyn_cast<mlir::LLVM::LLVMArrayType>(used.getGlobalType());
+ mlir::Block *init = used.getInitializerBlock();
+ if (!oldTy || oldTy.getElementType() != ptrTy || !init)
+ return used.emitError("unexpected form of ") << usedName;
+ oldReturn = mlir::cast<mlir::LLVM::ReturnOp>(init->getTerminator());
+ numOld = oldTy.getNumElements();
+ builder.setInsertionPoint(oldReturn);
+ } else {
+ if (symTab.lookup(usedName))
+ return mod.emitError("unexpected ") << usedName << " symbol";
+ builder.setInsertionPointToEnd(mod.getBody());
+ used = mlir::LLVM::GlobalOp::create(
+ builder, loc, mlir::LLVM::LLVMArrayType::get(ptrTy, 1),
+ /*isConstant=*/false, mlir::LLVM::Linkage::Appending, usedName,
+ mlir::Attribute{});
+ used.setSection("llvm.metadata");
+ builder.createBlock(&used.getInitializerRegion());
+ }
+
+ auto usedTy = mlir::LLVM::LLVMArrayType::get(ptrTy, numOld + 1);
+ mlir::Value value = mlir::LLVM::UndefOp::create(builder, loc, usedTy);
+ for (uint64_t i = 0; i < numOld; ++i) {
+ int64_t pos = i;
+ mlir::Value elt = mlir::LLVM::ExtractValueOp::create(
+ builder, loc, oldReturn.getArg(), llvm::ArrayRef<int64_t>(pos));
+ value = mlir::LLVM::InsertValueOp::create(builder, loc, value, elt,
+ llvm::ArrayRef<int64_t>(pos));
+ }
+ mlir::Value addr = mlir::LLVM::AddressOfOp::create(builder, loc, ptrTy, name);
+ int64_t markerPos = numOld;
+ value = mlir::LLVM::InsertValueOp::create(builder, loc, value, addr,
+ llvm::ArrayRef<int64_t>(markerPos));
+ if (oldReturn)
+ oldReturn.getArgMutable().assign(value);
+ else
+ mlir::LLVM::ReturnOp::create(builder, loc, value);
+ used.setGlobalTypeAttr(mlir::TypeAttr::get(usedTy));
+ return mlir::success();
+}
+
class CUFSharedTypeInfo
: public fir::impl::CUFSharedTypeInfoBase<CUFSharedTypeInfo> {
public:
@@ -60,6 +117,7 @@ class CUFSharedTypeInfo
mlir::StringAttr::get(ctx, cudaSharedTypeInfoSection);
llvm::SmallVector<fir::GlobalOp> typeDescs;
+ bool definesTypeInfo = false;
for (fir::GlobalOp global : mod.getOps<fir::GlobalOp>()) {
if (!cuf::isTypeInfoGlobal(global))
continue;
@@ -71,11 +129,41 @@ class CUFSharedTypeInfo
if (global.isInitialized()) {
global.removeConstantAttr();
global.setSectionAttr(section);
+ definesTypeInfo = true;
}
if (cuf::isTypeDescriptorGlobal(global))
typeDescs.push_back(global);
}
+ std::string unitTag = getUnitTag(mod);
+ mlir::OpBuilder builder(ctx);
+
+ // The CUDA Fortran constructor registers the whole pages of the section
+ // with the CUDA runtime, so no other data may share them: a copy of a
+ // variable that is only partly in registered memory fails. A page-aligned
+ // global makes the section start on a page boundary. Emitted last, it
+ // also follows the type information of this unit, so the copy from the
+ // last unit linked marks a page after all the type information, and the
+ // constructor registers only the pages before it. This assumes the type
+ // information is in COMDAT groups, which puts each global in its own
+ // section ahead of the one opened for this global.
+ // Nothing references the global, so it is kept by llvm.used. It is created
+ // in the LLVM dialect because llvm.used can only refer to LLVM globals.
+ if (definesTypeInfo) {
+ std::string markerName = cudaSharedTypeInfoEndPrefix.str() + unitTag;
+ mlir::Location loc = mod.getLoc();
+ builder.setInsertionPointToEnd(mod.getBody());
+ auto i8Ty = mlir::IntegerType::get(ctx, 8);
+ auto marker = mlir::LLVM::GlobalOp::create(
+ builder, loc, i8Ty, /*isConstant=*/false,
+ mlir::LLVM::Linkage::Internal, markerName,
+ builder.getIntegerAttr(i8Ty, 0),
+ getCudaSharedTypeInfoAlignment(fir::getTargetTriple(mod)));
+ marker.setSection(section.getValue());
+ if (mlir::failed(addToUsed(mod, symTab, markerName)))
+ return signalPassFailure();
+ }
+
auto gpuMod = symTab.lookup<mlir::gpu::GPUModuleOp>(cudaDeviceModuleName);
if (!gpuMod)
return;
@@ -91,8 +179,6 @@ class CUFSharedTypeInfo
// The pointer names and the dictionary keys use the assembly form of the
// type descriptor name. CompilerGeneratedNamesConversion does not update
// names nested in the dictionary, and leaves names in this form unchanged.
- std::string unitTag = getUnitTag(mod);
- mlir::OpBuilder builder(ctx);
auto ptrTy = fir::LLVMPointerType::get(ctx, mlir::IntegerType::get(ctx, 8));
auto managed =
cuf::DataAttributeAttr::get(ctx, cuf::DataAttribute::Managed);
diff --git a/flang/test/Fir/CUDA/cuda-shared-type-info-no-managed.mlir b/flang/test/Fir/CUDA/cuda-shared-type-info-no-managed.mlir
new file mode 100644
index 0000000000000..2c799a033692f
--- /dev/null
+++ b/flang/test/Fir/CUDA/cuda-shared-type-info-no-managed.mlir
@@ -0,0 +1,35 @@
+// A unit defining shared type information but no managed variable or device
+// code still registers the pages of the section. The registration runs in the
+// initialization constructor, which is terminated and listed in the global
+// constructors.
+
+// The global ending the type information is aligned to the largest page size
+// of the target.
+
+// RUN: fir-opt --cuf-shared-type-info --cuf-add-constructor="cuda-managed-type-info=true" %s | FileCheck %s --check-prefixes=CHECK,X86
+// RUN: sed 's/x86_64-unknown-linux-gnu/aarch64-unknown-linux-gnu/' %s | fir-opt --cuf-shared-type-info | FileCheck %s --check-prefix=LARGEPAGE
+// RUN: sed 's/x86_64-unknown-linux-gnu/powerpc64le-unknown-linux-gnu/' %s | fir-opt --cuf-shared-type-info | FileCheck %s --check-prefix=LARGEPAGE
+
+module attributes {dlti.dl_spec = #dlti.dl_spec<i8 = dense<8> : vector<2xi64>, i16 = dense<16> : vector<2xi64>, i1 = dense<8> : vector<2xi64>, !llvm.ptr = dense<64> : vector<4xi64>, f80 = dense<128> : vector<2xi64>, i128 = dense<128> : vector<2xi64>, i64 = dense<64> : vector<2xi64>, !llvm.ptr<271> = dense<32> : vector<4xi64>, !llvm.ptr<272> = dense<64> : vector<4xi64>, f128 = dense<128> : vector<2xi64>, !llvm.ptr<270> = dense<32> : vector<4xi64>, f16 = dense<16> : vector<2xi64>, f64 = dense<64> : vector<2xi64>, i32 = dense<32> : vector<2xi64>, "dlti.stack_alignment" = 128 : i64, "dlti.endianness" = "little">, fir.defaultkind = "a1c4d8i4l4r4", fir.kindmap = "", llvm.data_layout = "e-m:e-p270:32:32-p271:32:32-p272:64:64-i64:64-i128:128-f80:128-n8:16:32:64-S128", llvm.target_triple = "x86_64-unknown-linux-gnu"} {
+ fir.global linkonce_odr @_QMmE.n.t constant : !fir.char<1> {
+ %0 = fir.string_lit "t"(1) : !fir.char<1>
+ fir.has_value %0 : !fir.char<1>
+ }
+}
+
+// X86: llvm.mlir.global internal @__nv_type_info_end_{{[0-9a-f]+}}(0 : i8) {{.*}}alignment = 4096
+// LARGEPAGE: llvm.mlir.global internal @__nv_type_info_end_{{[0-9a-f]+}}(0 : i8) {{.*}}alignment = 65536
+
+// CHECK-LABEL: llvm.func internal @__cudaFortranConstructor()
+// CHECK-NOT: CUFRegisterHostMemoryRange
+// CHECK-NOT: CUFInitModule
+// CHECK: llvm.return
+
+// CHECK-LABEL: llvm.func internal @__cudaFortranInitConstructor()
+// CHECK-NOT: CUFInitModule
+// CHECK: %[[START:.*]] = llvm.mlir.addressof @__start___nv_type_info : !llvm.ptr
+// CHECK: %[[STOP:.*]] = llvm.mlir.addressof @__stop___nv_type_info : !llvm.ptr
+// CHECK: llvm.call @_FortranACUFRegisterHostMemoryRange(%[[START]], %[[STOP]]) : (!llvm.ptr, !llvm.ptr) -> ()
+// CHECK-NEXT: llvm.return
+
+// CHECK: llvm.mlir.global_ctors ctors = [@__cudaFortranConstructor, @__cudaFortranInitConstructor], priorities = [0 : i32, 1 : i32]
diff --git a/flang/test/Fir/CUDA/cuda-shared-type-info-used.mlir b/flang/test/Fir/CUDA/cuda-shared-type-info-used.mlir
new file mode 100644
index 0000000000000..cd0fadd1d104f
--- /dev/null
+++ b/flang/test/Fir/CUDA/cuda-shared-type-info-used.mlir
@@ -0,0 +1,31 @@
+// The global ending the shared type information is appended to an existing
+// llvm.used array.
+
+// RUN: fir-opt --cuf-shared-type-info %s | FileCheck %s
+
+module attributes {fir.defaultkind = "a1c4d8i4l4r4", fir.kindmap = "", llvm.target_triple = "x86_64-unknown-linux-gnu"} {
+ fir.global linkonce_odr @_QMmE.n.t constant : !fir.char<1> {
+ %0 = fir.string_lit "t"(1) : !fir.char<1>
+ fir.has_value %0 : !fir.char<1>
+ }
+ llvm.mlir.global internal @keep(0 : i32) : i32
+ llvm.mlir.global appending @llvm.used() {section = "llvm.metadata"} : !llvm.array<1 x ptr> {
+ %0 = llvm.mlir.addressof @keep : !llvm.ptr
+ %1 = llvm.mlir.undef : !llvm.array<1 x ptr>
+ %2 = llvm.insertvalue %0, %1[0] : !llvm.array<1 x ptr>
+ llvm.return %2 : !llvm.array<1 x ptr>
+ }
+}
+
+// CHECK-NOT: @llvm.used
+// CHECK: llvm.mlir.global appending @llvm.used() {{.*}}: !llvm.array<2 x ptr> {
+// CHECK: %[[KEEP:.*]] = llvm.mlir.addressof @keep : !llvm.ptr
+// CHECK: %[[OLD0:.*]] = llvm.insertvalue %[[KEEP]], %{{.*}}[0] : !llvm.array<1 x ptr>
+// CHECK: %[[NEW:.*]] = llvm.mlir.undef : !llvm.array<2 x ptr>
+// CHECK: %[[ELT:.*]] = llvm.extractvalue %[[OLD0]][0] : !llvm.array<1 x ptr>
+// CHECK: %[[NEW0:.*]] = llvm.insertvalue %[[ELT]], %[[NEW]][0] : !llvm.array<2 x ptr>
+// CHECK: %[[MARKER:.*]] = llvm.mlir.addressof @__nv_type_info_end_{{[0-9a-f]+}} : !llvm.ptr
+// CHECK: %[[NEW1:.*]] = llvm.insertvalue %[[MARKER]], %[[NEW0]][1] : !llvm.array<2 x ptr>
+// CHECK: llvm.return %[[NEW1]] : !llvm.array<2 x ptr>
+// CHECK: llvm.mlir.global internal @__nv_type_info_end_{{[0-9a-f]+}}(0 : i8)
+// CHECK-NOT: @llvm.used
diff --git a/flang/test/Fir/CUDA/cuda-shared-type-info.mlir b/flang/test/Fir/CUDA/cuda-shared-type-info.mlir
index 220b146802d2a..000be0fbdd2d1 100644
--- a/flang/test/Fir/CUDA/cuda-shared-type-info.mlir
+++ b/flang/test/Fir/CUDA/cuda-shared-type-info.mlir
@@ -63,25 +63,34 @@ module attributes {dlti.dl_spec = #dlti.dl_spec<i8 = dense<8> : vector<2xi64>, i
// CHECK-DAG: fir.global linkonce_odr @_QMmE.dt.t constant : !fir.llvm_ptr<i8> {
// CHECK-DAG: fir.global @_QMmEXdtXtXhostaddr[[TAG]] <{data_attr = #cuf.cuda<managed>}> : !fir.llvm_ptr<i8> {
+// The unit ends its type information with a page-aligned global, so no other
+// data shares the pages registered by the constructor. llvm.used keeps the
+// global, which nothing references.
+// CHECK: llvm.mlir.global internal @__nv_type_info_end_[[TAG]](0 : i8)
+// CHECK-SAME: alignment = 4096
+// CHECK-SAME: section = "__nv_type_info"
+// CHECK: llvm.mlir.global appending @llvm.used()
+// CHECK-SAME: section = "llvm.metadata"
+// CHECK-SAME: !llvm.array<1 x ptr>
+// CHECK: llvm.mlir.addressof @__nv_type_info_end_[[TAG]] : !llvm.ptr
+
// NODECLARE-NOT: acc.declare
-// The constructor registers the pages of the section and the managed pointer,
-// and saves the module handle. A second constructor, which runs after the
-// registration constructors of all the units, initializes the module and then
-// stores the host address of the type descriptor in the managed pointer,
-// unless the runtime left the pointer null.
+// The constructor registers the managed pointer and saves the module handle,
+// without calling the CUDA runtime. A second constructor, which runs after the
+// registration constructors of all the units, initializes the module, stores
+// the host address of the type descriptor in the managed pointer unless the
+// runtime left the pointer null, and registers the pages of the section.
// CTOR-LABEL: llvm.func internal @__cudaFortranConstructor()
-// CTOR: %[[START:.*]] = llvm.mlir.addressof @__start___nv_type_info : !llvm.ptr
-// CTOR: %[[STOP:.*]] = llvm.mlir.addressof @__stop___nv_type_info : !llvm.ptr
-// CTOR: llvm.call @_FortranACUFRegisterHostMemoryRange(%[[START]], %[[STOP]]) : (!llvm.ptr, !llvm.ptr) -> ()
+// CTOR-NOT: CUFRegisterHostMemoryRange
// CTOR: %[[MOD:.*]] = cuf.register_module @cuda_device_mod -> !llvm.ptr
// CTOR: fir.call @_FortranACUFRegisterManagedVariable
// CTOR-NOT: fir.call @_FortranACUFInitModule
+// CTOR-NOT: CUFRegisterHostMemoryRange
// CTOR: %[[HANDLEADDR:.*]] = llvm.mlir.addressof @__cudaFortranModuleHandle : !llvm.ptr
// CTOR: llvm.store %[[MOD]], %[[HANDLEADDR]] : !llvm.ptr, !llvm.ptr
+// CTOR-NOT: CUFRegisterHostMemoryRange
// CTOR: llvm.return
-// CTOR-DAG: llvm.mlir.global extern_weak hidden @__start___nv_type_info() {{.*}}: i8
-// CTOR-DAG: llvm.mlir.global extern_weak hidden @__stop___nv_type_info() {{.*}}: i8
// CTOR-DAG: llvm.mlir.global internal @__cudaFortranModuleHandle() {{.*}}: !llvm.ptr
// CTOR-LABEL: llvm.func internal @__cudaFortranInitConstructor()
@@ -100,6 +109,11 @@ module attributes {dlti.dl_spec = #dlti.dl_spec<i8 = dense<8> : vector<2xi64>, i
// CTOR: fir.store %[[VAL]] to %[[DEST]] : !fir.ref<!fir.llvm_ptr<i8>>
// CTOR: llvm.br ^[[NEXT]]
// CTOR: ^[[NEXT]]:
+// CTOR: %[[START:.*]] = llvm.mlir.addressof @__start___nv_type_info : !llvm.ptr
+// CTOR: %[[STOP:.*]] = llvm.mlir.addressof @__stop___nv_type_info : !llvm.ptr
+// CTOR: llvm.call @_FortranACUFRegisterHostMemoryRange(%[[START]], %[[STOP]]) : (!llvm.ptr, !llvm.ptr) -> ()
// CTOR: llvm.return
+// CTOR-DAG: llvm.mlir.global extern_weak hidden @__start___nv_type_info() {{.*}}: i8
+// CTOR-DAG: llvm.mlir.global extern_weak hidden @__stop___nv_type_info() {{.*}}: i8
// CTOR: llvm.mlir.global_ctors ctors = [@__cudaFortranConstructor, @__cudaFortranInitConstructor], priorities = [0 : i32, 1 : i32]
More information about the flang-commits
mailing list