linjamaki updated this revision to Diff 387500.
linjamaki added a comment.
Rebase, add asserts and move address space map reset for HIP from
SPIRVTargetInfo to BaseSPIRTargetInfo
Repository:
rG LLVM Github Monorepo
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D108621/new/
https://reviews.llvm.org/D108621
Files:
clang/lib/Basic/Targets/SPIR.h
clang/test/CodeGenHIP/hipspv-addr-spaces.cpp
Index: clang/test/CodeGenHIP/hipspv-addr-spaces.cpp
===================================================================
--- /dev/null
+++ 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;
+}
Index: clang/lib/Basic/Targets/SPIR.h
===================================================================
--- clang/lib/Basic/Targets/SPIR.h
+++ 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;
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
