[HIPSPV] Convert HIP kernels to SPIR-V kernels

[linjamaki]

Summary

This patch translates HIP kernels to SPIR-V kernels when the HIP
compilation mode is targeting SPIR-S. This involves:

• Setting Cuda calling convention to CC_OpenCLKernel (which maps to SPIR_KERNEL in LLVM IR later on).

• Coercing pointer arguments with default address space (AS) qualifier to CrossWorkGroup AS (__global in OpenCL). HIPSPV's device code is ultimately SPIR-V for OpenCL execution environment (as starter/default) where Generic or Function (OpenCL's private) is not supported as storage class for kernel pointer types. This leaves the CrossWorkGroup to be the only reasonable choice for HIP buffers.

[linjamaki]

Rebase.

[linjamaki]

Rebase.

[yaxunl]

LGTM. Thanks!

[Anastasia]

clang/lib/CodeGen/TargetInfo.cpp

10236

It feels like this needs to be in SPIRVABIInfo or something? Or can this be generalized to both - SPIR and SPIR-V?

10239

Can you explain why this mapping is needed? We already have an address space map to perform the mapping of address spaces b/w language and target. It would be good if we don't replicate similar logic in too many places.

10288

Again this feels like we need SPIRVTargetCodeGenInfo unless the logic is generalizable to both.

[linjamaki]

clang/lib/CodeGen/TargetInfo.cpp

10236

A comment was added in D109144 to state that the SPIRABIInfo is an ABI implementation for both the SPIR and SPIR-V. For now, there is not much difference between SPIR and SPIR-V for this class. Would it be satisfactory if the class is renamed to something more general (like CommonSPIRABIInfo)?

10239

HIP does not require address space qualifiers on kernel pointer arguments (e.g. see hipspv-kernel.cpp) nor there are AS qualifiers that can be placed on them. With the default logic, provided by SPIRVTargetInfo’s address space map, the kernel pointer arguments get converted to generic pointers which are not allowed by the OpenCL SPIR-V Environment Specification.

[Anastasia]

clang/lib/CodeGen/TargetInfo.cpp

10236

It might be reasonable to rename indeed or we can just amend its documentation to clarify what it is used for both...

However if you need to specialize different logic I think the conventional way in clang would be to create a subclass...

This is related to a wider discussion that is ongoing about the best way to reuse SPIR for SPIR-V and their evolution... CCing @bader here too.

10239

I feel that it is the same for SYCL... It might be good to check with @bader whether there is already a way to handle this that can be reused for HIP...

[bader]

clang/lib/CodeGen/TargetInfo.cpp

10239

We need to do similar transformation for SYCL, but it's not exactly the same.
For SYCL kernels, which represented as function objects, compiler generates SPIR kernel function and fixes up the address space for pointer arguments in compiler generated declaration. For more details, see the description of https://reviews.llvm.org/D71016 and handlePointerType function code in clang/lib/Sema/SemaSYCL.cpp of this review request (lines 848-876). As address space is fixed in Sema, it works for all targets SYCL currently supports SPIR, NVPTX and AMDGPU.

If I understand it correctly, we are trying to do minimal amount of work for convert HIP kernel function to SPIR kernel function, i.e. fix calling convention and address spaces.
Are these two fixes enough or we need more fixes to enable more sophisticated kernels?

[linjamaki]

clang/lib/CodeGen/TargetInfo.cpp

10239

This patch and D108621 covers only the calling convention and the address space aspects in HIP->SPIR-V conversion. There are still various HIP features [1] which need to be expanded or emulated afterwards. A fully linked device program is needed before the fixes can be applied, so these fixes won’t be implemented at Sema nor CodeGen which operate per translation unit. The full-program fixes are applied by the HIPSPV tool chain by applying LLVM passes provided by a HIPSPV runtime [2]. For the time being, we are not seeing a need to specialize SPIR target infos further.

[1]: “HIP code expansion” section of the “[RFC][HIPSPV] Emitting HIP device code as SPIR-V”
[2]: clang/lib/Driver/ToolChains/HIPSPV.cpp:constructEmitSpirvCommand() @ https://github.com/parmance/llvm-project/tree/hipspv-wip

[linjamaki]

Put HIPSPV logic into new SPIRV{ABI,TargetCodeGen}Info subclassed from
CommonSPIR{ABI,TargetCodeGen}Info.

[linjamaki]

Rebase.

[linjamaki]

Rebase.

[linjamaki]

The patch is ready to land. @Anastasia, @bader, could you commit this patch to the LLVM for us? Thanks.

[bader]

In D109818#3169531, @linjamaki wrote:

The patch is ready to land. @Anastasia, @bader, could you commit this patch to the LLVM for us? Thanks.

Could you rebase on the tip of the main branch, please? I see a couple of conflicts when I cherry-pick the patch.

[linjamaki]

Rebase.

[linjamaki]

Cherry-picking should work now, @bader.

[linjamaki]

Thanks, @bader.

clang/lib/CodeGen/TargetInfo.cpp

 namespace {
 class CommonSPIRABIInfo : public DefaultABIInfo {
 public:
   CommonSPIRABIInfo(CodeGenTypes &CGT) : DefaultABIInfo(CGT) { setCCs(); }
 
 private:
   void setCCs();
 };
+
+class SPIRVABIInfo : public CommonSPIRABIInfo {
+public:
+  SPIRVABIInfo(CodeGenTypes &CGT) : CommonSPIRABIInfo(CGT) {}
+  void computeInfo(CGFunctionInfo &FI) const override;
+
+private:
+  ABIArgInfo classifyKernelArgumentType(QualType Ty) const;
+};
 } // end anonymous namespace
 namespace {
 class CommonSPIRTargetCodeGenInfo : public TargetCodeGenInfo {
 public:
   CommonSPIRTargetCodeGenInfo(CodeGen::CodeGenTypes &CGT)
       : TargetCodeGenInfo(std::make_unique<CommonSPIRABIInfo>(CGT)) {}
+  CommonSPIRTargetCodeGenInfo(std::unique_ptr<ABIInfo> ABIInfo)
+      : TargetCodeGenInfo(std::move(ABIInfo)) {}
 
   LangAS getASTAllocaAddressSpace() const override {
     return getLangASFromTargetAS(
         getABIInfo().getDataLayout().getAllocaAddrSpace());
   }
 
   unsigned getOpenCLKernelCallingConv() const override;
 };
-
+class SPIRVTargetCodeGenInfo : public CommonSPIRTargetCodeGenInfo {
+public:
+  SPIRVTargetCodeGenInfo(CodeGen::CodeGenTypes &CGT)
+      : CommonSPIRTargetCodeGenInfo(std::make_unique<SPIRVABIInfo>(CGT)) {}
+  void setCUDAKernelCallingConvention(const FunctionType *&FT) const override;
+};
 } // End anonymous namespace.
+
 void CommonSPIRABIInfo::setCCs() {
   assert(getRuntimeCC() == llvm::CallingConv::C);
   RuntimeCC = llvm::CallingConv::SPIR_FUNC;
 }
 
+ABIArgInfo SPIRVABIInfo::classifyKernelArgumentType(QualType Ty) const {
+  if (getContext().getLangOpts().HIP) {
+    // Coerce pointer arguments with default address space to CrossWorkGroup
+    // pointers for HIPSPV. When the language mode is HIP, the SPIRTargetInfo
+    // maps cuda_device to SPIR-V's CrossWorkGroup address space.
+    llvm::Type *LTy = CGT.ConvertType(Ty);
+    auto DefaultAS = getContext().getTargetAddressSpace(LangAS::Default);
+    auto GlobalAS = getContext().getTargetAddressSpace(LangAS::cuda_device);
+    if (LTy->isPointerTy() && LTy->getPointerAddressSpace() == DefaultAS) {
+      LTy = llvm::PointerType::get(
+          cast<llvm::PointerType>(LTy)->getElementType(), GlobalAS);
+      return ABIArgInfo::getDirect(LTy, 0, nullptr, false);
+    }
+  }
+  return classifyArgumentType(Ty);
+}
+
+void SPIRVABIInfo::computeInfo(CGFunctionInfo &FI) const {
+  // The logic is same as in DefaultABIInfo with an exception on the kernel
+  // arguments handling.
+  llvm::CallingConv::ID CC = FI.getCallingConvention();
+
+  if (!getCXXABI().classifyReturnType(FI))
+    FI.getReturnInfo() = classifyReturnType(FI.getReturnType());
+
+  for (auto &I : FI.arguments()) {
+    if (CC == llvm::CallingConv::SPIR_KERNEL) {
+      I.info = classifyKernelArgumentType(I.type);
+    } else {
+      I.info = classifyArgumentType(I.type);
+    }
+  }
+}
+
 namespace clang {
 namespace CodeGen {
 void computeSPIRKernelABIInfo(CodeGenModule &CGM, CGFunctionInfo &FI) {
-  DefaultABIInfo SPIRABI(CGM.getTypes());
-  SPIRABI.computeInfo(FI);
+  if (CGM.getTarget().getTriple().isSPIRV())
+    SPIRVABIInfo(CGM.getTypes()).computeInfo(FI);
+  else
+    CommonSPIRABIInfo(CGM.getTypes()).computeInfo(FI);
 }
 }
 }
 
 unsigned CommonSPIRTargetCodeGenInfo::getOpenCLKernelCallingConv() const {
   return llvm::CallingConv::SPIR_KERNEL;
 }
 
+void SPIRVTargetCodeGenInfo::setCUDAKernelCallingConvention(
+    const FunctionType *&FT) const {
+  // Convert HIP kernels to SPIR-V kernels.
+  if (getABIInfo().getContext().getLangOpts().HIP) {
+    FT = getABIInfo().getContext().adjustFunctionType(
+        FT, FT->getExtInfo().withCallingConv(CC_OpenCLKernel));
+    return;
+  }
+}
+
 static bool appendType(SmallStringEnc &Enc, QualType QType,
                        const CodeGen::CodeGenModule &CGM,
                        TypeStringCache &TSC);
 
 /// Helper function for appendRecordType().
 /// Builds a SmallVector containing the encoded field types in declaration
 /// order.
 static bool extractFieldType(SmallVectorImpl<FieldEncoding> &FE,
const TargetCodeGenInfo &CodeGenModule::getTargetCodeGenInfo() {
   case llvm::Triple::sparcv9:
     return SetCGInfo(new SparcV9TargetCodeGenInfo(Types));
   case llvm::Triple::xcore:
     return SetCGInfo(new XCoreTargetCodeGenInfo(Types));
   case llvm::Triple::arc:
     return SetCGInfo(new ARCTargetCodeGenInfo(Types));
   case llvm::Triple::spir:
   case llvm::Triple::spir64:
+    return SetCGInfo(new CommonSPIRTargetCodeGenInfo(Types));
   case llvm::Triple::spirv32:
   case llvm::Triple::spirv64:
-    return SetCGInfo(new CommonSPIRTargetCodeGenInfo(Types));
+    return SetCGInfo(new SPIRVTargetCodeGenInfo(Types));
   case llvm::Triple::ve:
     return SetCGInfo(new VETargetCodeGenInfo(Types));
   }
 }
 
 /// Create an OpenCL kernel for an enqueued block.
 ///
 /// The kernel has the same function type as the block invoke function. Its

clang/test/CodeGenHIP/hipspv-kernel.cpp

+// RUN: %clang_cc1 -triple spirv64 -x hip -emit-llvm -fcuda-is-device \
+// RUN:   -o - %s | FileCheck %s
+
+#define __global__ __attribute__((global))
+
+// CHECK: define {{.*}}spir_kernel void @_Z3fooPff(float addrspace(1)* {{.*}}, float {{.*}})
+__global__ void foo(float *a, float b) {
+  *a = b;
+}