Skip to content

[Clang][OpenMP] New OpenMP 6.0 - Codegen support for groupprivate #169018

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

Open
Ritanya-B-Bharadwaj wants to merge 1 commit into
base: main
Choose a base branch
from
+42 −2
Open

[Clang][OpenMP] New OpenMP 6.0 - Codegen support for groupprivate #169018

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
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
6 changes: 5 additions & 1 deletion clang/lib/CodeGen/CGDecl.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -212,7 +212,11 @@ void CodeGenFunction::EmitVarDecl(const VarDecl &D) {
// Static sampler variables translated to function calls.
if (D.getType()->isSamplerT())
return;

if (D.hasAttr<OMPGroupPrivateDeclAttr>()) {
llvm::GlobalValue::LinkageTypes Linkage =
CGM.getLLVMLinkageVarDefinition(&D);
return EmitStaticVarDecl(D, Linkage);
}
llvm::GlobalValue::LinkageTypes Linkage =
CGM.getLLVMLinkageVarDefinition(&D);

Expand Down
11 changes: 11 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 @@ -5354,6 +5354,11 @@ CodeGenModule::GetOrCreateLLVMGlobal(StringRef MangledName, llvm::Type *Ty,
// Lookup the entry, lazily creating it if necessary.
llvm::GlobalValue *Entry = GetGlobalValue(MangledName);
unsigned TargetAS = getContext().getTargetAddressSpace(AddrSpace);
if (D && D->hasAttr<OMPGroupPrivateDeclAttr>() && getLangOpts().OpenMP &&
getTarget().getTriple().isGPU()) {
Entry->setLinkage(llvm::GlobalValue::InternalLinkage);
AddrSpace = LangAS::cuda_shared;
}
if (Entry) {
if (WeakRefReferences.erase(Entry)) {
if (D && !D->hasAttr<WeakAttr>())
Expand Down Expand Up @@ -5734,6 +5739,9 @@ LangAS CodeGenModule::GetGlobalVarAddressSpace(const VarDecl *D) {
LangAS AS;
if (OpenMPRuntime->hasAllocateAttributeForGlobalVar(D, AS))
return AS;
if (D && D->hasAttr<OMPGroupPrivateDeclAttr>()) {
return LangAS::cuda_shared; // maps to target addressspace 3 on NVPTX/AMD
}
}
return getTargetCodeGenInfo().getGlobalVarAddressSpace(*this, D);
}
Expand Down Expand Up @@ -7623,6 +7631,9 @@ void CodeGenModule::EmitTopLevelDecl(Decl *D) {
EmitOMPThreadPrivateDecl(cast<OMPThreadPrivateDecl>(D));
break;

case Decl::OMPGroupPrivate:
break;

case Decl::OMPAllocate:
EmitOMPAllocateDecl(cast<OMPAllocateDecl>(D));
break;
Expand Down
2 changes: 1 addition & 1 deletion clang/lib/Sema/SemaOpenMP.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -3275,7 +3275,7 @@ SemaOpenMP::ActOnOpenMPGroupPrivateDirective(SourceLocation Loc,
ArrayRef<Expr *> VarList) {
if (!getLangOpts().OpenMP || getLangOpts().OpenMP < 60) {
Diag(Loc, diag::err_omp_unexpected_directive)
<< getOpenMPDirectiveName(OMPD_groupprivate, getLangOpts().OpenMP);
<< 1 << getOpenMPDirectiveName(OMPD_groupprivate, getLangOpts().OpenMP);
return nullptr;
}
if (OMPGroupPrivateDecl *D = CheckOMPGroupPrivateDecl(Loc, VarList)) {
Expand Down
25 changes: 25 additions & 0 deletions clang/test/OpenMP/groupprivate_codegen.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,25 @@
// RUN: %clang_cc1 -fopenmp -fopenmp-version=60 -triple x86_64-unknown-unknown -fopenmp-targets=amdgcn-amd-amdhsa -fopenmp-is-device -emit-llvm -o - %s | FileCheck %s --check-prefix=CHECK
// RUN: %clang_cc1 -fopenmp -fopenmp-version=60 -triple x86_64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -fopenmp-is-device -emit-llvm -o - %s | FileCheck %s --check-prefix=CHECK
//
// This test ensures that variables marked 'groupprivate' are emitted as
// device globals in the GPU shared address space (addrspace(3)).
// The test is GPU-only and checks the LLVM IR for addrspace(3).
//

int group_var;

#pragma omp groupprivate(group_var)

void foo() {
#pragma omp target teams num_teams(4) thread_limit(100)
{
// simple use so the var is referenced in device codegen
group_var = group_var + 1;
}
}

// CHECK: @group_var = global i32 0, align 4, addrspace(3)
// CHECK: store i32 %{{.*}}, i32 addrspace(3)* @group_var, align 4

// CHECK: @group_var = global i32 0, align 4, addrspace(3)
// CHECK: store i32 %{{.*}}, i32 addrspace(3)* @group_var, align 4
Loading