Skip to content

[clang][OpenMP][SPIR-V] Fix addrspace of global constants #134399

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.

Sign up for GitHub

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 5 commits into from
Apr 9, 2025
Merged

[clang][OpenMP][SPIR-V] Fix addrspace of global constants #134399

Show file tree
Hide file tree
Changes from 1 commit
Commits
Show all changes
5 commits
Select commit Hold shift + click to select a range
3812b13
[clang][OpenMP][SPIR-V] Fix addrspace of globals and global constants
sarnex Apr 3, 2025
93ff204
do it in the target
sarnex Apr 7, 2025
06bfbba
whitespace
sarnex Apr 7, 2025
c4941ea
opencl_constant only for ocl
sarnex Apr 7, 2025
c9a24d0
revert global changes
sarnex Apr 8, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
9 changes: 9 additions & 0 deletions clang/lib/CodeGen/CodeGenModule.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -5384,6 +5384,11 @@ LangAS CodeGenModule::GetGlobalVarAddressSpace(const VarDecl *D) {
LangAS AS;
if (OpenMPRuntime->hasAllocateAttributeForGlobalVar(D, AS))
return AS;
if (LangOpts.OpenMPIsTargetDevice && getTriple().isSPIRV())
Copy link
Contributor

@jhuber6 jhuber6 Apr 4, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Is this really OpenMP specific? Sounds like a target info thing to me.

Copy link
Contributor

@AlexVlx AlexVlx Apr 4, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Yes I am a bit confused as to why this is necessary, DataLayout already encodes that global is AS1. If you're seeing globals end up in generic (I am excluding llvm.used and llvm.compiler.used here, since they are special and should be in generic/0) it might just be a case where CodeGen has a subtle bug. Could you please say a bit more as to what is motivating this change? Thank you!

Copy link
Member Author

@sarnex sarnex Apr 7, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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.

Copy link
Contributor

@AlexVlx AlexVlx Apr 7, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This feels like a spot where we are missing something in Clang - the string should've at least been AS1; some time ago I had a pop at fixing a bunch of places in CodeGen where we just used 0 / unqual rather than getting the GlobalVar AS or the Constant AS, but issues clearly remain - I think we should try to address this in Clang. Are you seeing the above with spirv64-unknown-unknown?

Copy link
Member Author

@sarnex sarnex Apr 7, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Yep, I see the problem with the string even with pure spirv64-unknown-unknown. Repro:

extern int printf(const char*);
int main() {
    printf("foo");
return 0;
}

clang++ -cc1 -triple spirv64-unknown-unknown -emit-llvm test.cpp -o -

; ModuleID = 'test.cpp'
source_filename = "test.cpp"
target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64-G1"
target triple = "spirv64-unknown-unknown"

@.str = private unnamed_addr constant [4 x i8] c"foo\00", align 1

; Function Attrs: mustprogress noinline norecurse nounwind optnone
define noundef i32 @main() #0 {
entry:
  %retval = alloca i32, align 4
  store i32 0, ptr %retval, align 4
  %call = call spir_func noundef i32 @_Z6printfPKc(ptr noundef @.str)
  ret i32 0
}

declare spir_func noundef i32 @_Z6printfPKc(ptr noundef) #1

attributes #0 = { mustprogress noinline norecurse nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
attributes #1 = { "no-trapping-math"="true" "stack-protector-buffer-size"="8" }

!llvm.module.flags = !{!0}
!llvm.ident = !{!1}

!0 = !{i32 1, !"wchar_size", i32 4}
!1 = !{!"clang version 21.0.0git (https://github.com/llvm/llvm-project.git 06bfbba877c26630b6c5b0ffef7f6623aa2e9ee8)"}

Here's where we get the addrspace from:

 unsigned AddrSpace = CGM.getContext().getTargetAddressSpace(
      CGM.GetGlobalConstantAddressSpace());

In GetGlobalConstantAddressSpace, we do

if (auto AS = getTarget().getConstantAddressSpace())
    return *AS;

and since there's no override for SPIR-V we just get the default:

 virtual std::optional<LangAS> getConstantAddressSpace() const {
    return LangAS::Default;
  }

If you see something wrong in this callstack let me know, I'm happy to fix it!

Copy link
Contributor

@AlexVlx AlexVlx Apr 7, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This seems ok. I would suggest that we consider, in the override, returning opencl_constant (2) only for OCL, and otherwise returning the global var AS (1), to prevent crashing into the invalid cast problem.

Copy link
Member Author

@sarnex sarnex Apr 7, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Sure, let me try that. Probably AS1 will be fine for my use case.

Copy link
Member Author

@sarnex sarnex Apr 7, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

BTW codegen already does the right thing because of this code in GetGlobalConstantAddressSpace, and CodeGenOpenCL/str_literals.cl already locks it down, so the OpenCL part of my change is basically NFC.

// SPIR-V globals should map to CrossWorkGroup instead of default
// AS, as generic/no address space is invalid. This is similar
// to what is done for HIPSPV.
return LangAS::opencl_global;
}
return getTargetCodeGenInfo().getGlobalVarAddressSpace(*this, D);
Copy link
Contributor

@jhuber6 jhuber6 Apr 4, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Wouldn't this logic fit in here?

}
Expand All @@ -5402,6 +5407,10 @@ LangAS CodeGenModule::GetGlobalConstantAddressSpace() const {
// UniformConstant storage class is not viable as pointers to it may not be
// casted to Generic pointers which are used to model HIP's "flat" pointers.
return LangAS::cuda_device;
if (LangOpts.OpenMPIsTargetDevice && getTriple().isSPIRV())
// OpenMP SPIR-V global constants should map to UniformConstant, different
// from the HIPSPV case above.
return LangAS::opencl_constant;
Copy link
Contributor

@AlexVlx AlexVlx Apr 4, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I believe that whilst this makes sense you might run into obnoxious issues where valid source ends up generating a SPIR-V invalid cast to/from generic/constant, which will fail in the translator / fail SPIR-V validation (HIP code runs into this). I think we need to relax this restriction at least in the translator / BE, or potentially extend SPIR-V itself in this direction.

Copy link
Member Author

@sarnex sarnex Apr 7, 2025 •
edited
Loading

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Yeah I hit something similar when trying to use addrspace(4) for generic. For constants it seems to be working okay for now, hopefully it's okay with you if I try this (but moved into the target) for now and then extend the solution if a problem comes up.

if (auto AS = getTarget().getConstantAddressSpace())
return *AS;
return LangAS::Default;
Expand Down
20 changes: 20 additions & 0 deletions clang/test/OpenMP/spirv_target_addrspace.c
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,20 @@
// RUN: %clang_cc1 -fopenmp -x c++ -triple x86_64-unknown-unknown -fopenmp-targets=spirv64 -emit-llvm-bc %s -o %t-host.bc
// RUN: %clang_cc1 -O0 -fopenmp -fopenmp-targets=spirv64 -fopenmp-is-target-device -triple spirv64 -fopenmp-host-ir-file-path %t-host.bc -emit-llvm %s -o - | FileCheck %s

extern int printf(char[]);

#pragma omp declare target
// CHECK: @global = addrspace(1) global i32 0, align 4
// CHECK: @.str = private unnamed_addr addrspace(2) constant [4 x i8] c"foo\00", align 1
int global = 0;
#pragma omp end declare target
int main() {
// CHECK: = call i32 @__kmpc_target_init(ptr addrspacecast (ptr addrspace(1) @__omp_offloading_{{.*}}_kernel_environment to ptr), ptr %{{.*}})
#pragma omp target
{
for(int i = 0; i < 5; i++)
global++;
printf("foo");
}
return global;
}
6 changes: 6 additions & 0 deletions llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -6295,6 +6295,12 @@ OpenMPIRBuilder::InsertPointTy OpenMPIRBuilder::createTargetInit(
: ConstantExpr::getAddrSpaceCast(KernelEnvironmentGV,
KernelEnvironmentPtr);
Value *KernelLaunchEnvironment = DebugKernelWrapper->getArg(0);
Type *KernelLaunchEnvParamTy = Fn->getFunctionType()->getParamType(1);
KernelLaunchEnvironment =
KernelLaunchEnvironment->getType() == KernelLaunchEnvParamTy
? KernelLaunchEnvironment
: Builder.CreateAddrSpaceCast(KernelLaunchEnvironment,
KernelLaunchEnvParamTy);
CallInst *ThreadKind =
Builder.CreateCall(Fn, {KernelEnvironment, KernelLaunchEnvironment});

Expand Down