================
@@ -346,3 +382,130 @@ mlir::Operation
*CIRGenNVCUDARuntime::getKernelHandle(cir::FuncOp fn,
return globalOp;
}
+
+void CIRGenNVCUDARuntime::internalizeDeviceSideVar(
+ const VarDecl *d, cir::GlobalLinkageKind &linkage) {
+ if (cgm.getLangOpts().GPURelocatableDeviceCode)
+ cgm.errorNYI(d->getSourceRange(),
+ "internalizeDeviceSideVar: GPU Relocatable Device Code
(RDC)");
+
+ // __shared__ variables are odd. Shadows do get created, but
+ // they are not registered with the CUDA runtime, so they
+ // can't really be used to access their device-side
+ // counterparts. It's not clear yet whether it's nvcc's bug or
+ // a feature, but we've got to do the same for compatibility.
+ if (d->hasAttr<CUDADeviceAttr>() || d->hasAttr<CUDAConstantAttr>() ||
+ d->hasAttr<CUDASharedAttr>()) {
+ linkage = cir::GlobalLinkageKind::InternalLinkage;
+ }
+
+ if (d->getType()->isCUDADeviceBuiltinSurfaceType() ||
+ d->getType()->isCUDADeviceBuiltinTextureType())
+ cgm.errorNYI(d->getSourceRange(),
+ "internalizeDeviceSideVar: CUDA Surface/Texture support");
+}
+
+std::string CIRGenNVCUDARuntime::getDeviceSideName(const NamedDecl *nd) {
+ GlobalDecl gd;
+ // nd could be either a kernel or a variable.
+ if (auto *fd = dyn_cast<FunctionDecl>(nd))
+ gd = GlobalDecl(fd, KernelReferenceKind::Kernel);
+ else
+ gd = GlobalDecl(nd);
+ std::string deviceSideName;
+ MangleContext *mc;
+ if (cgm.getLangOpts().CUDAIsDevice)
+ mc = &cgm.getCXXABI().getMangleContext();
+ else
+ mc = deviceMC.get();
+ if (mc->shouldMangleDeclName(nd)) {
+ SmallString<256> buffer;
+ llvm::raw_svector_ostream out(buffer);
+ mc->mangleName(gd, out);
+ deviceSideName = std::string(out.str());
+ } else
+ deviceSideName = std::string(nd->getIdentifier()->getName());
+
+ // Make unique name for device side static file-scope variable for HIP.
+ if (cgm.getASTContext().shouldExternalize(nd) &&
+ cgm.getLangOpts().GPURelocatableDeviceCode) {
+ SmallString<256> buffer;
+ llvm::raw_svector_ostream out(buffer);
+ out << deviceSideName;
+ cgm.printPostfixForExternalizedDecl(out, nd);
+ deviceSideName = std::string(out.str());
+ }
+ return deviceSideName;
+}
+
+void CIRGenNVCUDARuntime::handleVarRegistration(const VarDecl *vd,
+ cir::GlobalOp var) {
+ if (vd->hasAttr<CUDADeviceAttr>() || vd->hasAttr<CUDAConstantAttr>()) {
+ // Shadow variables and their properties must be registered with CUDA
+ // runtime. Skip Extern global variables, which will be registered in
+ // the TU where they are defined.
+ //
+ // Don't register a C++17 inline variable. The local symbol can be
+ // discarded and referencing a discarded local symbol from outside the
+ // comdat (__cuda_register_globals) is disallowed by the ELF spec.
+ //
+ // HIP managed variables need to be always recorded in device and host
+ // compilations for transformation.
+ //
+ // HIP managed variables and variables in CUDADeviceVarODRUsedByHost are
+ // added to llvm.compiler-used, therefore they are safe to be registered.
+ if ((!vd->hasExternalStorage() && !vd->isInline()) ||
+ cgm.getASTContext().CUDADeviceVarODRUsedByHost.contains(vd) ||
+ vd->hasAttr<HIPManagedAttr>()) {
+ registerDeviceVar(vd, var, !vd->hasDefinition(),
+ vd->hasAttr<CUDAConstantAttr>());
+ }
+ } else if (vd->getType()->isCUDADeviceBuiltinSurfaceType() ||
+ vd->getType()->isCUDADeviceBuiltinTextureType()) {
+ // Builtin surfaces and textures and their template arguments are
+ // also registered with CUDA runtime.
+ cgm.errorNYI(vd->getSourceRange(),
+ "handleVarRegistration: Surface and Texture registration");
+ }
+}
+
+void CIRGenNVCUDARuntime::handleGlobalReplace(cir::GlobalOp oldGV,
----------------
RiverDave wrote:
Functionality looks correct. My main concern is that this adds a CUDA-specific
hook inside CIRGenModule::replaceGlobal, which is language-agnostic. It handles
globals from any frontend.
The StringAttr + getGlobalValue approach I suggested earlier would let the hook
go away entirely (so symbolLookupCache already stays coherent across replaces).
I'd appreciate your thoughts here @koparasy.
https://github.com/llvm/llvm-project/pull/190087
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
