[all-commits] [llvm/llvm-project] f4d3cb: [HIPSPV] Add CUDA->SPIR-V address space mapping

[Anastasia Stulova via All-commits]

  Branch: refs/heads/main
  Home:   https://github.com/llvm/llvm-project
  Commit: f4d3cb4ca833d0b165d199e78ed8f1d59a8032eb
      https://github.com/llvm/llvm-project/commit/f4d3cb4ca833d0b165d199e78ed8f1d59a8032eb
  Author: Anastasia Stulova <anastasia.stulova at arm.com>
  Date:   2021-12-02 (Thu, 02 Dec 2021)

  Changed paths:
    M clang/lib/Basic/Targets/SPIR.h
    A clang/test/CodeGenHIP/hipspv-addr-spaces.cpp

  Log Message:
  -----------
  [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