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

[sarnex]

We should be able to use spirv64 as a device variant match and it should be considered a GPU.

Also add the triple to an RTTI check.

[llvmbot]

@llvm/pr-subscribers-flang-openmp
@llvm/pr-subscribers-clang

@llvm/pr-subscribers-clang-codegen

Author: Nick Sarnie (sarnex)

Changes

We should be able to use spirv64 as a device variant match and it should be considered a GPU.

Also add the triple to an RTTI check.

---

Full diff: https://github.com/llvm/llvm-project/pull/126801.diff

4 Files Affected:

• (modified) clang/lib/CodeGen/CodeGenModule.h (+2-1)
• (added) clang/test/OpenMP/spirv_variant_match.cpp (+46)
• (modified) llvm/include/llvm/Frontend/OpenMP/OMPKinds.def (+2)
• (modified) llvm/lib/Frontend/OpenMP/OMPContext.cpp (+2)

diff --git a/clang/lib/CodeGen/CodeGenModule.h b/clang/lib/CodeGen/CodeGenModule.h
index 0956296e2d5d8..c6f6fd5b9a7bd 100644
--- a/clang/lib/CodeGen/CodeGenModule.h
+++ b/clang/lib/CodeGen/CodeGenModule.h
@@ -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()));
   }
 
   /// Get the address of the RTTI descriptor for the given type.
diff --git a/clang/test/OpenMP/spirv_variant_match.cpp b/clang/test/OpenMP/spirv_variant_match.cpp
new file mode 100644
index 0000000000000..b37858bc3008b
--- /dev/null
+++ b/clang/test/OpenMP/spirv_variant_match.cpp
@@ -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;
+}
diff --git a/llvm/include/llvm/Frontend/OpenMP/OMPKinds.def b/llvm/include/llvm/Frontend/OpenMP/OMPKinds.def
index 44a9a37c70597..f974cfc78c8dd 100644
--- a/llvm/include/llvm/Frontend/OpenMP/OMPKinds.def
+++ b/llvm/include/llvm/Frontend/OpenMP/OMPKinds.def
@@ -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)
 
@@ -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)
 
diff --git a/llvm/lib/Frontend/OpenMP/OMPContext.cpp b/llvm/lib/Frontend/OpenMP/OMPContext.cpp
index 5e13da172d677..2edfd786c5c23 100644
--- a/llvm/lib/Frontend/OpenMP/OMPContext.cpp
+++ b/llvm/lib/Frontend/OpenMP/OMPContext.cpp
@@ -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:
@@ -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;

[jhuber6]

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.

[sarnex]

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 :)

[jhuber6]

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

[sarnex]

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.

[jhuber6]

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

[sarnex]

ah yeah