[clang] [llvm] [HIPSTDPAR][SPIR-V] Select accelerator code for AMD SPIR-V (PR #234542)

[via cfe-commits]

llvmorg-github-actions[bot] wrote:


<!--LLVM PR SUMMARY COMMENT-->

@llvm/pr-subscribers-llvm-transforms

Author: Arseniy Obolenskiy (aobolensk)

<details>
<summary>Changes</summary>

AMD SPIR-V skips the AMDGPU pipeline that runs these passes, so host code stayed in device SPIR-V and broke the translator

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


4 Files Affected:

- (modified) clang/lib/CodeGen/BackendUtil.cpp (+8) 
- (added) clang/test/CodeGenHipStdPar/spirv-select-accelerator-code.cpp (+18) 
- (modified) llvm/lib/Transforms/HipStdPar/HipStdPar.cpp (+1-1) 
- (added) llvm/test/Transforms/HipStdPar/accelerator-code-selection-spirv.ll (+18) 


``````````diff
diff --git a/clang/lib/CodeGen/BackendUtil.cpp b/clang/lib/CodeGen/BackendUtil.cpp
index da7aa0b4b253a8..2c981b6f152eb8 100644
--- a/clang/lib/CodeGen/BackendUtil.cpp
+++ b/clang/lib/CodeGen/BackendUtil.cpp
@@ -1182,6 +1182,14 @@ void EmitAssemblyHelper::RunOptimizationPipeline(
     MPM.addPass(AssignGUIDPass());
   }
 
+  // AMD SPIR-V does not run the AMDGPU pipeline that selects accelerator code.
+  if (LangOpts.HIPStdPar && LangOpts.CUDAIsDevice &&
+      !LangOpts.GPURelocatableDeviceCode && TargetTriple.isSPIRV() &&
+      TargetTriple.getVendor() == llvm::Triple::VendorType::AMD) {
+    MPM.addPass(HipStdParMathFixupPass());
+    MPM.addPass(HipStdParAcceleratorCodeSelectionPass());
+  }
+
   if (LangOpts.HIPStdPar && !LangOpts.CUDAIsDevice &&
       LangOpts.HIPStdParInterposeAlloc)
     MPM.addPass(HipStdParAllocationInterpositionPass());
diff --git a/clang/test/CodeGenHipStdPar/spirv-select-accelerator-code.cpp b/clang/test/CodeGenHipStdPar/spirv-select-accelerator-code.cpp
new file mode 100644
index 00000000000000..fd26178827bcce
--- /dev/null
+++ b/clang/test/CodeGenHipStdPar/spirv-select-accelerator-code.cpp
@@ -0,0 +1,18 @@
+// RUN: %clang_cc1 -triple spirv64-amd-amdhsa -x hip --hipstdpar -fcuda-is-device \
+// RUN:   -disable-llvm-optzns -emit-llvm -o - %s \
+// RUN:   | FileCheck --implicit-check-not='@foo(' %s
+
+// RUN: %clang_cc1 -triple spirv64-amd-amdhsa -x hip --hipstdpar -fcuda-is-device \
+// RUN:   -fgpu-rdc -disable-llvm-optzns -emit-llvm -o - %s \
+// RUN:   | FileCheck --check-prefix=RDC %s
+
+#define __global__ __attribute__((global))
+
+// RDC: define {{.*}} @foo(
+extern "C" void foo() {}
+
+// CHECK: define {{.*}} @bar(
+extern "C" void bar() {}
+
+// CHECK: define spir_kernel void @kern(
+extern "C" __global__ void kern() { bar(); }
diff --git a/llvm/lib/Transforms/HipStdPar/HipStdPar.cpp b/llvm/lib/Transforms/HipStdPar/HipStdPar.cpp
index f0a5da16a90a0e..9cdcf989ca1d5f 100644
--- a/llvm/lib/Transforms/HipStdPar/HipStdPar.cpp
+++ b/llvm/lib/Transforms/HipStdPar/HipStdPar.cpp
@@ -367,7 +367,7 @@ static inline bool isAcceleratorExecutionRoot(const Function *F) {
     if (!F)
       return false;
 
-    return F->getCallingConv() == CallingConv::AMDGPU_KERNEL;
+    return F->hasKernelCallingConv();
 }
 
 static inline bool isCXXExceptionRuntimeFunction(StringRef Name) {
diff --git a/llvm/test/Transforms/HipStdPar/accelerator-code-selection-spirv.ll b/llvm/test/Transforms/HipStdPar/accelerator-code-selection-spirv.ll
new file mode 100644
index 00000000000000..b456b93dc78ab8
--- /dev/null
+++ b/llvm/test/Transforms/HipStdPar/accelerator-code-selection-spirv.ll
@@ -0,0 +1,18 @@
+; RUN: opt -S -mtriple=spirv64-amd-amdhsa -passes=hipstdpar-select-accelerator-code \
+; RUN: %s | FileCheck %s
+
+; CHECK-NOT: define {{.*}} @should_be_removed
+define void @should_be_removed() {
+  ret void
+}
+
+; CHECK: define {{.*}} @directly_called
+define void @directly_called() {
+  ret void
+}
+
+; CHECK: define spir_kernel void @accelerator_execution_root
+define spir_kernel void @accelerator_execution_root() {
+  call void @directly_called()
+  ret void
+}

``````````

</details>


https://github.com/llvm/llvm-project/pull/234542