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.
Reviewed By: yaxunl
Differential Revision: https://reviews.llvm.org/D109818
Add mapping for CUDA address spaces for HIP to SPIR-V
translation. This change allows HIP device code to be
emitted as valid SPIR-V by mapping unqualified pointers
to generic address space and by mapping __device__ and
__shared__ AS to their equivalent AS in SPIR-V
(CrossWorkgroup and Workgroup, respectively).
Cuda's __constant__ AS is handled specially. In HIP
unqualified pointers (aka "flat" pointers) can point to
__constant__ objects. Mapping this AS to ConstantMemory
would produce to illegal address space casts to
generic AS. Therefore, __constant__ AS is mapped to
CrossWorkgroup.
Patch by linjamaki (Henry Linjamäki)!
Differential Revision: https://reviews.llvm.org/D108621
Summary:
This change implements the expansion in two parts:
- Add a utility function emitAMDGPUPrintfCall() in LLVM.
- Invoke the above function from Clang CodeGen, when processing a HIP
program for the AMDGPU target.
The printf expansion has undefined behaviour if the format string is
not a compile-time constant. As a sufficient condition, the HIP
ToolChain now emits -Werror=format-nonliteral.
Reviewed By: arsenm
Differential Revision: https://reviews.llvm.org/D71365
Henry Linjamäki