[clang] f3cd223 - [OpenMP][OpenMPIRBuilder] Add initial changes for SPIR-V target frontend support (#125920)

[via cfe-commits]

Author: Nick Sarnie
Date: 2025-02-10T16:16:40Z
New Revision: f3cd2238383f695c719e7eab6aebec828781ec91

URL: https://github.com/llvm/llvm-project/commit/f3cd2238383f695c719e7eab6aebec828781ec91
DIFF: https://github.com/llvm/llvm-project/commit/f3cd2238383f695c719e7eab6aebec828781ec91.diff

LOG: [OpenMP][OpenMPIRBuilder] Add initial changes for SPIR-V target frontend support (#125920)

As Intel is working to add support for SPIR-V OpenMP device offloading
in upstream clang/liboffload, we need to modify the OpenMP frontend to
allow SPIR-V as well as generate valid IR for SPIR-V. For example, we
need the frontend to generate code to define and interact with device
globals used in the DeviceRTL.

This is the beginning of what I expect will be (many) other changes, but
let's get started with something simple.

---------

Signed-off-by: Sarnie, Nick <nick.sarnie at intel.com>

Added: 
    clang/test/OpenMP/spirv_target_codegen_basic.cpp

Modified: 
    clang/include/clang/Basic/TargetInfo.h
    clang/lib/CodeGen/CodeGenModule.cpp
    llvm/include/llvm/Frontend/OpenMP/OMPGridValues.h
    llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp

Removed: 
    


################################################################################
diff  --git a/clang/include/clang/Basic/TargetInfo.h b/clang/include/clang/Basic/TargetInfo.h
index b9e46a5e7d1ca5e..070cc792ca7db62 100644
--- a/clang/include/clang/Basic/TargetInfo.h
+++ b/clang/include/clang/Basic/TargetInfo.h
@@ -1662,7 +1662,7 @@ class TargetInfo : public TransferrableTargetInfo,
   // access target-specific GPU grid values that must be consistent between
   // host RTL (plugin), deviceRTL and clang.
   virtual const llvm::omp::GV &getGridValue() const {
-    llvm_unreachable("getGridValue not implemented on this target");
+    return llvm::omp::SPIRVGridValues;
   }
 
   /// Retrieve the name of the platform as it is used in the

diff  --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp
index c056d103a7fe4f3..7924c32fcf633f5 100644
--- a/clang/lib/CodeGen/CodeGenModule.cpp
+++ b/clang/lib/CodeGen/CodeGenModule.cpp
@@ -486,8 +486,10 @@ void CodeGenModule::createOpenMPRuntime() {
   case llvm::Triple::nvptx:
   case llvm::Triple::nvptx64:
   case llvm::Triple::amdgcn:
-    assert(getLangOpts().OpenMPIsTargetDevice &&
-           "OpenMP AMDGPU/NVPTX is only prepared to deal with device code.");
+  case llvm::Triple::spirv64:
+    assert(
+        getLangOpts().OpenMPIsTargetDevice &&
+        "OpenMP AMDGPU/NVPTX/SPIRV is only prepared to deal with device code.");
     OpenMPRuntime.reset(new CGOpenMPRuntimeGPU(*this));
     break;
   default:

diff  --git a/clang/test/OpenMP/spirv_target_codegen_basic.cpp b/clang/test/OpenMP/spirv_target_codegen_basic.cpp
new file mode 100644
index 000000000000000..fb2810e88c06306
--- /dev/null
+++ b/clang/test/OpenMP/spirv_target_codegen_basic.cpp
@@ -0,0 +1,17 @@
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple x86_64-unknown-linux -fopenmp-targets=spirv64-intel -emit-llvm-bc %s -o %t-host.bc
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple spirv64-intel -fopenmp-targets=spirv64-intel -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-host.bc -o - | FileCheck %s
+
+// expected-no-diagnostics
+
+// CHECK: @__omp_offloading_{{.*}}_dynamic_environment = weak_odr protected addrspace(1) global %struct.DynamicEnvironmentTy zeroinitializer
+// CHECK: @__omp_offloading_{{.*}}_kernel_environment = weak_odr protected addrspace(1) constant %struct.KernelEnvironmentTy
+
+// CHECK: define weak_odr protected spir_kernel void @__omp_offloading_{{.*}}
+
+int main() {
+  int ret = 0;
+  #pragma omp target
+  for(int i = 0; i < 5; i++)
+    ret++;
+  return ret;
+}

diff  --git a/llvm/include/llvm/Frontend/OpenMP/OMPGridValues.h b/llvm/include/llvm/Frontend/OpenMP/OMPGridValues.h
index bfac2d734b81d8e..788a3c8a56f3806 100644
--- a/llvm/include/llvm/Frontend/OpenMP/OMPGridValues.h
+++ b/llvm/include/llvm/Frontend/OpenMP/OMPGridValues.h
@@ -120,6 +120,17 @@ static constexpr GV NVPTXGridValues = {
     128,       // GV_Default_WG_Size
 };
 
+/// For generic SPIR-V GPUs
+static constexpr GV SPIRVGridValues = {
+    256,       // GV_Slot_Size
+    64,        // GV_Warp_Size
+    (1 << 16), // GV_Max_Teams
+    440,       // GV_Default_Num_Teams
+    896,       // GV_SimpleBufferSize
+    1024,      // GV_Max_WG_Size,
+    256,       // GV_Default_WG_Size
+};
+
 } // namespace omp
 } // namespace llvm
 

diff  --git a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
index 91fc16e54c88fcc..f30eb64f1b4c93f 100644
--- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
+++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
@@ -159,6 +159,8 @@ static const omp::GV &getGridValue(const Triple &T, Function *Kernel) {
   }
   if (T.isNVPTX())
     return omp::NVPTXGridValues;
+  if (T.isSPIRV())
+    return omp::SPIRVGridValues;
   llvm_unreachable("No grid value available for this architecture!");
 }
 
@@ -6472,6 +6474,8 @@ void OpenMPIRBuilder::setOutlinedTargetRegionFunctionAttributes(
       OutlinedFn->setCallingConv(CallingConv::AMDGPU_KERNEL);
     else if (T.isNVPTX())
       OutlinedFn->setCallingConv(CallingConv::PTX_Kernel);
+    else if (T.isSPIRV())
+      OutlinedFn->setCallingConv(CallingConv::SPIR_KERNEL);
   }
 }