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

[linjamaki]

Summary

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.

Depends on D109144.

[bader]

LGTM.

clang/lib/Basic/Targets/SPIR.h

143

Minor: in my opinion, Opts.HIP check is unnecessary. I don't think CUDA can be compiled to SPIR target today, but when this flow is enabled I think it should set DefaultIsGeneric flag the same way as HIP.

[Anastasia]

clang/lib/Basic/Targets/SPIR.h

61

I am slightly confused as in the LLVM project those address spaces are for SPIR not SPIR-V though. It is however used outside of LLVM project by some tools like SPIRV-LLVM Translator as a path to SPIR-V, but it has only been done as a workaround since we had no SPIR-V support in the LLVM project yet. And if we are adding it let's do it clean to avoid/resolve any confusion.

I think we need to keep both because some vendors do target/use SPIR but not SPIR-V.

So if you are interested in SPIR-V and not SPIR you should probably add a new target that will make things cleaner.

[linjamaki]

Thanks. I will make a patch for adding spirv triples and new target info and update this.

[bader]

clang/lib/Basic/Targets/SPIR.h

61

I think we need to keep both because some vendors do target/use SPIR but not SPIR-V.

@Anastasia, could you elaborate more on the difference between SPIR and SPIR-V?
I would like to understand what these terms mean in the context of LLVM project.

[linjamaki]

Enable HIP-to-SPIR-V address space mapping only for SPIR-V targets.

Patch now depends on D109144.

[Anastasia]

clang/lib/Basic/Targets/SPIR.h

61

Their conceptual differences are just that they are two different intermediate formats.

The important thing to highlight is that it is not impossible that some vendors use SPIR (without using SPIR-V) even despite the fact it has been discontinued by Khronos.

Nobody has deprecated or discontinued SPIR in the LLVM project yet.

[Anastasia]

clang/lib/Basic/Targets/SPIR.h

233

My guess is that this is not only HIP specific but for example the same applies to SYCL.

I am not sure if it makes more sense to move this into a BaseSPIRTargetInfo since this is not really SPIR-V specific logic. It is just a clang design misalignment between two address space concepts that has to be addressed properly at some point.

[keryell]

clang/lib/Basic/Targets/SPIR.h

61

Their conceptual differences are just that they are two different intermediate formats.

The important thing to highlight is that it is not impossible that some vendors use SPIR (without using SPIR-V) even despite the fact it has been discontinued by Khronos.

Nobody has deprecated or discontinued SPIR in the LLVM project yet.

All the official Xilinx OpenCL stack is based on legacy SPIR (encoded in LLVM 6.x IR but this is another story) and I suspect this is the case for other companies.
So, do not deprecate or discontinue, please. :-)

[bader]

clang/lib/Basic/Targets/SPIR.h

61

The important thing to highlight is that it is not impossible that some vendors use SPIR (without using SPIR-V) even despite the fact it has been discontinued by Khronos.
Nobody has deprecated or discontinued SPIR in the LLVM project yet.

Strictly speaking SPIR is not defined as an intermediate language. Khronos defines SPIR-1.2 and SPIR-2.0 formats which are based on LLVM 3.2 and LLVM 3.4 version (https://www.khronos.org/spir/). There is no definition of SPIR format based on current version of LLVM IR. Another note is that metadata and intrinsics emitted for OpenCL with clang-14 doesn't follow neither SPIR-1.2 nor SPIR-2.0.

I always think of LLVM IR as leaving thing that is subject to change by LLVM community, so tools working with LLVM IR must adjust to the particular version (e.g. release version like LLVM 13 or ToT). We apply this logic to SPIRV-LLVM-Translator tool and update it according to LLVM format changes (e.g. kernel argument information defined in Khronos spec must be named metadata whereas clang emits function metadata).

I am slightly confused as in the LLVM project those address spaces are for SPIR not SPIR-V though.

[skip]

Their conceptual differences are just that they are two different intermediate formats.

If this is the only difference, I don't think it a good idea to create another LLVM target to separate SPIR and SPIR-V. From my point of view it creates logic ambiguity and code duplication with no additional value. @Anastasia, what problems do you see if we continue treating LLVM IR with spir* target triple as LLVM IR representation of SPIR-V format?

[linjamaki]

clang/lib/Basic/Targets/SPIR.h

233

The DefaultIsGeneric AS mapping is enabled for SYCL in the BaseSPIRTargetInfo::adjust (which also means the mapping is available for both the SPIR and SPIR-V targets). On the other hand, the AS mapping for HIPSPV is enabled in SPIRVTargetInfo::adjust only as we intend to emit SPIR-V only. I’m under the impression that this is what was wanted.

[linjamaki]

clang/lib/Basic/Targets/SPIR.h

61

The state of SPIR 1.2/2.0 in Clang seems to be that the SPIR target has transformed to mean “SPIR 1.2/2.0 derivative”, but that does not still make it SPIR-V, which is not based on LLVM IR. When one is targeting spir* there is ambiguity on whether one is aiming to produce the old-SPIR-derivative or SPIR-V. Considering that there are still SPIR-derivative consumers, in my opinion we should have separate LLVM targets for SPIR-V to have explicit disambiguation of intent for producing the SPIR-derivative vs SPIR-V.

[linjamaki]

Rebase.

[Anastasia]

clang/lib/Basic/Targets/SPIR.h

61

@bader, if you would like to migrate SPIR into SPIR-V properly then we should at least rename it. I would certainly prefer triple SPIR-V to SPIR which eliminates the need to explain what it actually is and especially considering that SPIR has existed as an alternative IR format for quite a while. It would at least make sense tpo eliminate the confusion.

However if you would like to go this route you should send a wider community messaging about it and then see if there are any objections. From my experience of previous conversations some years back there are tool developers using SPIR as a portable format even if it's LLVM release dependent however in practice it worked across the latest releases quite well. I would like to remind that not all vendors that support OpenCL or other accelerator API also support SPIR-V. There are also vendors that are migrating to SPIR-V but have older releases in maintenance that don't support SPIR-V. So my feeling is that SPIR has been and is still used as a portable format in tooling.

Regarding an extra triple/target, I don't see a lot of code duplication if we use inheritance/generic programming and other C++ features that will allow us to share the code effectively between both.

[bader]

clang/lib/Basic/Targets/SPIR.h

61

if you would like to migrate SPIR into SPIR-V properly then we should at least rename it.

I have an impression that existing SPIR target should work for both use cases: tools working with "SPIR 1.2/2.0 derivatives" and LLVM -> SPIR-V translation tool(s). I'm trying to clarify why adding mapping for CUDA address spaces works for SPIR-V, but doesn't work for "SPIR 1.2/2.0 derivatives".

[Anastasia]

clang/lib/Basic/Targets/SPIR.h

61

I have an impression that existing SPIR target should work for both use cases: tools working with "SPIR 1.2/2.0 derivatives" and LLVM -> SPIR-V translation tool(s).

Ok, I have two concerns if we take this route:

1. What do we do about documentation and messaging if we use one target for both? I imagine some updates will be needed somewhere to make it clear that SPIR is SPIR-V and SPIR-V is SPIR and that they will evolve the same way if we decide to go this route... Then at least we probably need a new triple for SPIR-V?
2. What happens if we need different behavior for SPIR-V than what SPIR currently has? For example, my impression is that for SPIR-V backend some OpenCL builtins will be represented differently. Btw developers working on SPIR-V backend should probably be included into this discussion...

Overall I feel adding a new target with code reuse from SPIR will probably make things clearer in a long run, but this should probably be discussed elsewhere either in https://reviews.llvm.org/D109144 or as a wider discussion perhaps via a new RFC thread about the best approach of adding SPIR-V target and the future evolution of SPIR. Then we can make sure this can reach the right audience... Then we can collect a list of requirements about different use cases that developers targets and where they are heading with those in the future and define a suitable direction.

[linjamaki]

Rebase.

[linjamaki]

Rebase.

[bader]

clang/lib/Basic/Targets/SPIR.h

61

I have an impression that existing SPIR target should work for both use cases: tools working with "SPIR 1.2/2.0 derivatives" and LLVM -> SPIR-V translation tool(s).

Ok, I have two concerns if we take this route:

This route has been taken starting with LLVM 3.4+ after SPIR switched from LLVM-based format to SPIR-V, so adding another target and deviating LLVM IR format for SPIR-V from "SPIR 1.2/2.0 derivatives" can be disruptive for the tools like SPIR-V translator. How do you see the transition for these tools to LLVM IR for another target?

1. What do we do about documentation and messaging if we use one target for both? I imagine some updates will be needed somewhere to make it clear that SPIR is SPIR-V and SPIR-V is SPIR and that they will evolve the same way if we decide to go this route... Then at least we probably need a new triple for SPIR-V?

I'm not sure if there is a confusion about the difference between LLVM IR for SPIR target and SPIR-V format. As noted above, SPIR target has been used "for both" from the start (i.e. as soon as SPIR-V has been introduced). Additional SPIR-related restrictions/additions for LLVM IR format are not documented anywhere (except a short section in the SPIR-V translator documentation), so it seems to be a good idea to document the format and how to use it (e.g. https://llvm.org/docs/AMDGPUUsage.html).

2. What happens if we need different behavior for SPIR-V than what SPIR currently has? For example, my impression is that for SPIR-V backend some OpenCL builtins will be represented differently. Btw developers working on SPIR-V backend should probably be included into this discussion...

OpenCL defines built-ins representation in high-level language and SPIR-V defines it for the binary format. How built-ins are represented in LLVM IR is not defined, so implementers has freedom to design it. I think SPIR-V backend developers are trying to design it so multiple languages can target SPIR-V format in addition to OpenCL.

Overall I feel adding a new target with code reuse from SPIR will probably make things clearer in a long run, but this should probably be discussed elsewhere either in https://reviews.llvm.org/D109144 or as a wider discussion perhaps via a new RFC thread about the best approach of adding SPIR-V target and the future evolution of SPIR. Then we can make sure this can reach the right audience... Then we can collect a list of requirements about different use cases that developers targets and where they are heading with those in the future and define a suitable direction.

[Anastasia]

clang/lib/Basic/Targets/SPIR.h

61

    I have an impression that existing SPIR target should work for both use cases: tools working with "SPIR 1.2/2.0 derivatives" and LLVM -> SPIR-V translation tool(s).

Ok, I have two concerns if we take this route:

This route has been taken starting with LLVM 3.4+ after SPIR switched from LLVM-based format to SPIR-V, so adding another target and deviating LLVM IR format for SPIR-V from "SPIR 1.2/2.0 derivatives" can be disruptive for the tools like SPIR-V translator. How do you see the transition for these tools to LLVM IR for another target?

My understanding is that the tools were designed with reuse of SPIR target because we haven't been able to add SPIR-V target into LLVM. If we were able to do it earlier I am not sure that it would have been done this way.

At this point I would like to draw attention to the fact that in OpenCL we would like to revise and improve the tooling for SPIR-V in comparison to what they were in SPIR. One example is a redesign of builtin function support. However there are a lot of tools that do rely on SPIR target and changing the design for SPIR would cause ABI changes for them which we would like to avoid. So at least in OpenCL we would need to maintain old SPIR format but also migrate to more optimal SPIR-V tailored tooling support. In general, I see adding SPIR-V target explicitly as an opportunity to reset and optimize tooling architecture...

[bader]

clang/lib/Basic/Targets/SPIR.h

61

At this point I would like to draw attention to the fact that in OpenCL we would like to revise and improve the tooling for SPIR-V in comparison to what they were in SPIR. One example is a redesign of builtin function support. However there are a lot of tools that do rely on SPIR target and changing the design for SPIR would cause ABI changes for them which we would like to avoid. So at least in OpenCL we would need to maintain old SPIR format but also migrate to more optimal SPIR-V tailored tooling support. In general, I see adding SPIR-V target explicitly as an opportunity to reset and optimize tooling architecture...

Does this patch break any ABIs for OpenCL? I think it's specific to HIP/CUDA language and doesn't impact OpenCL compiler. Please, let me know if I get it wrong.

I fully support improving tooling for SPIR-V, but in my opinion some of such improvements should be done separately from the work done by Henry as they require additional discussions.

[Anastasia]

clang/lib/Basic/Targets/SPIR.h

61

I feel that our discussion has diverged from this patch as right now we are discussing how to add SPIR-V target while this patch needs a change of address spaces for SPIR-V.

I fully support improving tooling for SPIR-V, but in my opinion some of such improvements should be done separately from the work done by Henry as they require additional discussions.

I don't think we can or even should add everything in one commit. The improvements we can do however will likely depend on the route that is taken.

[linjamaki]

Rebase.

[linjamaki]

Gentle ping. Is anything needed to be addressed to get this patch accepted?

[Anastasia]

clang/lib/Basic/Targets/SPIR.h

233

I think the issues here is not related to the target but to the flaw in the address space design in clang. So right now all languages that don't derive the address space semantic from embedded C (SYCL/CUDA/HIP) would need to reset the address space map. See FIXME comment in the definition of adjust.

So the right thing to do would be to set the address space map correctly straight away based on the language being compiled for which would avoid overriding this in adjust. But if we do override it then it makes more sense to at least unify the logic among targets.

On the other hand, the AS mapping for HIPSPV is enabled in SPIRVTargetInfo::adjust only as we intend to emit SPIR-V only.

I am not really sure how you would support one target only.
Clang architecture (at least originally) assumes that all languages can map to all targets which in practice is not true in some cases. This is why we need to provide an address space map even for targets that have no memory segmented language compiled to it.

[linjamaki]

Rebase, add asserts and move address space map reset for HIP from SPIRVTargetInfo to BaseSPIRTargetInfo

[linjamaki]

clang/lib/Basic/Targets/SPIR.h

233

So the right thing to do would be to set the address space map correctly straight away based on the language being compiled for which would avoid overriding this in adjust. But if we do override it then it makes more sense to at least unify the logic among targets.

Since we are not sure how we would solve this issue optimally, we adjusted the patch to avoid adding more overrides for the adjust method and the logic previously in the SPIRVTargetInfo::adjust is moved to BaseSPIRTargetInfo::adjust with the SYCL. Would this be sufficient for the functionality added by this patch?

I am not really sure how you would support one target only.
Clang architecture (at least originally) assumes that all languages can map to all targets which in practice is not true in some cases. This is why we need to provide an address space map even for targets that have no memory segmented language compiled to it.

“HIPSPV” is not meant to be a new language. We are just adjusting the address space mapping from the HIP language (for device code) to SPIR-V that suits better than the default mapping where all the HIP address spaces would be mapped to target address space zero. We map the address spaces to the suitable ones in the OpenCL standard, which works both for HIPCL (which uses the OpenCL-based runtime and the OpenCL SPIR-V profile) and HIPLZ (which uses the LZ-based runtime and also the OpenCL SPIR-V profile).

[Anastasia]

LGTM! Thanks

[linjamaki]

@Anastasia, Could you please commit this patch to the LLVM for us? Thanks.

[linjamaki]

Thanks, @Anastasia.

clang/lib/Basic/Targets/SPIR.h

static const unsigned SPIRDefIsGenMap[] = {
     // OpenCL address space values for this map are dummy and they can't be used
     0, // opencl_global
     0, // opencl_local
     0, // opencl_constant
     0, // opencl_private
     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
     3, // sycl_local
     0, // sycl_private
     0, // ptr32_sptr
     0, // ptr32_uptr
     0  // ptr64
public:
 
   void setAddressSpaceMap(bool DefaultIsGeneric) {
     AddrSpaceMap = DefaultIsGeneric ? &SPIRDefIsGenMap : &SPIRDefIsPrivMap;
   }
 
   void adjust(DiagnosticsEngine &Diags, LangOptions &Opts) override {
     TargetInfo::adjust(Diags, Opts);
     // 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);
   }
 
public:
   }
 
   void getTargetDefines(const LangOptions &Opts,
                         MacroBuilder &Builder) const override;
 
   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 {
 public:
   SPIRV32TargetInfo(const llvm::Triple &Triple, const TargetOptions &Opts)
       : SPIRVTargetInfo(Triple, Opts) {
     PointerWidth = PointerAlign = 32;
     SizeType = TargetInfo::UnsignedInt;

clang/test/CodeGenHIP/hipspv-addr-spaces.cpp

+// 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;
+}