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

[Arseniy Obolenskiy via cfe-commits]

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

>From d30dff94798e367aad14aa35266fb0e1200e0611 Mon Sep 17 00:00:00 2001
From: Arseniy Obolenskiy <arseniy.obolenskiy at amd.com>
Date: Thu, 22 Oct 2026 11:12:46 +0200
Subject: [PATCH] [HIPSTDPAR][SPIR-V] Select accelerator code for AMD SPIR-V

AMD SPIR-V skips the AMDGPU pipeline that runs these passes, so host code stayed in device SPIR-V and broke the translator
---
 clang/lib/CodeGen/BackendUtil.cpp              |  8 ++++++++
 .../spirv-select-accelerator-code.cpp          | 18 ++++++++++++++++++
 llvm/lib/Transforms/HipStdPar/HipStdPar.cpp    |  2 +-
 .../accelerator-code-selection-spirv.ll        | 18 ++++++++++++++++++
 4 files changed, 45 insertions(+), 1 deletion(-)
 create mode 100644 clang/test/CodeGenHipStdPar/spirv-select-accelerator-code.cpp
 create mode 100644 llvm/test/Transforms/HipStdPar/accelerator-code-selection-spirv.ll

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
+}