From f4d3cb4ca833d0b165d199e78ed8f1d59a8032eb Mon Sep 17 00:00:00 2001 From: Anastasia Stulova Date: Thu, 2 Dec 2021 13:30:05 +0000 Subject: [PATCH] [HIPSPV] Add CUDA->SPIR-V address space mapping MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit 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 --- clang/lib/Basic/Targets/SPIR.h | 38 ++++++++++++---- clang/test/CodeGenHIP/hipspv-addr-spaces.cpp | 46 ++++++++++++++++++++ 2 files changed, 76 insertions(+), 8 deletions(-) create mode 100644 clang/test/CodeGenHIP/hipspv-addr-spaces.cpp diff --git a/clang/lib/Basic/Targets/SPIR.h b/clang/lib/Basic/Targets/SPIR.h index 704b1843dfed..8cf18b6c20f1 100644 --- a/clang/lib/Basic/Targets/SPIR.h +++ b/clang/lib/Basic/Targets/SPIR.h @@ -56,9 +56,14 @@ static const unsigned SPIRDefIsGenMap[] = { 0, // opencl_generic 0, // opencl_global_device 0, // opencl_global_host - 0, // cuda_device - 0, // cuda_constant - 0, // cuda_shared + // cuda_* address space mapping is intended for HIPSPV (HIP to SPIR-V + // translation). This mapping is enabled when the language mode is HIP. + 1, // cuda_device + // cuda_constant pointer can be casted to default/"flat" pointer, but in + // SPIR-V casts between constant and generic pointers are not allowed. For + // this reason cuda_constant is mapped to SPIR-V CrossWorkgroup. + 1, // cuda_constant + 3, // cuda_shared 1, // sycl_global 5, // sycl_global_device 6, // sycl_global_host @@ -74,6 +79,8 @@ class LLVM_LIBRARY_VISIBILITY BaseSPIRTargetInfo : public TargetInfo { protected: BaseSPIRTargetInfo(const llvm::Triple &Triple, const TargetOptions &) : TargetInfo(Triple) { + assert((Triple.isSPIR() || Triple.isSPIRV()) && + "Invalid architecture for SPIR or SPIR-V."); assert(getTriple().getOS() == llvm::Triple::UnknownOS && "SPIR(-V) target must use unknown OS"); assert(getTriple().getEnvironment() == llvm::Triple::UnknownEnvironment && @@ -137,11 +144,16 @@ public: // FIXME: SYCL specification considers unannotated pointers and references // to be pointing to the generic address space. See section 5.9.3 of // SYCL 2020 specification. - // Currently, there is no way of representing SYCL's default address space - // language semantic along with the semantics of embedded C's default - // address space in the same address space map. Hence the map needs to be - // reset to allow mapping to the desired value of 'Default' entry for SYCL. - setAddressSpaceMap(/*DefaultIsGeneric=*/Opts.SYCLIsDevice); + // Currently, there is no way of representing SYCL's and HIP's default + // address space language semantic along with the semantics of embedded C's + // default address space in the same address space map. Hence the map needs + // to be reset to allow mapping to the desired value of 'Default' entry for + // SYCL and HIP. + setAddressSpaceMap( + /*DefaultIsGeneric=*/Opts.SYCLIsDevice || + // The address mapping from HIP language for device code is only defined + // for SPIR-V. + (getTriple().isSPIRV() && Opts.HIP && Opts.CUDAIsDevice)); } void setSupportedOpenCLOpts() override { @@ -159,6 +171,7 @@ class LLVM_LIBRARY_VISIBILITY SPIRTargetInfo : public BaseSPIRTargetInfo { public: SPIRTargetInfo(const llvm::Triple &Triple, const TargetOptions &Opts) : BaseSPIRTargetInfo(Triple, Opts) { + assert(Triple.isSPIR() && "Invalid architecture for SPIR."); assert(getTriple().getOS() == llvm::Triple::UnknownOS && "SPIR target must use unknown OS"); assert(getTriple().getEnvironment() == llvm::Triple::UnknownEnvironment && @@ -177,6 +190,8 @@ class LLVM_LIBRARY_VISIBILITY SPIR32TargetInfo : public SPIRTargetInfo { public: SPIR32TargetInfo(const llvm::Triple &Triple, const TargetOptions &Opts) : SPIRTargetInfo(Triple, Opts) { + assert(Triple.getArch() == llvm::Triple::spir && + "Invalid architecture for 32-bit SPIR."); PointerWidth = PointerAlign = 32; SizeType = TargetInfo::UnsignedInt; PtrDiffType = IntPtrType = TargetInfo::SignedInt; @@ -192,6 +207,8 @@ class LLVM_LIBRARY_VISIBILITY SPIR64TargetInfo : public SPIRTargetInfo { public: SPIR64TargetInfo(const llvm::Triple &Triple, const TargetOptions &Opts) : SPIRTargetInfo(Triple, Opts) { + assert(Triple.getArch() == llvm::Triple::spir64 && + "Invalid architecture for 64-bit SPIR."); PointerWidth = PointerAlign = 64; SizeType = TargetInfo::UnsignedLong; PtrDiffType = IntPtrType = TargetInfo::SignedLong; @@ -207,6 +224,7 @@ class LLVM_LIBRARY_VISIBILITY SPIRVTargetInfo : public BaseSPIRTargetInfo { public: SPIRVTargetInfo(const llvm::Triple &Triple, const TargetOptions &Opts) : BaseSPIRTargetInfo(Triple, Opts) { + assert(Triple.isSPIRV() && "Invalid architecture for SPIR-V."); assert(getTriple().getOS() == llvm::Triple::UnknownOS && "SPIR-V target must use unknown OS"); assert(getTriple().getEnvironment() == llvm::Triple::UnknownEnvironment && @@ -225,6 +243,8 @@ class LLVM_LIBRARY_VISIBILITY SPIRV32TargetInfo : public SPIRVTargetInfo { public: SPIRV32TargetInfo(const llvm::Triple &Triple, const TargetOptions &Opts) : SPIRVTargetInfo(Triple, Opts) { + assert(Triple.getArch() == llvm::Triple::spirv32 && + "Invalid architecture for 32-bit SPIR-V."); PointerWidth = PointerAlign = 32; SizeType = TargetInfo::UnsignedInt; PtrDiffType = IntPtrType = TargetInfo::SignedInt; @@ -240,6 +260,8 @@ class LLVM_LIBRARY_VISIBILITY SPIRV64TargetInfo : public SPIRVTargetInfo { public: SPIRV64TargetInfo(const llvm::Triple &Triple, const TargetOptions &Opts) : SPIRVTargetInfo(Triple, Opts) { + assert(Triple.getArch() == llvm::Triple::spirv64 && + "Invalid architecture for 64-bit SPIR-V."); PointerWidth = PointerAlign = 64; SizeType = TargetInfo::UnsignedLong; PtrDiffType = IntPtrType = TargetInfo::SignedLong; diff --git a/clang/test/CodeGenHIP/hipspv-addr-spaces.cpp b/clang/test/CodeGenHIP/hipspv-addr-spaces.cpp new file mode 100644 index 000000000000..44289866e58a --- /dev/null +++ b/clang/test/CodeGenHIP/hipspv-addr-spaces.cpp @@ -0,0 +1,46 @@ +// RUN: %clang_cc1 -triple spirv64 -x hip -emit-llvm -fcuda-is-device \ +// RUN: -o - %s | FileCheck %s + +#define __device__ __attribute__((device)) +#define __shared__ __attribute__((shared)) +#define __constant__ __attribute__((constant)) + +// CHECK: %struct.foo_t = type { i32, i32 addrspace(4)* } + +// CHECK: @d ={{.*}} addrspace(1) externally_initialized global +__device__ int d; + +// CHECK: @c ={{.*}} addrspace(1) externally_initialized global +__constant__ int c; + +// CHECK: @s ={{.*}} addrspace(3) global +__shared__ int s; + +// CHECK: @foo ={{.*}} addrspace(1) externally_initialized global %struct.foo_t +__device__ struct foo_t { + int i; + int* pi; +} foo; + +// CHECK: define{{.*}} spir_func i32 addrspace(4)* @_Z3barPi(i32 addrspace(4)* +__device__ int* bar(int *x) { + return x; +} + +// CHECK: define{{.*}} spir_func i32 addrspace(4)* @_Z5baz_dv() +__device__ int* baz_d() { + // CHECK: ret i32 addrspace(4)* addrspacecast (i32 addrspace(1)* @d to i32 addrspace(4)* + return &d; +} + +// CHECK: define{{.*}} spir_func i32 addrspace(4)* @_Z5baz_cv() +__device__ int* baz_c() { + // CHECK: ret i32 addrspace(4)* addrspacecast (i32 addrspace(1)* @c to i32 addrspace(4)* + return &c; +} + +// CHECK: define{{.*}} spir_func i32 addrspace(4)* @_Z5baz_sv() +__device__ int* baz_s() { + // CHECK: ret i32 addrspace(4)* addrspacecast (i32 addrspace(3)* @s to i32 addrspace(4)* + return &s; +}