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 @@ -219,6 +224,16 @@ bool hasFeature(StringRef Feature) const override { return Feature == "spirv"; } + + void adjust(DiagnosticsEngine &Diags, LangOptions &Opts) override { + BaseSPIRTargetInfo::adjust(Diags, Opts); + // Guarded so we don't override address space map setting set by + // BaseSPIRTargetInfo::adjust. + if (Opts.HIP && Opts.CUDAIsDevice) + // Enable address space mapping from HIP to SPIR-V. + // See comment on the SPIRDefIsGenMap table. + setAddressSpaceMap(/*DefaultIsGeneric=*/true); + } }; class LLVM_LIBRARY_VISIBILITY SPIRV32TargetInfo : public SPIRVTargetInfo { 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; +}