[clang] [llvm] [clang][OpenMP][SPIR-V] Fix addrspace of globals and global constants (PR #134399)
Nick Sarnie via llvm-commits
llvm-commits at lists.llvm.org
Mon Apr 7 11:28:51 PDT 2025
================
@@ -5384,6 +5384,11 @@ LangAS CodeGenModule::GetGlobalVarAddressSpace(const VarDecl *D) {
LangAS AS;
if (OpenMPRuntime->hasAllocateAttributeForGlobalVar(D, AS))
return AS;
+ if (LangOpts.OpenMPIsTargetDevice && getTriple().isSPIRV())
----------------
sarnex wrote:
Thanks for the feedback guys, the SPIR-V address space stuff is a total nightmare so I'll take any feedback I can get.
Here's the problem I'm trying to solve. For the code in the test I have:
```
extern int printf(char[]);
#pragma omp declare target
int global = 0;
#pragma omp end declare target
int main() {
#pragma omp target
{
for(int i = 0; i < 5; i++)
global++;
printf("foo");
}
return global;
}
```
Currently we get this IR
```
@global = global i32 0, align 4
@.str = private unnamed_addr constant [4 x i8] c"foo\00", align 1
```
Clearly the address space of both is wrong, `addrspace(0)` is not valid in SPIR-V for globals.
I think doing it in the target itself is much better, let me update the PR doing that, thanks.
https://github.com/llvm/llvm-project/pull/134399
More information about the llvm-commits
mailing list