[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

216

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

216

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

216

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

216

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/include/clang/Basic/DiagnosticGroups.td

 
 def InvalidOrNonExistentDirectory : DiagGroup<"invalid-or-nonexistent-directory">;
 
 def OptionIgnored : DiagGroup<"option-ignored">;
 
 def UnknownArgument : DiagGroup<"unknown-argument">;
 
 // A warning group for warnings about code that clang accepts when
-// compiling OpenCL C/C++ but which is not compatible with the SPIR spec.
+// compiling OpenCL C/C++ but which is not compatible with the SPIR(-V) spec.
 def SpirCompat : DiagGroup<"spir-compat">;
+def : DiagGroup<"spirv-compat", [SpirCompat]>; // Alias.
 
 // Warning for the GlobalISel options.
 def GlobalISel : DiagGroup<"global-isel">;
 
 // A warning group specifically for warnings related to function
 // multiversioning.
 def FunctionMultiVersioning : DiagGroup<"function-multiversion">;
 

clang/lib/Basic/Targets.cpp

case llvm::Triple::spir: {
     return new SPIR32TargetInfo(Triple, Opts);
   }
   case llvm::Triple::spir64: {
     if (os != llvm::Triple::UnknownOS ||
         Triple.getEnvironment() != llvm::Triple::UnknownEnvironment)
       return nullptr;
     return new SPIR64TargetInfo(Triple, Opts);
   }
+  case llvm::Triple::spirv32: {
+    if (os != llvm::Triple::UnknownOS ||
+        Triple.getEnvironment() != llvm::Triple::UnknownEnvironment)
+      return nullptr;
+    return new SPIRV32TargetInfo(Triple, Opts);
+  }
+  case llvm::Triple::spirv64: {
+    if (os != llvm::Triple::UnknownOS ||
+        Triple.getEnvironment() != llvm::Triple::UnknownEnvironment)
+      return nullptr;
+    return new SPIRV64TargetInfo(Triple, Opts);
+  }
   case llvm::Triple::wasm32:
     if (Triple.getSubArch() != llvm::Triple::NoSubArch ||
         Triple.getVendor() != llvm::Triple::UnknownVendor ||
         !Triple.isOSBinFormatWasm())
       return nullptr;
     switch (os) {
       case llvm::Triple::WASI:
         return new WASITargetInfo<WebAssembly32TargetInfo>(Triple, Opts);

clang/lib/Basic/Targets/SPIR.h

-//===--- SPIR.h - Declare SPIR target feature support -----------*- C++ -*-===//
+//===--- SPIR.h - Declare SPIR and SPIR-V target feature support *- C++ -*-===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 //
-// This file declares SPIR TargetInfo objects.
+// This file declares SPIR and SPIR-V TargetInfo objects.
 //
 //===----------------------------------------------------------------------===//
 
 #ifndef LLVM_CLANG_LIB_BASIC_TARGETS_SPIR_H
 #define LLVM_CLANG_LIB_BASIC_TARGETS_SPIR_H
 
 #include "clang/Basic/TargetInfo.h"
 #include "clang/Basic/TargetOptions.h"
 #include "llvm/ADT/Triple.h"
 #include "llvm/Support/Compiler.h"
 
 namespace clang {
 namespace targets {
 
+// Used by both the SPIR and SPIR-V targets.
 static const unsigned SPIRDefIsPrivMap[] = {
     0, // Default
     1, // opencl_global
     3, // opencl_local
     2, // opencl_constant
     0, // opencl_private
     4, // opencl_generic
     5, // opencl_global_device
     6, // opencl_global_host
     0, // cuda_device
     0, // cuda_constant
     0, // cuda_shared
     // SYCL address space values for this map are dummy
     0, // sycl_global
     0, // sycl_global_device
     0, // sycl_global_host
     0, // sycl_local
     0, // sycl_private
     0, // ptr32_sptr
     0, // ptr32_uptr
     0  // ptr64
 };
 
+// Used by both the SPIR and SPIR-V targets.
 static const unsigned SPIRDefIsGenMap[] = {
     4, // Default
     // 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
 };
 
-class LLVM_LIBRARY_VISIBILITY SPIRTargetInfo : public TargetInfo {
-public:
-  SPIRTargetInfo(const llvm::Triple &Triple, const TargetOptions &)
+// Base class for SPIR and SPIR-V target info.
+class LLVM_LIBRARY_VISIBILITY BaseSPIRTargetInfo : public TargetInfo {
+protected:
+  BaseSPIRTargetInfo(const llvm::Triple &Triple, const TargetOptions &)
       : TargetInfo(Triple) {
     assert(getTriple().getOS() == llvm::Triple::UnknownOS &&
-           "SPIR target must use unknown OS");
+           "SPIR(-V) target must use unknown OS");
     assert(getTriple().getEnvironment() == llvm::Triple::UnknownEnvironment &&
-           "SPIR target must use unknown environment type");
+           "SPIR(-V) target must use unknown environment type");
     TLSSupported = false;
     VLASupported = false;
     LongWidth = LongAlign = 64;
     AddrSpaceMap = &SPIRDefIsPrivMap;
     UseAddrSpaceMapMangling = true;
     HasLegalHalfType = true;
     HasFloat16 = true;
     // Define available target features
     // These must be defined in sorted order!
     NoAsmVariants = true;
   }
 
-  void getTargetDefines(const LangOptions &Opts,
-                        MacroBuilder &Builder) const override;
-
-  bool hasFeature(StringRef Feature) const override {
-    return Feature == "spir";
-  }
-
+public:
   // SPIR supports the half type and the only llvm intrinsic allowed in SPIR is
   // memcpy as per section 3 of the SPIR spec.
   bool useFP16ConversionIntrinsics() const override { return false; }
 
   ArrayRef<Builtin::Info> getTargetBuiltins() const override { return None; }
 
   const char *getClobbers() const override { return ""; }
 
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);
   }
 
   void setSupportedOpenCLOpts() override {
     // Assume all OpenCL extensions and optional core features are supported
-    // for SPIR since it is a generic target.
+    // for SPIR and SPIR-V since they are generic targets.
     supportAllOpenCLOpts();
   }
 
   bool hasExtIntType() const override { return true; }
 
   bool hasInt128Type() const override { return false; }
 };
 
+class LLVM_LIBRARY_VISIBILITY SPIRTargetInfo : public BaseSPIRTargetInfo {
+public:
+  SPIRTargetInfo(const llvm::Triple &Triple, const TargetOptions &Opts)
+      : BaseSPIRTargetInfo(Triple, Opts) {
+    assert(getTriple().getOS() == llvm::Triple::UnknownOS &&
+           "SPIR target must use unknown OS");
+    assert(getTriple().getEnvironment() == llvm::Triple::UnknownEnvironment &&
+           "SPIR target must use unknown environment type");
+  }
+
+  void getTargetDefines(const LangOptions &Opts,
+                        MacroBuilder &Builder) const override;
+
+  bool hasFeature(StringRef Feature) const override {
+    return Feature == "spir";
+  }
+};
+
 class LLVM_LIBRARY_VISIBILITY SPIR32TargetInfo : public SPIRTargetInfo {
 public:
   SPIR32TargetInfo(const llvm::Triple &Triple, const TargetOptions &Opts)
       : SPIRTargetInfo(Triple, Opts) {
     PointerWidth = PointerAlign = 32;
     SizeType = TargetInfo::UnsignedInt;
     PtrDiffType = IntPtrType = TargetInfo::SignedInt;
     resetDataLayout("e-p:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-"
SPIR64TargetInfo(const llvm::Triple &Triple, const TargetOptions &Opts)
     PtrDiffType = IntPtrType = TargetInfo::SignedLong;
     resetDataLayout("e-i64:64-v16:16-v24:32-v32:32-v48:64-"
                     "v96:128-v192:256-v256:256-v512:512-v1024:1024");
   }
 
   void getTargetDefines(const LangOptions &Opts,
                         MacroBuilder &Builder) const override;
 };
+
+class LLVM_LIBRARY_VISIBILITY SPIRVTargetInfo : public BaseSPIRTargetInfo {
+public:
+  SPIRVTargetInfo(const llvm::Triple &Triple, const TargetOptions &Opts)
+      : BaseSPIRTargetInfo(Triple, Opts) {
+    assert(getTriple().getOS() == llvm::Triple::UnknownOS &&
+           "SPIR-V target must use unknown OS");
+    assert(getTriple().getEnvironment() == llvm::Triple::UnknownEnvironment &&
+           "SPIR-V target must use unknown environment type");
+  }
+
+  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;
+    PtrDiffType = IntPtrType = TargetInfo::SignedInt;
+    resetDataLayout("e-p:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-"
+                    "v96:128-v192:256-v256:256-v512:512-v1024:1024");
+  }
+
+  void getTargetDefines(const LangOptions &Opts,
+                        MacroBuilder &Builder) const override;
+};
+
+class LLVM_LIBRARY_VISIBILITY SPIRV64TargetInfo : public SPIRVTargetInfo {
+public:
+  SPIRV64TargetInfo(const llvm::Triple &Triple, const TargetOptions &Opts)
+      : SPIRVTargetInfo(Triple, Opts) {
+    PointerWidth = PointerAlign = 64;
+    SizeType = TargetInfo::UnsignedLong;
+    PtrDiffType = IntPtrType = TargetInfo::SignedLong;
+    resetDataLayout("e-i64:64-v16:16-v24:32-v32:32-v48:64-"
+                    "v96:128-v192:256-v256:256-v512:512-v1024:1024");
+  }
+
+  void getTargetDefines(const LangOptions &Opts,
+                        MacroBuilder &Builder) const override;
+};
+
 } // namespace targets
 } // namespace clang
 #endif // LLVM_CLANG_LIB_BASIC_TARGETS_SPIR_H

clang/lib/Basic/Targets/SPIR.cpp

-//===--- SPIR.cpp - Implement SPIR target feature support -----------------===//
+//===--- SPIR.cpp - Implement SPIR and SPIR-V target feature support ------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 //
-// This file implements SPIR TargetInfo objects.
+// This file implements SPIR and SPIR-V TargetInfo objects.
 //
 //===----------------------------------------------------------------------===//
 
 #include "SPIR.h"
 #include "Targets.h"
 
 using namespace clang;
 using namespace clang::targets;
void SPIR32TargetInfo::getTargetDefines(const LangOptions &Opts,
   DefineStd(Builder, "SPIR32", Opts);
 }
 
 void SPIR64TargetInfo::getTargetDefines(const LangOptions &Opts,
                                         MacroBuilder &Builder) const {
   SPIRTargetInfo::getTargetDefines(Opts, Builder);
   DefineStd(Builder, "SPIR64", Opts);
 }
+
+void SPIRVTargetInfo::getTargetDefines(const LangOptions &Opts,
+                                       MacroBuilder &Builder) const {
+  DefineStd(Builder, "SPIRV", Opts);
+}
+
+void SPIRV32TargetInfo::getTargetDefines(const LangOptions &Opts,
+                                         MacroBuilder &Builder) const {
+  SPIRVTargetInfo::getTargetDefines(Opts, Builder);
+  DefineStd(Builder, "SPIRV32", Opts);
+}
+
+void SPIRV64TargetInfo::getTargetDefines(const LangOptions &Opts,
+                                         MacroBuilder &Builder) const {
+  SPIRVTargetInfo::getTargetDefines(Opts, Builder);
+  DefineStd(Builder, "SPIRV64", Opts);
+}

clang/lib/CodeGen/TargetInfo.cpp

for (unsigned I = 0; I != MangledDeclNames.size(); ++I) {
     auto Val = *(MangledDeclNames.begin() + I);
     llvm::GlobalValue *GV = CGM.GetGlobalValue(Val.second);
     if (GV) {
       const Decl *D = Val.first.getDecl()->getMostRecentDecl();
       emitTargetMD(D, GV, CGM);
     }
   }
 }
+
 //===----------------------------------------------------------------------===//
-// SPIR ABI Implementation
+// Base ABI and target codegen info implementation common between SPIR and
+// SPIR-V.
 //===----------------------------------------------------------------------===//
 
 namespace {
-class SPIRABIInfo : public DefaultABIInfo {
+class CommonSPIRABIInfo : public DefaultABIInfo {
 public:
-  SPIRABIInfo(CodeGenTypes &CGT) : DefaultABIInfo(CGT) { setCCs(); }
+  CommonSPIRABIInfo(CodeGenTypes &CGT) : DefaultABIInfo(CGT) { setCCs(); }
 
 private:
   void setCCs();
 };
 } // end anonymous namespace
 namespace {
-class SPIRTargetCodeGenInfo : public TargetCodeGenInfo {
+class CommonSPIRTargetCodeGenInfo : public TargetCodeGenInfo {
 public:
-  SPIRTargetCodeGenInfo(CodeGen::CodeGenTypes &CGT)
-      : TargetCodeGenInfo(std::make_unique<SPIRABIInfo>(CGT)) {}
+  CommonSPIRTargetCodeGenInfo(CodeGen::CodeGenTypes &CGT)
+      : TargetCodeGenInfo(std::make_unique<CommonSPIRABIInfo>(CGT)) {}
 
   LangAS getASTAllocaAddressSpace() const override {
     return getLangASFromTargetAS(
         getABIInfo().getDataLayout().getAllocaAddrSpace());
   }
 
   unsigned getOpenCLKernelCallingConv() const override;
 };
 
 } // End anonymous namespace.
-void SPIRABIInfo::setCCs() {
+void CommonSPIRABIInfo::setCCs() {
   assert(getRuntimeCC() == llvm::CallingConv::C);
   RuntimeCC = llvm::CallingConv::SPIR_FUNC;
 }
 
 namespace clang {
 namespace CodeGen {
 void computeSPIRKernelABIInfo(CodeGenModule &CGM, CGFunctionInfo &FI) {
   DefaultABIInfo SPIRABI(CGM.getTypes());
   SPIRABI.computeInfo(FI);
 }
 }
 }
 
-unsigned SPIRTargetCodeGenInfo::getOpenCLKernelCallingConv() const {
+unsigned CommonSPIRTargetCodeGenInfo::getOpenCLKernelCallingConv() const {
   return llvm::CallingConv::SPIR_KERNEL;
 }
 
 static bool appendType(SmallStringEnc &Enc, QualType QType,
                        const CodeGen::CodeGenModule &CGM,
                        TypeStringCache &TSC);
 
 /// Helper function for appendRecordType().
const TargetCodeGenInfo &CodeGenModule::getTargetCodeGenInfo() {
   case llvm::Triple::sparcv9:
     return SetCGInfo(new SparcV9TargetCodeGenInfo(Types));
   case llvm::Triple::xcore:
     return SetCGInfo(new XCoreTargetCodeGenInfo(Types));
   case llvm::Triple::arc:
     return SetCGInfo(new ARCTargetCodeGenInfo(Types));
   case llvm::Triple::spir:
   case llvm::Triple::spir64:
-    return SetCGInfo(new SPIRTargetCodeGenInfo(Types));
+  case llvm::Triple::spirv32:
+  case llvm::Triple::spirv64:
+    return SetCGInfo(new CommonSPIRTargetCodeGenInfo(Types));
   case llvm::Triple::ve:
     return SetCGInfo(new VETargetCodeGenInfo(Types));
   }
 }
 
 /// Create an OpenCL kernel for an enqueued block.
 ///
 /// The kernel has the same function type as the block invoke function. Its

clang/lib/Driver/ToolChains/Clang.cpp

void Clang::ConstructJob(Compilation &C, const JobAction &JA,
   }
 
   // Push all default warning arguments that are specific to
   // the given target.  These come before user provided warning options
   // are provided.
   TC.addClangWarningOptions(CmdArgs);
 
   // FIXME: Subclass ToolChain for SPIR and move this to addClangWarningOptions.
-  if (Triple.isSPIR())
+  if (Triple.isSPIR() || Triple.isSPIRV())
     CmdArgs.push_back("-Wspir-compat");
 
   // Select the appropriate action.
   RewriteKind rewriteKind = RK_None;
 
   // If CollectArgsForIntegratedAssembler() isn't called below, claim the args
   // it claims when not running an assembler. Otherwise, clang would emit
   // "argument unused" warnings for assembler flags when e.g. adding "-E" to

clang/lib/Frontend/InitPreprocessor.cpp

#undef DEFINE_LOCK_FREE_MACRO
   if (LangOpts.SYCLIsDevice) {
     Builder.defineMacro("__SYCL_DEVICE_ONLY__", "1");
   }
 
   // OpenCL definitions.
   if (LangOpts.OpenCL) {
     InitializeOpenCLFeatureTestMacros(TI, LangOpts, Builder);
 
-    if (TI.getTriple().isSPIR())
+    if (TI.getTriple().isSPIR() || TI.getTriple().isSPIRV())
       Builder.defineMacro("__IMAGE_SUPPORT__");
   }
 
   if (TI.hasInt128Type() && LangOpts.CPlusPlus && LangOpts.GNUMode) {
     // For each extended integer type, g++ defines a macro mapping the
     // index of the type (0 in this case) in some list of extended types
     // to the type.
     Builder.defineMacro("__GLIBCXX_TYPE_INT_N_0", "__int128");

clang/lib/Headers/opencl-c-base.h

 //===----- opencl-c-base.h - OpenCL C language base definitions -----------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
 #ifndef _OPENCL_BASE_H_
 #define _OPENCL_BASE_H_
 
 // Define extension macros
 
 #if (defined(__OPENCL_CPP_VERSION__) || __OPENCL_C_VERSION__ >= 200)
-// For SPIR all extensions are supported.
-#if defined(__SPIR__)
+// For SPIR and SPIR-V all extensions are supported.
+#if defined(__SPIR__) || defined(__SPIRV__)
 #define cl_khr_subgroup_extended_types 1
 #define cl_khr_subgroup_non_uniform_vote 1
 #define cl_khr_subgroup_ballot 1
 #define cl_khr_subgroup_non_uniform_arithmetic 1
 #define cl_khr_subgroup_shuffle 1
 #define cl_khr_subgroup_shuffle_relative 1
 #define cl_khr_subgroup_clustered_reduce 1
 #define cl_khr_extended_bit_ops 1

 #define __opencl_c_ext_fp64_global_atomic_min_max 1
 #define __opencl_c_ext_fp64_local_atomic_min_max 1
 #endif
 #define __opencl_c_ext_fp32_global_atomic_add 1
 #define __opencl_c_ext_fp32_local_atomic_add 1
 #define __opencl_c_ext_fp32_global_atomic_min_max 1
 #define __opencl_c_ext_fp32_local_atomic_min_max 1
 
-#endif // defined(__SPIR__)
+#endif // defined(__SPIR__) || defined(__SPIRV__)
 #endif // (defined(__OPENCL_CPP_VERSION__) || __OPENCL_C_VERSION__ >= 200)
 
 // Define feature macros for OpenCL C 2.0
 #if (__OPENCL_CPP_VERSION__ == 100 || __OPENCL_C_VERSION__ == 200)
 #define __opencl_c_pipes 1
 #define __opencl_c_generic_address_space 1
 #define __opencl_c_work_group_collective_functions 1
 #define __opencl_c_atomic_order_acq_rel 1
 #define __opencl_c_atomic_order_seq_cst 1
 #define __opencl_c_atomic_scope_device 1
 #define __opencl_c_atomic_scope_all_devices 1
 #define __opencl_c_device_enqueue 1
 #define __opencl_c_read_write_images 1
 #define __opencl_c_program_scope_global_variables 1
 #define __opencl_c_images 1
 #endif
 
 // Define header-only feature macros for OpenCL C 3.0.
 #if (__OPENCL_CPP_VERSION__ == 202100 || __OPENCL_C_VERSION__ == 300)
-// For the SPIR target all features are supported.
-#if defined(__SPIR__)
+// For the SPIR and SPIR-V target all features are supported.
+#if defined(__SPIR__) || defined(__SPIRV__)
 #define __opencl_c_atomic_scope_all_devices 1
 #endif // defined(__SPIR__)
 #endif // (__OPENCL_CPP_VERSION__ == 202100 || __OPENCL_C_VERSION__ == 300)
 
 // built-in scalar data types:
 
 /**
  * An unsigned 8-bit integer.

clang/lib/Headers/opencl-c.h

 #endif //defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 
 #if __OPENCL_C_VERSION__ < CL_VERSION_2_0
 #ifdef cl_khr_3d_image_writes
 #pragma OPENCL EXTENSION cl_khr_3d_image_writes : enable
 #endif //cl_khr_3d_image_writes
 #endif //__OPENCL_C_VERSION__ < CL_VERSION_2_0
 
-
-#if (defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_1_2)) && defined(__SPIR__)
+#if (defined(__OPENCL_CPP_VERSION__) ||                                        \
+     (__OPENCL_C_VERSION__ >= CL_VERSION_1_2)) &&                              \
+    (defined(__SPIR__) || defined(__SPIRV__))
 #pragma OPENCL EXTENSION cl_intel_planar_yuv : begin
 #pragma OPENCL EXTENSION cl_intel_planar_yuv : end
-#endif // (defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_1_2)) && defined(__SPIR__)
+#endif // (defined(__OPENCL_CPP_VERSION__) ||
+       //  (__OPENCL_C_VERSION__ >= CL_VERSION_1_2)) &&
+       // (defined(__SPIR__) || defined(__SPIRV__))
 
 #define __ovld __attribute__((overloadable))
 #define __conv __attribute__((convergent))
 
 // Optimizations
 #define __purefn __attribute__((pure))
 #define __cnfn __attribute__((const))
 

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

clang/test/CodeGenOpenCL/spirv_target.cl

+// RUN: %clang_cc1 %s -triple "spirv32-unknown-unknown" -verify -emit-llvm -o - | FileCheck %s -check-prefix=SPIRV32
+// RUN: %clang_cc1 %s -triple "spirv64-unknown-unknown" -verify -emit-llvm -o - | FileCheck %s -check-prefix=SPIRV64
+
+// SPIRV32: target triple = "spirv32-unknown-unknown"
+// SPIRV64: target triple = "spirv64-unknown-unknown"
+
+typedef struct {
+  char c;
+  void *v;
+  void *v2;
+} my_st;
+
+kernel void foo(global long *arg) {
+#if __SPIRV32__ == 1
+  int res1[sizeof(my_st)  == 12 ? 1 : -1]; // expected-no-diagnostics
+  int res2[sizeof(void *) ==  4 ? 1 : -1]; // expected-no-diagnostics
+  int res3[sizeof(arg)    ==  4 ? 1 : -1]; // expected-no-diagnostics
+#elif __SPIRV64__ == 1
+  int res1[sizeof(my_st)  == 24 ? 1 : -1]; // expected-no-diagnostics
+  int res2[sizeof(void *) ==  8 ? 1 : -1]; // expected-no-diagnostics
+  int res3[sizeof(arg)    ==  8 ? 1 : -1]; // expected-no-diagnostics
+#endif
+  my_st *tmp = 0;
+
+  // SPIRV32: store i64 4, i64 addrspace(1)*
+  // SPIRV64: store i64 8, i64 addrspace(1)*
+  arg[0] = (long)(&tmp->v);
+  // SPIRV32: store i64 8, i64 addrspace(1)*
+  // SPIRV64: store i64 16, i64 addrspace(1)*
+  arg[1] = (long)(&tmp->v2);
+}

clang/test/Headers/opencl-c-header.cl

 // RUN: %clang_cc1 -O0 -triple spir-unknown-unknown -internal-isystem ../../lib/Headers -include opencl-c.h -emit-llvm -o - %s -verify | FileCheck %s
 // RUN: %clang_cc1 -O0 -triple spir-unknown-unknown -internal-isystem ../../lib/Headers -include opencl-c.h -emit-llvm -o - %s -verify -cl-std=CL1.1 | FileCheck %s
 // RUN: %clang_cc1 -O0 -triple spir-unknown-unknown -internal-isystem ../../lib/Headers -include opencl-c.h -emit-llvm -o - %s -verify -cl-std=CL1.2 | FileCheck %s
 // RUN: %clang_cc1 -O0 -triple spir-unknown-unknown -internal-isystem ../../lib/Headers -include opencl-c.h -emit-llvm -o - %s -verify -cl-std=clc++ | FileCheck %s --check-prefix=CHECK20
 // RUN: %clang_cc1 -O0 -triple spir-unknown-unknown -internal-isystem ../../lib/Headers -include opencl-c.h -emit-llvm -o - %s -verify -cl-std=CL3.0 | FileCheck %s
 
+// RUN: %clang_cc1 -O0 -triple spirv32-unknown-unknown -internal-isystem ../../lib/Headers -include opencl-c.h -emit-llvm -o - %s -verify | FileCheck %s
+
+
 // Test including the default header as a module.
 // The module should be compiled only once and loaded from cache afterwards.
 // Change the directory mode to read only to make sure no new modules are created.
 // Check time report to make sure module is used.
 // Check that some builtins occur in the generated IR when called.
 
 // ===
 // Clear current directory.

 #if defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ == CL_VERSION_2_0)
 global atomic_int z = ATOMIC_VAR_INIT(99);
 #endif //__OPENCL_C_VERSION__
 // CHECK-MOD: Reading modules
 
 // Check that extension macros are defined correctly.
 
 // For SPIR all extensions are supported.
-#if defined(__SPIR__)
+#if defined(__SPIR__) || defined(__SPIRV__)
 
 // Verify that cl_intel_planar_yuv extension is defined from OpenCL 1.2 onwards.
 #if defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_1_2)
 // expected-no-diagnostics
 #else //__OPENCL_C_VERSION__
 // expected-warning@+2{{unknown OpenCL extension 'cl_intel_planar_yuv' - ignoring}}
 #endif //__OPENCL_C_VERSION__
 #pragma OPENCL EXTENSION cl_intel_planar_yuv : enable

clang/test/Preprocessor/predefined-macros.c

 
 // RUN: %clang_cc1 %s -E -dM -o - -x cl -triple spir64-unknown-unknown \
 // RUN:   | FileCheck -match-full-lines %s --check-prefix=CHECK-SPIR64
 // CHECK-SPIR64-DAG: #define __IMAGE_SUPPORT__ 1
 // CHECK-SPIR64-DAG: #define __SPIR__ 1
 // CHECK-SPIR64-DAG: #define __SPIR64__ 1
 // CHECK-SPIR64-NOT: #define __SPIR32__ 1
 
-// RUN: %clang_cc1 %s -E -dM -o - -x hip -triple amdgcn-amd-amdhsa \
+// RUN: %clang_cc1 %s -E -dM -o - -x cl -triple spirv32-unknown-unknown \
+// RUN:   | FileCheck -match-full-lines %s --check-prefix=CHECK-SPIRV32
+// CHECK-SPIRV32-DAG: #define __IMAGE_SUPPORT__ 1
+// CHECK-SPIRV32-DAG: #define __SPIRV__ 1
+// CHECK-SPIRV32-DAG: #define __SPIRV32__ 1
+// CHECK-SPIRV32-NOT: #define __SPIRV64__ 1
+
+// RUN: %clang_cc1 %s -E -dM -o - -x cl -triple spirv64-unknown-unknown \
+// RUN:   | FileCheck -match-full-lines %s --check-prefix=CHECK-SPIRV64
+// CHECK-SPIRV64-DAG: #define __IMAGE_SUPPORT__ 1
+// CHECK-SPIRV64-DAG: #define __SPIRV__ 1
+// CHECK-SPIRV64-DAG: #define __SPIRV64__ 1
+// CHECK-SPIRV64-NOT: #define __SPIRV32__ 1
+
+// RUN: %clang_cc1 %s -E -dM -o - -x hip -triple amdgcn-amd-amdhsa      \
 // RUN:   | FileCheck -match-full-lines %s --check-prefix=CHECK-HIP
 // CHECK-HIP-NOT: #define __CUDA_ARCH__
 // CHECK-HIP: #define __HIPCC__ 1
 // CHECK-HIP-NOT: #define __HIP_DEVICE_COMPILE__ 1
 // CHECK-HIP: #define __HIP__ 1
 
 // RUN: %clang_cc1 %s -E -dM -o - -x hip -triple amdgcn-amd-amdhsa \
 // RUN:   -fcuda-is-device \
 // RUN:   | FileCheck -match-full-lines %s --check-prefix=CHECK-HIP-DEV
 // CHECK-HIP-DEV-NOT: #define __CUDA_ARCH__
 // CHECK-HIP-DEV: #define __HIPCC__ 1
 // CHECK-HIP-DEV: #define __HIP_DEVICE_COMPILE__ 1
 // CHECK-HIP-DEV: #define __HIP__ 1

llvm/include/llvm/ADT/Triple.h

enum ArchType {
     le32,           // le32: generic little-endian 32-bit CPU (PNaCl)
     le64,           // le64: generic little-endian 64-bit CPU (PNaCl)
     amdil,          // AMDIL
     amdil64,        // AMDIL with 64-bit pointers
     hsail,          // AMD HSAIL
     hsail64,        // AMD HSAIL with 64-bit pointers
     spir,           // SPIR: standard portable IR for OpenCL 32-bit version
     spir64,         // SPIR: standard portable IR for OpenCL 64-bit version
+    spirv32,        // SPIR-V with 32-bit pointers
+    spirv64,        // SPIR-V with 64-bit pointers
     kalimba,        // Kalimba: generic kalimba
     shave,          // SHAVE: Movidius vector VLIW processors
     lanai,          // Lanai: Lanai 32-bit
     wasm32,         // WebAssembly with 32-bit pointers
     wasm64,         // WebAssembly with 64-bit pointers
     renderscript32, // 32-bit RenderScript
     renderscript64, // 64-bit RenderScript
     ve,             // NEC SX-Aurora Vector Engine
return getEnvironment() == Triple::Musl ||
            getEnvironment() == Triple::MuslX32;
   }
 
   /// Tests whether the target is SPIR (32- or 64-bit).
   bool isSPIR() const {
     return getArch() == Triple::spir || getArch() == Triple::spir64;
   }
 
+  /// Tests whether the target is SPIR-V (32/64-bit).
+  bool isSPIRV() const {
+    return getArch() == Triple::spirv32 || getArch() == Triple::spirv64;
+  }
+
   /// Tests whether the target is NVPTX (32- or 64-bit).
   bool isNVPTX() const {
     return getArch() == Triple::nvptx || getArch() == Triple::nvptx64;
   }
 
   /// Tests whether the target is AMDGCN
   bool isAMDGCN() const { return getArch() == Triple::amdgcn; }
 

llvm/lib/Support/Triple.cpp

StringRef Triple::getArchTypeName(ArchType Kind) {
   case riscv32:        return "riscv32";
   case riscv64:        return "riscv64";
   case shave:          return "shave";
   case sparc:          return "sparc";
   case sparcel:        return "sparcel";
   case sparcv9:        return "sparcv9";
   case spir64:         return "spir64";
   case spir:           return "spir";
+  case spirv32:        return "spirv32";
+  case spirv64:        return "spirv64";
   case systemz:        return "s390x";
   case tce:            return "tce";
   case tcele:          return "tcele";
   case thumb:          return "thumb";
   case thumbeb:        return "thumbeb";
   case ve:             return "ve";
   case wasm32:         return "wasm32";
   case wasm64:         return "wasm64";
StringRef Triple::getArchTypePrefix(ArchType Kind) {
   case amdil:
   case amdil64:     return "amdil";
 
   case hsail:
   case hsail64:     return "hsail";
 
   case spir:
   case spir64:      return "spir";
+
+  case spirv32:
+  case spirv64:     return "spirv";
+
   case kalimba:     return "kalimba";
   case lanai:       return "lanai";
   case shave:       return "shave";
   case wasm32:
   case wasm64:      return "wasm";
 
   case riscv32:
   case riscv64:     return "riscv";
return StringSwitch<Triple::ArchType>(Name)
     .Case("le32", le32)
     .Case("le64", le64)
     .Case("amdil", amdil)
     .Case("amdil64", amdil64)
     .Case("hsail", hsail)
     .Case("hsail64", hsail64)
     .Case("spir", spir)
     .Case("spir64", spir64)
+    .Case("spirv32", spirv32)
+    .Case("spirv64", spirv64)
     .Case("kalimba", kalimba)
     .Case("lanai", lanai)
     .Case("shave", shave)
     .Case("wasm32", wasm32)
     .Case("wasm64", wasm64)
     .Case("renderscript32", renderscript32)
     .Case("renderscript64", renderscript64)
     .Case("ve", ve)
auto AT = StringSwitch<Triple::ArchType>(ArchName)
     .Case("le32", Triple::le32)
     .Case("le64", Triple::le64)
     .Case("amdil", Triple::amdil)
     .Case("amdil64", Triple::amdil64)
     .Case("hsail", Triple::hsail)
     .Case("hsail64", Triple::hsail64)
     .Case("spir", Triple::spir)
     .Case("spir64", Triple::spir64)
+    .Case("spirv32", Triple::spirv32)
+    .Case("spirv64", Triple::spirv64)
     .StartsWith("kalimba", Triple::kalimba)
     .Case("lanai", Triple::lanai)
     .Case("renderscript32", Triple::renderscript32)
     .Case("renderscript64", Triple::renderscript64)
     .Case("shave", Triple::shave)
     .Case("ve", Triple::ve)
     .Case("wasm32", Triple::wasm32)
     .Case("wasm64", Triple::wasm64)
static Triple::ObjectFormatType getDefaultFormat(const Triple &T) {
   case Triple::systemz:
     if (T.isOSzOS())
       return Triple::GOFF;
     return Triple::ELF;
 
   case Triple::wasm32:
   case Triple::wasm64:
     return Triple::Wasm;
+
+  case Triple::spirv32:
+  case Triple::spirv64:
+    // TODO: In future this will be Triple::SPIRV.
+    return Triple::UnknownObjectFormat;
   }
   llvm_unreachable("unknown architecture");
 }
 
 /// Construct a triple from the string representation provided.
 ///
 /// This stores the string representation and parses the various pieces into
 /// enum members.
static unsigned getArchPointerBitWidth(llvm::Triple::ArchType Arch) {
   case llvm::Triple::ppcle:
   case llvm::Triple::r600:
   case llvm::Triple::renderscript32:
   case llvm::Triple::riscv32:
   case llvm::Triple::shave:
   case llvm::Triple::sparc:
   case llvm::Triple::sparcel:
   case llvm::Triple::spir:
+  case llvm::Triple::spirv32:
   case llvm::Triple::tce:
   case llvm::Triple::tcele:
   case llvm::Triple::thumb:
   case llvm::Triple::thumbeb:
   case llvm::Triple::wasm32:
   case llvm::Triple::x86:
   case llvm::Triple::xcore:
     return 32;
static unsigned getArchPointerBitWidth(llvm::Triple::ArchType Arch) {
   case llvm::Triple::mips64el:
   case llvm::Triple::nvptx64:
   case llvm::Triple::ppc64:
   case llvm::Triple::ppc64le:
   case llvm::Triple::renderscript64:
   case llvm::Triple::riscv64:
   case llvm::Triple::sparcv9:
   case llvm::Triple::spir64:
+  case llvm::Triple::spirv64:
   case llvm::Triple::systemz:
   case llvm::Triple::ve:
   case llvm::Triple::wasm64:
   case llvm::Triple::x86_64:
     return 64;
   }
   llvm_unreachable("Invalid architecture value");
 }
Triple Triple::get32BitArchVariant() const {
   case Triple::ppcle:
   case Triple::r600:
   case Triple::renderscript32:
   case Triple::riscv32:
   case Triple::shave:
   case Triple::sparc:
   case Triple::sparcel:
   case Triple::spir:
+  case Triple::spirv32:
   case Triple::tce:
   case Triple::tcele:
   case Triple::thumb:
   case Triple::thumbeb:
   case Triple::wasm32:
   case Triple::x86:
   case Triple::xcore:
     // Already 32-bit.
     break;
 
   case Triple::aarch64:        T.setArch(Triple::arm);     break;
   case Triple::aarch64_be:     T.setArch(Triple::armeb);   break;
   case Triple::amdil64:        T.setArch(Triple::amdil);   break;
   case Triple::hsail64:        T.setArch(Triple::hsail);   break;
   case Triple::le64:           T.setArch(Triple::le32);    break;
   case Triple::mips64:         T.setArch(Triple::mips);    break;
   case Triple::mips64el:       T.setArch(Triple::mipsel);  break;
   case Triple::nvptx64:        T.setArch(Triple::nvptx);   break;
   case Triple::ppc64:          T.setArch(Triple::ppc);     break;
   case Triple::ppc64le:        T.setArch(Triple::ppcle);   break;
   case Triple::renderscript64: T.setArch(Triple::renderscript32); break;
   case Triple::riscv64:        T.setArch(Triple::riscv32); break;
   case Triple::sparcv9:        T.setArch(Triple::sparc);   break;
   case Triple::spir64:         T.setArch(Triple::spir);    break;
+  case Triple::spirv64:        T.setArch(Triple::spirv32); break;
   case Triple::wasm64:         T.setArch(Triple::wasm32);  break;
   case Triple::x86_64:         T.setArch(Triple::x86);     break;
   }
   return T;
 }
 
 Triple Triple::get64BitArchVariant() const {
   Triple T(*this);
Triple Triple::get64BitArchVariant() const {
   case Triple::mips64el:
   case Triple::nvptx64:
   case Triple::ppc64:
   case Triple::ppc64le:
   case Triple::renderscript64:
   case Triple::riscv64:
   case Triple::sparcv9:
   case Triple::spir64:
+  case Triple::spirv64:
   case Triple::systemz:
   case Triple::ve:
   case Triple::wasm64:
   case Triple::x86_64:
     // Already 64-bit.
     break;
 
   case Triple::aarch64_32:      T.setArch(Triple::aarch64);    break;
   case Triple::amdil:           T.setArch(Triple::amdil64);    break;
   case Triple::arm:             T.setArch(Triple::aarch64);    break;
   case Triple::armeb:           T.setArch(Triple::aarch64_be); break;
   case Triple::hsail:           T.setArch(Triple::hsail64);    break;
   case Triple::le32:            T.setArch(Triple::le64);       break;
   case Triple::mips:            T.setArch(Triple::mips64);     break;
   case Triple::mipsel:          T.setArch(Triple::mips64el);   break;
   case Triple::nvptx:           T.setArch(Triple::nvptx64);    break;
   case Triple::ppc:             T.setArch(Triple::ppc64);      break;
   case Triple::ppcle:           T.setArch(Triple::ppc64le);    break;
   case Triple::renderscript32:  T.setArch(Triple::renderscript64);     break;
   case Triple::riscv32:         T.setArch(Triple::riscv64);    break;
   case Triple::sparc:           T.setArch(Triple::sparcv9);    break;
   case Triple::spir:            T.setArch(Triple::spir64);     break;
+  case Triple::spirv32:         T.setArch(Triple::spirv64);    break;
   case Triple::thumb:           T.setArch(Triple::aarch64);    break;
   case Triple::thumbeb:         T.setArch(Triple::aarch64_be); break;
   case Triple::wasm32:          T.setArch(Triple::wasm64);     break;
   case Triple::x86:             T.setArch(Triple::x86_64);     break;
   }
   return T;
 }
 
Triple Triple::getBigEndianArchVariant() const {
   case Triple::r600:
   case Triple::renderscript32:
   case Triple::renderscript64:
   case Triple::riscv32:
   case Triple::riscv64:
   case Triple::shave:
   case Triple::spir64:
   case Triple::spir:
+  case Triple::spirv32:
+  case Triple::spirv64:
   case Triple::wasm32:
   case Triple::wasm64:
   case Triple::x86:
   case Triple::x86_64:
   case Triple::xcore:
   case Triple::ve:
   case Triple::csky:
 
bool Triple::isLittleEndian() const {
   case Triple::renderscript32:
   case Triple::renderscript64:
   case Triple::riscv32:
   case Triple::riscv64:
   case Triple::shave:
   case Triple::sparcel:
   case Triple::spir64:
   case Triple::spir:
+  case Triple::spirv32:
+  case Triple::spirv64:
   case Triple::tcele:
   case Triple::thumb:
   case Triple::ve:
   case Triple::wasm32:
   case Triple::wasm64:
   case Triple::x86:
   case Triple::x86_64:
   case Triple::xcore:

llvm/unittests/ADT/TripleTest.cpp

TEST(TripleTest, ParsedIDs) {
   EXPECT_EQ(Triple::UnknownVendor, T.getVendor());
   EXPECT_EQ(Triple::UnknownOS, T.getOS());
 
   T = Triple("spir64-unknown-unknown");
   EXPECT_EQ(Triple::spir64, T.getArch());
   EXPECT_EQ(Triple::UnknownVendor, T.getVendor());
   EXPECT_EQ(Triple::UnknownOS, T.getOS());
 
+  T = Triple("spirv32-unknown-unknown");
+  EXPECT_EQ(Triple::spirv32, T.getArch());
+  EXPECT_EQ(Triple::UnknownVendor, T.getVendor());
+  EXPECT_EQ(Triple::UnknownOS, T.getOS());
+
+  T = Triple("spirv64-unknown-unknown");
+  EXPECT_EQ(Triple::spirv64, T.getArch());
+  EXPECT_EQ(Triple::UnknownVendor, T.getVendor());
+  EXPECT_EQ(Triple::UnknownOS, T.getOS());
+
   T = Triple("x86_64-unknown-ananas");
   EXPECT_EQ(Triple::x86_64, T.getArch());
   EXPECT_EQ(Triple::UnknownVendor, T.getVendor());
   EXPECT_EQ(Triple::Ananas, T.getOS());
   EXPECT_EQ(Triple::UnknownEnvironment, T.getEnvironment());
 
   T = Triple("x86_64-unknown-cloudabi");
   EXPECT_EQ(Triple::x86_64, T.getArch());
TEST(TripleTest, BitWidthPredicates) {
   EXPECT_TRUE(T.isArch32Bit());
   EXPECT_FALSE(T.isArch64Bit());
 
   T.setArch(Triple::spir64);
   EXPECT_FALSE(T.isArch16Bit());
   EXPECT_FALSE(T.isArch32Bit());
   EXPECT_TRUE(T.isArch64Bit());
 
+  T.setArch(Triple::spirv32);
+  EXPECT_FALSE(T.isArch16Bit());
+  EXPECT_TRUE(T.isArch32Bit());
+  EXPECT_FALSE(T.isArch64Bit());
+
+  T.setArch(Triple::spirv64);
+  EXPECT_FALSE(T.isArch16Bit());
+  EXPECT_FALSE(T.isArch32Bit());
+  EXPECT_TRUE(T.isArch64Bit());
+
   T.setArch(Triple::sparc);
   EXPECT_FALSE(T.isArch16Bit());
   EXPECT_TRUE(T.isArch32Bit());
   EXPECT_FALSE(T.isArch64Bit());
 
   T.setArch(Triple::sparcel);
   EXPECT_FALSE(T.isArch16Bit());
   EXPECT_TRUE(T.isArch32Bit());
TEST(TripleTest, BitWidthArchVariants) {
   T.setArch(Triple::spir);
   EXPECT_EQ(Triple::spir, T.get32BitArchVariant().getArch());
   EXPECT_EQ(Triple::spir64, T.get64BitArchVariant().getArch());
 
   T.setArch(Triple::spir64);
   EXPECT_EQ(Triple::spir, T.get32BitArchVariant().getArch());
   EXPECT_EQ(Triple::spir64, T.get64BitArchVariant().getArch());
 
+  T.setArch(Triple::spirv32);
+  EXPECT_EQ(Triple::spirv32, T.get32BitArchVariant().getArch());
+  EXPECT_EQ(Triple::spirv64, T.get64BitArchVariant().getArch());
+
+  T.setArch(Triple::spirv64);
+  EXPECT_EQ(Triple::spirv32, T.get32BitArchVariant().getArch());
+  EXPECT_EQ(Triple::spirv64, T.get64BitArchVariant().getArch());
+
   T.setArch(Triple::wasm32);
   EXPECT_EQ(Triple::wasm32, T.get32BitArchVariant().getArch());
   EXPECT_EQ(Triple::wasm64, T.get64BitArchVariant().getArch());
 
   T.setArch(Triple::wasm64);
   EXPECT_EQ(Triple::wasm32, T.get32BitArchVariant().getArch());
   EXPECT_EQ(Triple::wasm64, T.get64BitArchVariant().getArch());