Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

[CIR][CUDA] Lowering device and shared variables #1438

Merged
merged 1 commit into from
Mar 5, 2025

Conversation

AdUhTkJm
Copy link
Contributor

@AdUhTkJm AdUhTkJm commented Mar 4, 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.

@bcardosolopes bcardosolopes merged commit 0e7b6c7 into llvm:main Mar 5, 2025
7 checks passed
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
None yet
Projects
None yet
Development

Successfully merging this pull request may close these issues.

2 participants