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 @@ -54,9 +54,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 @@ -137,6 +142,8 @@ void adjust(DiagnosticsEngine &Diags, LangOptions &Opts) override { TargetInfo::adjust(Diags, Opts); + // See comment on the SPIRDefIsGenMap table. + bool IsHIPSPV = Opts.HIP && Opts.CUDAIsDevice; // 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. @@ -144,7 +151,7 @@ // 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); + setAddressSpaceMap(/*DefaultIsGeneric=*/Opts.SYCLIsDevice || IsHIPSPV); } void setSupportedOpenCLOpts() override { 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 spir64 -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; +}