diff --git a/clang/lib/Basic/Targets/SPIR.h b/clang/lib/Basic/Targets/SPIR.h --- a/clang/lib/Basic/Targets/SPIR.h +++ b/clang/lib/Basic/Targets/SPIR.h @@ -56,9 +56,14 @@ 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 @@ 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 @@ // 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 --- /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; +}