maxjhandsome
/
llvm-project
forked from OSchip/llvm-project
1
0
Fork 0
Code Issues Pull Requests Packages Releases Wiki Activity

[HIPSPV] Add CUDA->SPIR-V address space mapping 

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
This commit is contained in:
Anastasia Stulova 2021-12-02 13:30:05 +00:00
parent 4e9e2f2417
commit f4d3cb4ca8
2 changed files with 76 additions and 8 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

38
clang/lib/Basic/Targets/SPIR.h
View File

@ -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;

46
clang/test/CodeGenHIP/hipspv-addr-spaces.cpp Normal file
View File

@ -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;
}