[CIR][CUDA] initial support for __constant__ variables #1436
+18
−19
Conversation
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
getValueFromLangAS: implement cuda constant fix shouldEmitCUDAGlobalVar - copies og emitGlobalVarDefinition: set constant for cudaconstant
- also fixed __device__ check
You can test this locally with the following command:git-clang-format --diff 2ab0704058381f1c8c79d67e6304b0f5d771b4c8 76901bbf6b0c8b770f5e81d563f2cbdb4daa2dfd --extensions cpp -- clang/lib/CIR/CodeGen/CIRGenModule.cpp clang/lib/CIR/Dialect/IR/CIRAttrs.cppView the diff from clang-format here.diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.cpp b/clang/lib/CIR/CodeGen/CIRGenModule.cpp
index e9afbdf9d5..8b3896a7ba 100644
--- a/clang/lib/CIR/CodeGen/CIRGenModule.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenModule.cpp
@@ -518,10 +518,10 @@ bool CIRGenModule::shouldEmitCUDAGlobalVar(const VarDecl *global) const {
// size and host-side address in order to provide access to
// their device-side incarnations.
return !langOpts.CUDAIsDevice || global->hasAttr<CUDADeviceAttr>() ||
- global->hasAttr<CUDAConstantAttr>() ||
- global->hasAttr<CUDASharedAttr>() ||
- global->getType()->isCUDADeviceBuiltinSurfaceType() ||
- global->getType()->isCUDADeviceBuiltinTextureType();
+ global->hasAttr<CUDAConstantAttr>() ||
+ global->hasAttr<CUDASharedAttr>() ||
+ global->getType()->isCUDADeviceBuiltinSurfaceType() ||
+ global->getType()->isCUDADeviceBuiltinTextureType();
}
void CIRGenModule::emitGlobal(GlobalDecl gd) {
@@ -1448,7 +1448,7 @@ void CIRGenModule::emitGlobalVarDefinition(const clang::VarDecl *d,
// TODO(cir): If it is safe to mark the global 'constant', do so now.
gv.setConstant((d->hasAttr<CUDAConstantAttr>() && langOpts.CUDAIsDevice) ||
(!needsGlobalCtor && !needsGlobalDtor &&
- isTypeConstant(d->getType(), true, true)));
+ isTypeConstant(d->getType(), true, true)));
// If it is in a read-only section, mark it 'constant'.
if (const SectionAttr *sa = d->getAttr<SectionAttr>())
|
PikachuHyA
reviewed
Mar 5, 2025
| __constant__ int b; | ||
|
|
||
| // CIR-DEVICE: cir.global constant external addrspace(offload_constant) @b ={{.*}} |
There was a problem hiding this comment.
Could you please add a checker for the corresponding LLVM IR?
|
The CI check for Check Code Formatting has failed. Please format your code using clang-format or adjust it manually. |
bcardosolopes
requested changes
Mar 5, 2025
bcardosolopes
pushed a commit
that referenced
this pull request
Mar 5, 2025
Currently `__shared__` and `__constant__` variables are ignored by CodeGen. This patch fixes this. (It is also fixed in #1436 .) Device and constant variables should be marked as `externally_initialized`, as they might be initialized by host, rather than on device. We can't identify which variables are device ones at lowering stage, so this patch adds a new attribute for it in CodeGen. Similar to `__global__` functions, global variables on device corresponds to "shadow" variables on host, and they must be registered to their counterpart. I added a `CUDAShadowNameAttr` in this patch for later use, but I didn't insert code to actually generate it.
|
#1438 only implements __device__ and __shared__variables, but not __constant__ variables which this PR does @bcardosolopes. See #1444 |
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
This PR adds support for __constant__ variables on devices.