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

[Henry Linjamäki via Phabricator via cfe-commits]

linjamaki updated this revision to Diff 375221.
linjamaki added a comment.

Rebase.


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


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D108621.375221.patch
Type: text/x-patch
Size: 3137 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20210927/c2860b1a/attachment.bin>