<html>
<head>
<meta http-equiv="Content-Type" content="text/html; charset=Windows-1252">
<style type="text/css" style="display:none;"> P {margin-top:0;margin-bottom:0;} </style>
</head>
<body dir="ltr">
<div style="font-family: Calibri, Arial, Helvetica, sans-serif; font-size: 12pt; color: rgb(0, 0, 0);">
Hi Pekka,</div>
<div style="font-family: Calibri, Arial, Helvetica, sans-serif; font-size: 12pt; color: rgb(0, 0, 0);">
<br>
</div>
<div style="font-family: Calibri, Arial, Helvetica, sans-serif; font-size: 12pt; color: rgb(0, 0, 0);">
Can you also clarify how do you plan to test the SPIR-V generation?</div>
<div style="font-family: Calibri, Arial, Helvetica, sans-serif; font-size: 12pt; color: rgb(0, 0, 0);">
<br>
</div>
<div style="font-family: Calibri, Arial, Helvetica, sans-serif; font-size: 12pt; color: rgb(0, 0, 0);">
Will you be adding IR only tests or also SPIR-V assembly/binary tests?</div>
<div style="font-family: Calibri, Arial, Helvetica, sans-serif; font-size: 12pt; color: rgb(0, 0, 0);">
If the latter one, does it mean that some LLVM test bots will need an</div>
<div style="font-family: Calibri, Arial, Helvetica, sans-serif; font-size: 12pt; color: rgb(0, 0, 0);">
installation and invocation of the translator?</div>
<div style="font-family: Calibri, Arial, Helvetica, sans-serif; font-size: 12pt; color: rgb(0, 0, 0);">
<br>
</div>
<div style="font-family: Calibri, Arial, Helvetica, sans-serif; font-size: 12pt; color: rgb(0, 0, 0);">
In general, it feels like at least we will likely be ending up with duplicate</div>
<div style="font-family: Calibri, Arial, Helvetica, sans-serif; font-size: 12pt; color: rgb(0, 0, 0);">
testing up until the transition phase is finalized.</div>
<div style="font-family: Calibri, Arial, Helvetica, sans-serif; font-size: 12pt; color: rgb(0, 0, 0);">
<br>
</div>
<div style="font-family: Calibri, Arial, Helvetica, sans-serif; font-size: 12pt; color: rgb(0, 0, 0);">
Cheers,<br>
Anastasia<br>
</div>
<div id="appendonsend"></div>
<div style="font-family:Calibri,Arial,Helvetica,sans-serif; font-size:12pt; color:rgb(0,0,0)">
<br>
</div>
<hr tabindex="-1" style="display:inline-block; width:98%">
<div id="divRplyFwdMsg" dir="ltr"><font style="font-size:11pt" face="Calibri, sans-serif" color="#000000"><b>From:</b> Pekka Jääskeläinen <pekka@parmance.com><br>
<b>Sent:</b> 25 August 2021 07:56<br>
<b>To:</b> Trifunovic, Konrad <konrad.trifunovic@intel.com><br>
<b>Cc:</b> Henry Linjamäki <henry.linjamaki@parmance.com>; Anastasia Stulova <Anastasia.Stulova@arm.com>; llvm-dev@lists.llvm.org <llvm-dev@lists.llvm.org>; yaxun.liu@amd.com <yaxun.liu@amd.com>; cfe-dev@lists.llvm.org <cfe-dev@lists.llvm.org>; Videau, Brice
<bvideau@anl.gov>; Harms, Kevin <harms@alcf.anl.gov><br>
<b>Subject:</b> Re: [llvm-dev] [RFC][HIPSPV] Emitting HIP device code as SPIR-V</font>
<div> </div>
</div>
<div class="BodyFragment"><font size="2"><span style="font-size:11pt">
<div class="PlainText">Hi Konrad,<br>
<br>
OK, thanks for the status update. We discussed this internally, and<br>
our current understanding is that the change from using llvm-spriv for<br>
the SPIR-V emission requires a) to change the target (in triple) from<br>
spir64 to spirv64 b) calling the llc that invokes the LLVM backend<br>
(BE) instead of the llvm-spriv translator tool.<br>
<br>
Thus, to make the transition and testing while developing the BE<br>
easier, we could contribute a patch to add a (hidden) command line<br>
option that switches between these modes: 1) The initial default: emit<br>
spir64 LLVM IR and call llvm-spirv. 2) Emit spirv64 LLVM IR instead<br>
and call a user-specified llc binary.<br>
<br>
When the LLVM BE lands upstream and is considered at least as<br>
useful/stable as llvm-spriv, we can then switch the default of that<br>
option to (2) instead of (1). This should allow easy testing of the<br>
BE-based toolchain even before it lands to the upstream repo if we<br>
allow redefining the llc binary location to the externally built LLVM<br>
BE. How does that sound?<br>
<br>
BR,<br>
Pekka and Henry<br>
<br>
On Tue, Aug 24, 2021 at 4:27 PM Trifunovic, Konrad via llvm-dev<br>
<llvm-dev@lists.llvm.org> wrote:<br>
><br>
> Hi,<br>
><br>
> Regarding the timeline for SPIR-V backend upstream, there is a tentative plan to land the base patches until end of this year (2021). Nevertheless, at that moment, SPIR-V backend will be in experimental state. After that, we need to keep on adding missing
functionality, fix bugs etc. so that we can pass a Khronos CTS test-suite with using a SPIR-V backend (instead of LLVM-SPIRV translator). At that moment, we might claim it is a production quality and not anymore experimental. That might take next several months.<br>
><br>
> I'm wondering - with the approach You have proposed - how much code that You plan to contribute as the intermediate step will become obsolete once there is a full switch to SPIR-V backend target?<br>
><br>
> konrad<br>
><br>
> > -----Original Message-----<br>
> > From: Henry Linjamäki <henry.linjamaki@parmance.com><br>
> > Sent: Monday, August 23, 2021 12:59 PM<br>
> > To: Anastasia Stulova <Anastasia.Stulova@arm.com><br>
> > Cc: cfe-dev@lists.llvm.org; llvm-dev@lists.llvm.org; yaxun.liu@amd.com;<br>
> > Trifunovic, Konrad <konrad.trifunovic@intel.com><br>
> > Subject: Re: [llvm-dev] [RFC][HIPSPV] Emitting HIP device code as SPIR-V<br>
> ><br>
> > On Mon, 23 Aug 2021 at 12:53, Anastasia Stulova<br>
> > <Anastasia.Stulova@arm.com> wrote:<br>
> > ><br>
> > > > Our planned HIP-SPIR-V tool chain calls the command-line tool<br>
> > > > ‘llvm-spirv’ for translating LLVM IR to SPIR-V, which is then<br>
> > > > embedded in the HIP binary as a byte array. The call to the<br>
> > > > llvm-spirv tool is used as a temporary solution until the SPIR-V backend lands<br>
> > in the?<br>
> > > > LLVM code base in the future. When the LLVM SPIR-V backend is usable<br>
> > > > in the upstream repo, we plan to simply switch to calling the LLVM’s<br>
> > > > internal ‘llc’ code generator tool for generating the SPIR-V instead of llvm-<br>
> > spirv.<br>
> > > > Can you immediately spot problems with that approach?<br>
> > ><br>
> > > Since the translator is not part of the LLVM project do you plan to<br>
> > > add any command-line options to set its location or any other<br>
> > > interface for this and if so would this be advertised as temporary<br>
> > > functionality? I imagine once you switch to the backend such<br>
> > > functionality would not be needed any longer so I am wondering how do you<br>
> > see exactly the transition path...<br>
> ><br>
> > We believe it might be enough to locate the tool (llvm-spirv) in PATH for now: A<br>
> > CLI option would be a nice addition, but perhaps not worth it for a temporary<br>
> > solution - we expect the SPIR-V backend (and thus<br>
> > llc) become usable sooner than later.<br>
> ><br>
> > ><br>
> > > And another aspect to clarify is the migration path i.e. what do you<br>
> > > see as criteria for changing to the backend and how soon could this happen?<br>
> ><br>
> > Calling the llvm-spirv in the PATH is not ideal: The tool’s version might be too old<br>
> > or new with respect to the calling LLVM’s version which might cause issues if<br>
> > there are incompatibilities with the LLVM IR. So, we don’t want to keep using<br>
> > the tool longer than necessary, and prefer to switch to the backend shipped with<br>
> > the LLVM installation as soon as it lands upstream. I think the criterion for<br>
> > switching to the BE instead of the llvm-spriv tool is when our internal HIP test<br>
> > suite (which will be published soon) passes on the upstream SPIR-V BE.<br>
> ><br>
> > ><br>
> > > Thanks,<br>
> > > Anastasia<br>
> > ><br>
> > > ________________________________<br>
> > > From: Henry Linjamäki <henry.linjamaki@parmance.com><br>
> > > Sent: 23 August 2021 07:32<br>
> > > To: Anastasia Stulova <Anastasia.Stulova@arm.com><br>
> > > Cc: cfe-dev@lists.llvm.org <cfe-dev@lists.llvm.org>;<br>
> > > llvm-dev@lists.llvm.org <llvm-dev@lists.llvm.org>; yaxun.liu@amd.com<br>
> > > <yaxun.liu@amd.com>; Trifunovic, Konrad <konrad.trifunovic@intel.com><br>
> > > Subject: Re: [llvm-dev] [RFC][HIPSPV] Emitting HIP device code as<br>
> > > SPIR-V<br>
> > ><br>
> > > Hi Anastasia,<br>
> > ><br>
> > > On Fri, 20 Aug 2021 at 12:44, Anastasia Stulova<br>
> > > <Anastasia.Stulova@arm.com> wrote:<br>
> > > ><br>
> > > > Hi Henry,<br>
> > > ><br>
> > > > Just to provide some background - we had a discussion about the<br>
> > > > integration of SPIRV-LLVM Translator some years back [1] and here<br>
> > > > is the design of our user interface that has been proposed at that<br>
> > > > time [2]. I appreciate we might not be able to unify the interfaces<br>
> > > > completely but it would make sense to provide common mechanisms for<br>
> > > > different languages to use SPIR-V even if it might not be achievable straight<br>
> > away we should aim for it as an end goal.<br>
> > > ><br>
> > > > Considering that there seems to be a lot of interest in this from<br>
> > > > different language communities, can we clarify the full plan? My<br>
> > > > understanding is that you are proposing to add the translator<br>
> > > > temporarily and it will be replaced by the backend in the future.<br>
> > > > How do you envision this transition? Do you plan to provide<br>
> > > > command-line options for the translator to be used that would be<br>
> > > > deprecated at some point later or would they be added as temporary from<br>
> > the start?<br>
> > ><br>
> > > Our planned HIP-SPIR-V tool chain calls the command-line tool<br>
> > > ‘llvm-spirv’ for translating LLVM IR to SPIR-V, which is then embedded<br>
> > > in the HIP binary as a byte array. The call to the llvm-spirv tool is<br>
> > > used as a temporary solution until the SPIR-V backend lands in the<br>
> > > LLVM code base in the future. When the LLVM SPIR-V backend is usable<br>
> > > in the upstream repo, we plan to simply switch to calling the LLVM’s<br>
> > > internal ‘llc’ code generator tool for generating the SPIR-V instead<br>
> > > of llvm-spirv. Can you immediately spot problems with that approach?<br>
> > ><br>
> > > We believe this is the best solution to integrate with the toolchain<br>
> > > infrastructure. Another would be to call the code generator at LLVM<br>
> > > API level, but it seems out of place for the toolchain framework: The<br>
> > > SPIR-V code generation path is not exposed so that clang frontends<br>
> > > could use it to emit SPIR-V for themselves.<br>
> > ><br>
> > > ><br>
> > > > Another consideration is that we have invested quite a lot of effort<br>
> > > > in the alternative approach i.e using the SPIR-V backend because<br>
> > > > this was highlighted as the best viable approach for SPIR-V support<br>
> > > > in Clang and LLVM when we had our discussion some years back. The<br>
> > > > situation is likely different now and your proposal isn�t<br>
> > > > identical, also we haven�t made a lot of progress with the backend<br>
> > > > yet. However, the integration of alternative SPIR-V translation<br>
> > > > might negatively impact the adoption of the backend. It might also<br>
> > > > result in either suboptimal design flow or code duplication in<br>
> > > > Clang. For example, we might need to redesign the OpenCL builtins<br>
> > representation and mapping to SPIR-V instructions.<br>
> > > ><br>
> > > > Considering that we might not be too far from integrating the<br>
> > > > backend into the LLVM, would it be reasonable to synchronize with<br>
> > > > the backend developers and see if the backend could be used straight<br>
> > > > away? I am looping in Konrad here who has been discussing the<br>
> > > > backend integration earlier this year [3]. Perhaps he can provide some<br>
> > insights about the backend work and the timeline for it.<br>
> > > > Maybe you could start working on some parts that are not related to<br>
> > > > IR consumption first and then add the SPIR-V emission later on and<br>
> > > > hopefully, the timing can align with the backend work too.<br>
> > > ><br>
> > > > [1]<br>
> > > > <a href="https://lists.llvm.org/pipermail/llvm-dev/2018-September/125948.html">
https://lists.llvm.org/pipermail/llvm-dev/2018-September/125948.html</a><br>
> > > > [2]<br>
> > > > <a href="https://github.com/KhronosGroup/SPIRV-LLVM-Translator/wiki/SPIRV-Too">
https://github.com/KhronosGroup/SPIRV-LLVM-Translator/wiki/SPIRV-Too</a><br>
> > > > lchain-for-Clang [3]<br>
> > > > <a href="https://lists.llvm.org/pipermail/llvm-dev/2021-March/148905.html">
https://lists.llvm.org/pipermail/llvm-dev/2021-March/148905.html</a><br>
> > > ><br>
> > > > Cheers,<br>
> > > > Anastasia<br>
> > > > ________________________________<br>
> > > > From: Henry Linjam�ki <henry.linjamaki@parmance.com><br>
> > > > Sent: 17 August 2021 17:16<br>
> > > > To: Anastasia Stulova <Anastasia.Stulova@arm.com><br>
> > > > Cc: cfe-dev@lists.llvm.org <cfe-dev@lists.llvm.org>;<br>
> > > > llvm-dev@lists.llvm.org <llvm-dev@lists.llvm.org>; yaxun.liu@amd.com<br>
> > > > <yaxun.liu@amd.com><br>
> > > > Subject: Re: [llvm-dev] [RFC][HIPSPV] Emitting HIP device code as<br>
> > > > SPIR-V<br>
> > > ><br>
> > > ><br>
> > > ><br>
> > > > Hi Anastasia,<br>
> > > ><br>
> > > > On Tue, 17 Aug 2021 at 12:53, Anastasia Stulova<br>
> > > > <Anastasia.Stulova@arm.com> wrote:<br>
> > > > ><br>
> > > > > Hi Henry,<br>
> > > > ><br>
> > > > > > Since the SPIR-V BE might not land in LLVM soon, we will set up<br>
> > > > > > the compilation flow to produce the SPIR-V binary by using the<br>
> > > > > > LLVM-SPIR-V translator [7] which is used in our experimental branch.<br>
> > > > ><br>
> > > > > Can you provide more details regarding this? Do you plan to<br>
> > > > > integrate the translator as an external tool?<br>
> > > > ><br>
> > > > The intention is to use the SPIRV-LLVM translator as a tool outside<br>
> > > > LLVM: either the tool is found in PATH or an error is emitted instead.<br>
> > > > Since I�m assuming that the new SPIR-V BE will eventually land on<br>
> > > > LLVM and supersede the translator, I don�t see much motivation for<br>
> > > > cleaner integration of it to the LLVM project.<br>
> > > ><br>
> > > > > Overall, there seem to be a huge overlap with what we need for<br>
> > > > > OpenCL so it would be good to make sure we are aligned and the new<br>
> > > > > functionality is reusable for OpenCL too.<br>
> > > > ><br>
> > > > Sure. I�m not so familiar with the OpenCL infra in the LLVM<br>
> > > > currently, so can you elaborate on any major overlap that OpenCL and<br>
> > HIPSPV have?<br>
> > > > We are planning to start submitting patches for review, perhaps<br>
> > > > it�s easier to point the overlaps on per-patch basis in the review<br>
> > > > system then?<br>
> > > ><br>
> > > > > Cheers,<br>
> > > > > Anastasia<br>
> > > > ><br>
> > > > ><br>
> > > > > ________________________________<br>
> > > > > From: llvm-dev <llvm-dev-bounces@lists.llvm.org> on behalf of<br>
> > > > > Henry Linjam�ki via llvm-dev <llvm-dev@lists.llvm.org><br>
> > > > > Sent: 09 August 2021 07:57<br>
> > > > > To: cfe-dev@lists.llvm.org <cfe-dev@lists.llvm.org><br>
> > > > > Cc: llvm-dev@lists.llvm.org <llvm-dev@lists.llvm.org>;<br>
> > > > > yaxun.liu@amd.com <yaxun.liu@amd.com><br>
> > > > > Subject: [llvm-dev] [RFC][HIPSPV] Emitting HIP device code as<br>
> > > > > SPIR-V<br>
> > > > ><br>
> > > > > Hi all,<br>
> > > > ><br>
> > > > > HIP is a C++ Runtime API and kernel language that allows<br>
> > > > > developers to create portable applications for AMD and NVIDIA GPUs<br>
> > > > > from a single source code [0]. There are also projects for running<br>
> > > > > HIP code on Intel GPU platforms via the Intel Level Zero API [1]<br>
> > > > > called HIPLZ [3] and HIPCL [2], which runs HIP programs in OpenCL<br>
> > > > > devices with certain advanced features supported. Both of these<br>
> > > > > backends consume SPIR-V binaries.<br>
> > > > ><br>
> > > > > We are proposing a patch set to be upstreamed that enables SPIR-V<br>
> > > > > emission through the HIP code path. The end goal of the patches to<br>
> > > > > be submitted is to emit SPIR-V binaries from HIP device code so it<br>
> > > > > can be embedded into executables for OpenCL-like environments (at<br>
> > > > > least for starters). Our current focus is on the two<br>
> > > > > above-mentioned projects, HIPCL and HIPLZ which are both<br>
> > > > > work-in-progress HIP implementations. They itself do not consume<br>
> > > > > SPIR-V, but the device binaries are handed over to the OpenCL and<br>
> > > > > Intel Level Zero APIs, respectively.<br>
> > > > ><br>
> > > > > Coarsely, the current process of translating the HIP code to<br>
> > > > > SPIR-V in LLVM/Clang involves:<br>
> > > > ><br>
> > > > > * Retargeting HIP device code generation to the SPIR-V target.<br>
> > > > > * Mapping address spaces in HIP to corresponding ones in SPIR-V.<br>
> > > > > * Expanding HIP features, which can not be directly modeled in SPIR-V<br>
> > > > > (e.g. dynamic shared memory).<br>
> > > > ><br>
> > > > > The HIPSPIRV experimental branch is available at [4]. Note that it<br>
> > > > > is not yet in a state we intend to propose for upstreaming, but<br>
> > > > > shaping up the patches is a work in progress. Before proceeding to<br>
> > > > > shape up and submit the patches, we would like to get feedback for<br>
> > > > > the plans we have for upstreaming. In the following sections, we<br>
> > > > > open up the above points further and sketch our plans for changes<br>
> > > > > to LLVM (mostly to the Clang tool) to achieve the goal.<br>
> > > > ><br>
> > > > > Retargeting device codegen<br>
> > > > > ==========================<br>
> > > > ><br>
> > > > > For making the HIP toolchain to emit and embed SPIR-V we are<br>
> > > > > tentatively planning the following changes to the LLVM/Clang:<br>
> > > > ><br>
> > > > > * Introduce, at minimum, a 'spirv64' architecture type in Triple. This<br>
> > > > > is what the SPIR-V backend [5] (SPIR-V BE) effort is planning to<br>
> > > > > upstream. We would like to upstream this change in advance to<br>
> > > > > specify the HIP SPIR-V device code target, potentially before the<br>
> > > > > SPIR-V BE work lands.<br>
> > > > ><br>
> > > > > * Implement a new SPIRVTargetInfo and fill it with necessary<br>
> > > > > information. For HIPCL/-LZ we are planning to adjust the address<br>
> > > > > space mapping in a way which is discussed later in the �address<br>
> > > > > space mapping� section.<br>
> > > > ><br>
> > > > > * Introduce a clang option to override the HIP device code target. We<br>
> > > > > are interested in the option �--offload=<target>� discussed in the<br>
> > > > > 'Unified offload option for CUDA/HIP/OpenMP'-thread [6]. This option<br>
> > > > > would suit this use case well. As far as we know, the subject has<br>
> > > > > not advanced further from the discussion - is anyone working on it?<br>
> > > > ><br>
> > > > > * Compilation driver:<br>
> > > > ><br>
> > > > > HIP offload builder is changed to retrieve the offload device target<br>
> > > > > from the --offload option. If it is not present, it can fall back to<br>
> > > > > AMD's default target for avoiding changing the current default HIP<br>
> > > > > compilation behavior.<br>
> > > > ><br>
> > > > > Temporarily change Driver to force clang to emit LLVM bitcode for<br>
> > > > > SPIR-V targets in the backend compilation phase. Otherwise, the<br>
> > > > > compilation will fail due to the lack of the real SPIR-V BE in many<br>
> > > > > parts of the code. Reworked HIPToolChain takes care of translating<br>
> > > > > the bitcode to SPIR-V during the linking phase. When the SPIR-V BE<br>
> > > > > lands in LLVM, we can revert this change.<br>
> > > > ><br>
> > > > > * Introduce �hipspv� as an OS or environment type in Triple. The<br>
> > > > > primary and the current use of the type is to select device offload<br>
> > > > > toolchain for HIPCL/-LZ.<br>
> > > > ><br>
> > > > > * Implement a new toolchain class 'HIPSPVToolChain' in clang which is<br>
> > > > > selected when the HIP device target is specified to be<br>
> > > > > �spirv64-unknown-hipspv� with the --offload option. Since the SPIR-V<br>
> > > > > BE might not land in LLVM soon, we will set up the compilation flow<br>
> > > > > to produce the SPIR-V binary by using the LLVM-SPIR-V translator [7]<br>
> > > > > which is used in our experimental branch.<br>
> > > > ><br>
> > > > > One important thing the toolchain does is to run one or several LLVM<br>
> > > > > IR passes, which are needed by the HIPCL/LZ runtime, on the final<br>
> > > > > fully linked device bitcode. The passes are required to be run<br>
> > > > > during link time - all user specified device code and HIPCL/LZ<br>
> > > > > device library routines have to be visible when the passes are<br>
> > > > > run. The reason for the requirement is explained in the 'HIP code<br>
> > > > > expansion' section. HIPSPVToolChain will use the opt tool for<br>
> > > > > running the passes at link time.<br>
> > > > ><br>
> > > > > * Currently, HIPToolChain is derived from ROCmToolchain and its long<br>
> > > > > chain of super classes (AMDGPUToolChain, Generic_ELF and<br>
> > > > > Generic_GCC). The new upstreamed target would not logically belong<br>
> > > > > under the AMDGPU/ROCm family so it does not make sense to derive the<br>
> > > > > HIPCL toolchain from the HIP toolchain. Therefore, we propose to:<br>
> > > > ><br>
> > > > > - Create a new base HIP tool chain, 'BaseHIPToolChain' or just<br>
> > > > > 'HIPToolChain', derived directly from ToolChain and put any<br>
> > > > > HIP-related code that is common or that can be reused in the<br>
> > > > > derived toolchains there.<br>
> > > > ><br>
> > > > > - Derive a new HIPSPVToolChain from HIPToolChain.<br>
> > > > ><br>
> > > > > - Rebase the HIPToolChain under the HIPToolChain and rename it to<br>
> > > > > HIPAMDToolChain. Since the current HIPToolChain depends on methods<br>
> > > > > in the super classes (e.g. AMDGPUToolChain�s getParsedTargetID)<br>
> > > > > the rebased class is planned to be a proxy class to avoid code<br>
> > > > > duplication and to reduce the amount of changes. Another option to<br>
> > > > > refactor the current HIPToolChain would be to use multiple<br>
> > > > > heritance but that leads to dreaded diamond class structure which<br>
> > > > > probably is not a great choice.<br>
> > > > ><br>
> > > > > With the current plan, HIPToolChain is not going to have much code<br>
> > > > > to be shared with the derived classes - so far only a bit of the<br>
> > > > > �fat binary� construction code is in sight for sharing, so the<br>
> > > > > immediate gains for the effort seems small. However, The TC�s layout<br>
> > > > > is more logical and it may spark more HIP implementations, as well<br>
> > > > > as help refactoring when going forward.<br>
> > > > ><br>
> > > > ><br>
> > > > > Address space mapping<br>
> > > > > =====================<br>
> > > > ><br>
> > > > > Translating HIP device code to valid SPIR-V binary requires tweaks<br>
> > > > > on<br>
> > > > > pointers:<br>
> > > > ><br>
> > > > > Pointers without address space (AS) qualification in HIP programs<br>
> > > > > are considered �flat� pointers - they can point to function<br>
> > > > > local, __device__, __shared__ and __constant__ memory space<br>
> > > > > dynamically, which matches the idea of �generic� pointers<br>
> > > > > introduced in OpenCL 2.0. Therefore, the logical choice for the<br>
> > > > > flat pointers is to map them to generic pointers of SPIR-V�s<br>
> > > > > OpenCL environment. HIPCL�s and HIPLZ�s SPIR-V environment<br>
> > > > > mandates that the kernel pointer parameters must point to<br>
> > > > > __global, __local or __constant memory (these are named<br>
> > > > > differently in SPIR-V; using OpenCL names as they are more<br>
> > > > > familiar). So HIP pointer parameters in the HIP kernel<br>
> > > > > (__global__) functions would be mapped to global pointers. Otherwise, HIP<br>
> > pointers with AS qualifiers are mapped to SPIR-V equivalent, if suitable.<br>
> > > > ><br>
> > > > > Now, there are significant differences between HIP�s<br>
> > > > > __constant__ and SPIR-V/OpenCL�s constant address space:<br>
> > > > ><br>
> > > > > * In HIP, __constant__ globals can be altered on the host side with<br>
> > > > > the hipMemcpyToSymbol() API function. In the OpenCL�s host API you<br>
> > > > > cannot do this.<br>
> > > > ><br>
> > > > > (Side-note: OpenCL host API does not have an equivalent method for<br>
> > > > > hipMemcpyToSymbol but HIPCL currently supports hipMemcpyToSymbol<br>
> > for<br>
> > > > > the global __global variables via Intel�s<br>
> > > > > clGetDeviceGlobalVariablePointerINTEL API extension, but we are<br>
> > > > > planning to inject shadow kernel commands that access the global<br>
> > > > > variables instead for portability.)<br>
> > > > ><br>
> > > > > * In HIP flat pointers can point to __constant__ memory. In OpenCL<br>
> > > > > this is not the case with __generic pointers, which means __constant<br>
> > > > > pointers cannot be casted to __generic pointers and vice versa.<br>
> > > > ><br>
> > > > > There are a couple ways to deal with constants:<br>
> > > > ><br>
> > > > > * Map __constant__ to __global space in SPIR-V. That way we can<br>
> > > > > generate code that works and is simple to implement. Of course, we<br>
> > > > > lose the optimization/placing benefits of constant memory.<br>
> > > > ><br>
> > > > > * Transform the code after clang codegen (by an LLVM pass) by<br>
> > > > > converting the __constant objects to kernel arguments. This covers<br>
> > > > > the hipMemcpyToSymbol() case. There is still the constant-to-generic<br>
> > > > > cast issue, so we would have to use the previous point as the<br>
> > > > > fallback.<br>
> > > > ><br>
> > > > > We plan to start by upstreaming the first option, and time<br>
> > > > > permitting, improve by implementing the second option.<br>
> > > > ><br>
> > > > > The planned changes to Clang to achieve the aforementioned AS<br>
> > > > > mapping are as follows:<br>
> > > > ><br>
> > > > > * Define address space mapping in the new, aforementioned<br>
> > > > > SPIRVTargetInfo to map CUDA address spaces (which the HIP reuses) to<br>
> > > > > do the mapping mentioned earlier. Default AS (0) used for the flat<br>
> > > > > pointers are mapped to the SPIR-V�s �generic�. We intend this<br>
> > > > > mapping being enabled when the language mode is HIP.<br>
> > > > ><br>
> > > > > * Change SPIRABIInfo to coerce kernel AS-unqualified pointer arguments<br>
> > > > > to __global ones. Pointer arguments in regular device functions<br>
> > > > > receive the __generic AS qualifier via the address space mapping<br>
> > > > > defined in SPIRVTargetInfo in the above point.<br>
> > > > ><br>
> > > > ><br>
> > > > > HIP code expansion<br>
> > > > > ==================<br>
> > > > ><br>
> > > > > There are features in HIP language which do not have direct<br>
> > > > > counterparts in SPIR-V�s OpenCL environment and those features<br>
> > > > > need to be rewritten before translation to SPIR-V (in the future,<br>
> > > > > lowering to SPIR-V machine code through the new BE). The<br>
> > > > > non-exhaustive list of features that need to be expanded includes:<br>
> > > > ><br>
> > > > > * Dynamic shared memory allocation (DSM): It is an array which is<br>
> > > > > declared globally in LLVM IR and its actual size determined at<br>
> > > > > kernel launch. OpTypeRuntimeArray in SPIR-V is the closest thing to<br>
> > > > > model this object, alas, it requires shader capability.<br>
> > > > ><br>
> > > > > * abort() builtin: No counterpart in SPIR-V/OpenCL.<br>
> > > > > (Note: the behavior is not well specified in the HIP spec<br>
> > > > > either. Assuming it terminates the whole grid if any work item<br>
> > > > > reaches it. AMD�s abort definition calls __builtin_trap).<br>
> > > > ><br>
> > > > > * printf(): OpenCL�s printf takes the format string as �__constant__<br>
> > > > > char*� while in HIP the format string does not have to reside in<br>
> > > > > constant memory.<br>
> > > > ><br>
> > > > > * Texture objects. These roughly correspond to image and sampler<br>
> > > > > objects of OpenCL combined. Also, texture objects carry more<br>
> > > > > information for the texture functions than image+sampler objects do.<br>
> > > > ><br>
> > > > > * Texture references. Same as above but these are program global<br>
> > > > > objects. In OpenCL, image objects cannot reside in the program<br>
> > > > > global space.<br>
> > > > ><br>
> > > > > HIPCL/-LZ�s solution to the DSM allocation case is that the<br>
> > > > > runtime allocates a shared buffer and passes it to the kernel as<br>
> > > > > an additional argument (which is hidden from the user). The device<br>
> > > > > code is modified so that the DSM object is replaced with the new<br>
> > > > > kernel argument. Various other cases listed will be handled similarly:<br>
> > > > ><br>
> > > > > * For the printf case we tentatively replace the printf calls with a<br>
> > > > > function that packs their arguments to an additional buffer passed<br>
> > > > > as additional kernel argument and do the printing on the host side.<br>
> > > > ><br>
> > > > > * Texture objects will be tentatively split to image and sampler<br>
> > > > > objects and possibly auxiliary struct to carry texture<br>
> > > > > settings. This means at least that the kernel parameter listing<br>
> > > > > needs to be rewritten for the Texture objects.<br>
> > > > ><br>
> > > > > * For the texture reference we tentatively planned replacing the<br>
> > > > > global texture objects also with a number of additional kernel<br>
> > > > > arguments.<br>
> > > > ><br>
> > > > > For this and other HIP features we need to apply LLVM IR passes to<br>
> > > > > perform modifications on the device code. In many cases the passes<br>
> > > > > should be run when the device code (as LLVM bitcode) is fully<br>
> > > > > linked. This is simply achieved as the HIP offload mechanism<br>
> > > > > already emits device code as LLVM bitcode in RDC mode (-fgpu-rdc),<br>
> > > > > so during linking we do receive the device code as LLVM bitcode<br>
> > > > > where to apply these expansions with full view of the device code.<br>
> > > > ><br>
> > > > > The current plan for implementing this is to make the<br>
> > > > > HIPSPVToolChain to build a linker that uses llvm-link for linking<br>
> > > > > device code, opt for running the IR passes needed and the external<br>
> > > > > llvm-spirv tool (llc in the future when the SPIR-V BE lands) for<br>
> > > > > emitting the SPIR-V binary. We load the passes from a path the<br>
> > > > > user provides via --hip-link-pass-path (name pending) or<br>
> > > > > automatically from HIP runtime�s installation location by using<br>
> > > > > the search logic provided by ROCmInstallationDetector.<br>
> > > > ><br>
> > > > > There is interest in upstreaming the HIPCL/-LZ passes from the<br>
> > > > > HIPCL/-LZ repositories in the future for reduced maintenance<br>
> > > > > burden. However, we are not attempting to upstream them initially,<br>
> > > > > as they are not yet completed and are subject to rapid changes.<br>
> > > > > Question<br>
> > > > > is: Where should the passes eventually be put in within the LLVM<br>
> > > > > project tree? Could it be OK to add a new directory under Clang<br>
> > > > > for tool chain passes?<br>
> > > > ><br>
> > > > ><br>
> > > > > Testing<br>
> > > > > =======<br>
> > > > ><br>
> > > > > We will provide llvm-lit tests for our toolchain in the upstream.<br>
> > > > > We also want to add tests to make sure clang who will run the<br>
> > > > > HIPCL/-LZ runtime passes get run at device code link time. For<br>
> > > > > this we need a dummy pass plugin that the clang loads during the test.<br>
> > > > ><br>
> > > > > When the new LLVM SPIR-V BE work lands on LLVM, we will add SPIR-V<br>
> > > > > assembly checks that are relevant for HIPSPV.<br>
> > > > ><br>
> > > > ><br>
> > > > > References<br>
> > > > > ==========<br>
> > > > ><br>
> > > > > [0]:<br>
> > > > ><br>
> > <a href="https://rocmdocs.amd.com/en/latest/Programming_Guides/Programming-">
https://rocmdocs.amd.com/en/latest/Programming_Guides/Programming-</a><br>
> > > > > Guides.html<br>
> > > > > [1]: <a href="https://spec.oneapi.com/level-zero/latest/index.html">https://spec.oneapi.com/level-zero/latest/index.html</a><br>
> > > > > [2]: <a href="https://github.com/cpc/hipcl">https://github.com/cpc/hipcl</a><br>
> > > > > [3]: <a href="https://github.com/jz10/anl-gt-gpu">https://github.com/jz10/anl-gt-gpu</a><br>
> > > > > [4]: <a href="https://github.com/parmance/llvm-project/tree/hip2spirv-v5">
https://github.com/parmance/llvm-project/tree/hip2spirv-v5</a><br>
> > > > > [5]: <a href="https://github.com/KhronosGroup/LLVM-SPIRV-Backend">https://github.com/KhronosGroup/LLVM-SPIRV-Backend</a><br>
> > > > > [6]:<br>
> > > > > <a href="https://lists.llvm.org/pipermail/cfe-dev/2020-December/067362.html">
https://lists.llvm.org/pipermail/cfe-dev/2020-December/067362.html</a><br>
> > > > > [7]: <a href="https://github.com/KhronosGroup/SPIRV-LLVM-Translator">https://github.com/KhronosGroup/SPIRV-LLVM-Translator</a><br>
> > > > > _______________________________________________<br>
> > > > > LLVM Developers mailing list<br>
> > > > > llvm-dev@lists.llvm.org<br>
> > > > > <a href="https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-dev">https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-dev</a><br>
> > > ><br>
> > > ><br>
> > > ><br>
> > > > --<br>
> > > > BR,<br>
> > > > Henry Linjam�ki<br>
> > > ><br>
> > ><br>
> > > BR,<br>
> > > Henry and Pekka<br>
> ><br>
> > BR,<br>
> > Henry and Pekka<br>
> _______________________________________________<br>
> LLVM Developers mailing list<br>
> llvm-dev@lists.llvm.org<br>
> <a href="https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-dev">https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-dev</a><br>
</div>
</span></font></div>
</body>
</html>