Skip to content

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

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

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
bcardosolopes pushed a commit that referenced this pull request Mar 11, 2025
This PR implements \_\_constant\_\_ variables.

#1438 only implements \_\_device\_\_ and \_\_shared\_\_ variables, 

~~This PR depends on #1445~~
lanza pushed a commit that referenced this pull request Mar 18, 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.
lanza pushed a commit that referenced this pull request Mar 18, 2025
This PR implements \_\_constant\_\_ variables.

#1438 only implements \_\_device\_\_ and \_\_shared\_\_ variables, 

~~This PR depends on #1445~~
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