Skip to content

[OpenMP][OpenMPIRBuilder] Support SPIR-V device variant matches #126801

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 1 commit into from
Feb 12, 2025
Merged

[OpenMP][OpenMPIRBuilder] Support SPIR-V device variant matches #126801

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
3 changes: 2 additions & 1 deletion clang/lib/CodeGen/CodeGenModule.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1067,7 +1067,8 @@ class CodeGenModule : public CodeGenTypeCache {
bool shouldEmitRTTI(bool ForEH = false) {
return (ForEH || getLangOpts().RTTI) && !getLangOpts().CUDAIsDevice &&
!(getLangOpts().OpenMP && getLangOpts().OpenMPIsTargetDevice &&
(getTriple().isNVPTX() || getTriple().isAMDGPU()));
(getTriple().isNVPTX() || getTriple().isAMDGPU() ||
getTriple().isSPIRV()));
Comment on lines +1070 to +1071
Copy link
Contributor

Choose a reason for hiding this comment

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

I really hate this pattern, would be nice to have a single helper that's like isGPUTarget but probably out of scope here. Also we could probably rework this from the frontend to just pass -fno-rtti for the device -cc1 job.

Copy link
Member Author

Choose a reason for hiding this comment

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

yeah, after this pr ill try to find all spots that check for the gpu offload triples and replace them with a single helper, wasting hours tracking down bugs because i missed one of these spots is not exactly a great use of time :)

Copy link
Contributor

Choose a reason for hiding this comment

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

Probably a helper for whether or not it's an offloading target GPU device, so SYCL, OpenMP, CUDA, etc.

Copy link
Member Author

Choose a reason for hiding this comment

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

as in the same helper should be callable for all those languages? if so makes sense, hope my grep skills are state of the art.

Copy link
Contributor

Choose a reason for hiding this comment

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

I just mean most of these nonstandard hacks are only for offloading languages.

Copy link
Member Author

Choose a reason for hiding this comment

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

ah yeah

}

/// Get the address of the RTTI descriptor for the given type.
Expand Down
46 changes: 46 additions & 0 deletions clang/test/OpenMP/spirv_variant_match.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,46 @@
// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple x86_64-unknown-unknown -fopenmp-targets=spirv64-intel -emit-llvm-bc %s -o %t-host.bc -DDEVICE
// RUN: %clang_cc1 -verify -triple spirv64-intel -aux-triple x86_64-unknown-unknown -fopenmp -fopenmp-is-target-device \
// RUN:-fopenmp-host-ir-file-path %t-host.bc -nogpulib %s -emit-llvm -DDEVICE -o - | FileCheck %s

// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple x86_64-unknown-unknown -fopenmp-targets=spirv64-intel -emit-llvm-bc %s -o %t-host.bc -DTARGET
// RUN: %clang_cc1 -verify -triple spirv64-intel -aux-triple x86_64-unknown-unknown -fopenmp -fopenmp-is-target-device \
// RUN: -fopenmp-host-ir-file-path %t-host.bc -nogpulib %s -emit-llvm -DTARGET -o - | FileCheck %s

// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple x86_64-unknown-unknown -fopenmp-targets=spirv64-intel -emit-llvm-bc %s -o %t-host.bc -DTARGET_KIND
// RUN: %clang_cc1 -verify -triple spirv64-intel -aux-triple x86_64-unknown-unknown -fopenmp -fopenmp-is-target-device \
// RUN: -fopenmp-host-ir-file-path %t-host.bc -nogpulib %s -emit-llvm -DTARGET_KIND -o - | FileCheck %s

// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple x86_64-unknown-unknown -fopenmp-targets=spirv64-intel -emit-llvm-bc %s -o %t-host.bc
// RUN: %clang_cc1 -verify -triple spirv64-intel -aux-triple x86_64-unknown-unknown -fopenmp -fopenmp-is-target-device \
// RUN: -fopenmp-host-ir-file-path %t-host.bc -nogpulib %s -emit-llvm -o - | FileCheck %s

// expected-no-diagnostics

#pragma omp declare target
int foo() { return 0; }

#ifdef DEVICE
#pragma omp begin declare variant match(device = {arch(spirv64)})
#elif defined(TARGET)
#pragma omp begin declare variant match(target_device = {arch(spirv64)})
#elif defined(TARGET_KIND)
#pragma omp begin declare variant match(target_device = {kind(gpu)})
#else
#pragma omp begin declare variant match(device = {kind(gpu)})
#endif

int foo() { return 1; }
#pragma omp end declare variant
#pragma omp end declare target

// CHECK-DAG: define{{.*}} @{{"_Z[0-9]+foo\$ompvariant\$.*"}}()

// CHECK-DAG: call spir_func noundef i32 @{{"_Z[0-9]+foo\$ompvariant\$.*"}}()

int main() {
int res;
#pragma omp target map(from \
: res)
res = foo();
return res;
}
2 changes: 2 additions & 0 deletions llvm/include/llvm/Frontend/OpenMP/OMPKinds.def
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1269,6 +1269,7 @@ __OMP_TRAIT_PROPERTY(device, arch, x86_64)
__OMP_TRAIT_PROPERTY(device, arch, amdgcn)
__OMP_TRAIT_PROPERTY(device, arch, nvptx)
__OMP_TRAIT_PROPERTY(device, arch, nvptx64)
__OMP_TRAIT_PROPERTY(device, arch, spirv64)

__OMP_TRAIT_SET(target_device)

Expand Down Expand Up @@ -1301,6 +1302,7 @@ __OMP_TRAIT_PROPERTY(target_device, arch, x86_64)
__OMP_TRAIT_PROPERTY(target_device, arch, amdgcn)
__OMP_TRAIT_PROPERTY(target_device, arch, nvptx)
__OMP_TRAIT_PROPERTY(target_device, arch, nvptx64)
__OMP_TRAIT_PROPERTY(target_device, arch, spirv64)

__OMP_TRAIT_SET(implementation)

Expand Down
2 changes: 2 additions & 0 deletions llvm/lib/Frontend/OpenMP/OMPContext.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -52,6 +52,7 @@ OMPContext::OMPContext(bool IsDeviceCompilation, Triple TargetTriple,
case Triple::amdgcn:
case Triple::nvptx:
case Triple::nvptx64:
case Triple::spirv64:
ActiveTraits.set(unsigned(TraitProperty::target_device_kind_gpu));
break;
default:
Expand Down Expand Up @@ -98,6 +99,7 @@ OMPContext::OMPContext(bool IsDeviceCompilation, Triple TargetTriple,
case Triple::amdgcn:
case Triple::nvptx:
case Triple::nvptx64:
case Triple::spirv64:
ActiveTraits.set(unsigned(TraitProperty::device_kind_gpu));
ActiveTraits.set(unsigned(TraitProperty::target_device_kind_gpu));
break;
Expand Down