[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