<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);">
<font size="2"><span style="font-size:11pt">> 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 of llvm-spirv.</span></font></div>
<div style="font-family: Calibri, Arial, Helvetica, sans-serif; font-size: 12pt; color: rgb(0, 0, 0);">
<font size="2"><span style="font-size:11pt">> Can you immediately spot problems with that approach?</span></font></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);">
Since the translator is not part of the LLVM project do you plan to add any</div>
<div style="font-family: Calibri, Arial, Helvetica, sans-serif; font-size: 12pt; color: rgb(0, 0, 0);">
command-line options to set its location or any other interface for this and</div>
<div style="font-family: Calibri, Arial, Helvetica, sans-serif; font-size: 12pt; color: rgb(0, 0, 0);">
if so would this be advertised as temporary functionality? I imagine once <br>
</div>
<div style="font-family: Calibri, Arial, Helvetica, sans-serif; font-size: 12pt; color: rgb(0, 0, 0);">
you switch to the backend such functionality would not be needed any</div>
<div style="font-family: Calibri, Arial, Helvetica, sans-serif; font-size: 12pt; color: rgb(0, 0, 0);">
longer so I am wondering how do you see exactly the transition path...<br>
</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);">
And another aspect to clarify is the migration path i.e. what do you see as</div>
<div style="font-family: Calibri, Arial, Helvetica, sans-serif; font-size: 12pt; color: rgb(0, 0, 0);">
criteria for changing to the backend and how soon could this happen?</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);">
Thanks,<br>
Anastasia<br>
</div>
<div style="font-family: Calibri, Arial, Helvetica, sans-serif; font-size: 12pt; color: rgb(0, 0, 0);">
<br>
</div>
<div id="appendonsend"></div>
<hr style="display:inline-block;width:98%" tabindex="-1">
<div id="divRplyFwdMsg" dir="ltr"><font face="Calibri, sans-serif" style="font-size:11pt" color="#000000"><b>From:</b> Henry Linjamäki <henry.linjamaki@parmance.com><br>
<b>Sent:</b> 23 August 2021 07:32<br>
<b>To:</b> Anastasia Stulova <Anastasia.Stulova@arm.com><br>
<b>Cc:</b> cfe-dev@lists.llvm.org <cfe-dev@lists.llvm.org>; llvm-dev@lists.llvm.org <llvm-dev@lists.llvm.org>; yaxun.liu@amd.com <yaxun.liu@amd.com>; Trifunovic, Konrad <konrad.trifunovic@intel.com><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 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 integration<br>
> of SPIRV-LLVM Translator some years back [1] and here is the design of our<br>
> user interface that has been proposed at that time [2]. I appreciate we might<br>
> not be able to unify the interfaces completely but it would make sense to<br>
> provide common mechanisms for different languages to use SPIR-V even if it<br>
> might not be achievable straight 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 different<br>
> language communities, can we clarify the full plan? My understanding is that<br>
> you are proposing to add the translator temporarily and it will be replaced by<br>
> the backend in the future. How do you envision this transition? Do you plan to<br>
> provide 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 the<br>
> 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 in the<br>
> alternative approach i.e using the SPIR-V backend because this was highlighted<br>
> as the best viable approach for SPIR-V support in Clang and LLVM when we<br>
> had our discussion some years back. The situation is likely different now and<br>
> your proposal isn�t identical, also we haven�t made a lot of progress with the<br>
> backend yet. However, the integration of alternative SPIR-V translation might<br>
> negatively impact the adoption of the backend. It might also result in either<br>
> suboptimal design flow or code duplication in Clang. For example, we might need<br>
> to redesign the OpenCL builtins representation and mapping to SPIR-V<br>
> instructions.<br>
><br>
> Considering that we might not be too far from integrating the backend into the<br>
> LLVM, would it be reasonable to synchronize with the backend developers and<br>
> see if the backend could be used straight away? I am looping in Konrad here<br>
> who has been discussing the backend integration earlier this year [3]. Perhaps<br>
> he can provide some insights about the backend work and the timeline for it.<br>
> Maybe you could start working on some parts that are not related to IR<br>
> consumption first and then add the SPIR-V emission later on and hopefully, the<br>
> timing can align with the backend work too.<br>
><br>
> [1] <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] <a href="https://github.com/KhronosGroup/SPIRV-LLVM-Translator/wiki/SPIRV-Toolchain-for-Clang">
https://github.com/KhronosGroup/SPIRV-LLVM-Translator/wiki/SPIRV-Toolchain-for-Clang</a><br>
> [3] <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>; llvm-dev@lists.llvm.org <llvm-dev@lists.llvm.org>; yaxun.liu@amd.com <yaxun.liu@amd.com><br>
> Subject: Re: [llvm-dev] [RFC][HIPSPV] Emitting HIP device code as 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 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>
> > Can you provide more details regarding this? Do you plan to integrate the<br>
> > 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 LLVM<br>
> and supersede the translator, I don�t see much motivation for cleaner<br>
> integration of it to the LLVM project.<br>
><br>
> > Overall, there seem to be a huge overlap with what we need for OpenCL so it would<br>
> > be good to make sure we are aligned and the new functionality is reusable for OpenCL<br>
> > too.<br>
> ><br>
> Sure. I�m not so familiar with the OpenCL infra in the LLVM currently,<br>
> so can you elaborate on any major overlap that OpenCL and HIPSPV have?<br>
> We are planning to start submitting patches for review, perhaps it�s<br>
> easier to point the overlaps on per-patch basis in the review system<br>
> then?<br>
><br>
> > Cheers,<br>
> > Anastasia<br>
> ><br>
> ><br>
> > ________________________________<br>
> > From: llvm-dev <llvm-dev-bounces@lists.llvm.org> on behalf of 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>; yaxun.liu@amd.com <yaxun.liu@amd.com><br>
> > Subject: [llvm-dev] [RFC][HIPSPV] Emitting HIP device code as SPIR-V<br>
> ><br>
> > Hi all,<br>
> ><br>
> > HIP is a C++ Runtime API and kernel language that allows developers to<br>
> > create portable applications for AMD and NVIDIA GPUs from a single<br>
> > source code [0]. There are also projects for running HIP code on Intel<br>
> > GPU platforms via the Intel Level Zero API [1] called HIPLZ [3] and<br>
> > HIPCL [2], which runs HIP programs in OpenCL devices with certain<br>
> > advanced features supported. Both of these backends consume SPIR-V<br>
> > 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 be<br>
> > submitted is to emit SPIR-V binaries from HIP device code so it can be<br>
> > embedded into executables for OpenCL-like environments (at least for<br>
> > starters). Our current focus is on the two above-mentioned projects,<br>
> > HIPCL and HIPLZ which are both work-in-progress HIP<br>
> > implementations. They itself do not consume SPIR-V, but the device<br>
> > binaries are handed over to the OpenCL and Intel Level Zero APIs,<br>
> > respectively.<br>
> ><br>
> > Coarsely, the current process of translating the HIP code to SPIR-V in<br>
> > 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 is<br>
> > not yet in a state we intend to propose for upstreaming, but shaping<br>
> > up the patches is a work in progress. Before proceeding to shape up<br>
> > and submit the patches, we would like to get feedback for the plans we<br>
> > have for upstreaming. In the following sections, we open up the above<br>
> > points further and sketch our plans for changes to LLVM (mostly to the<br>
> > 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 on<br>
> > pointers:<br>
> ><br>
> > Pointers without address space (AS) qualification in HIP programs are<br>
> > considered �flat� pointers - they can point to function local,<br>
> > __device__, __shared__ and __constant__ memory space dynamically,<br>
> > which matches the idea of �generic� pointers introduced in OpenCL<br>
> > 2.0. Therefore, the logical choice for the flat pointers is to map<br>
> > them to generic pointers of SPIR-V�s OpenCL environment. HIPCL�s and<br>
> > HIPLZ�s SPIR-V environment mandates that the kernel pointer parameters<br>
> > must point to __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 (__global__)<br>
> > functions would be mapped to global pointers. Otherwise, HIP pointers<br>
> > with AS qualifiers are mapped to SPIR-V equivalent, if suitable.<br>
> ><br>
> > Now, there are significant differences between HIP�s __constant__ and<br>
> > 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 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 permitting,<br>
> > improve by implementing the second option.<br>
> ><br>
> > The planned changes to Clang to achieve the aforementioned AS mapping<br>
> > 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 need to<br>
> > be rewritten before translation to SPIR-V (in the future, lowering to<br>
> > SPIR-V machine code through the new BE). The non-exhaustive list of<br>
> > 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 runtime<br>
> > allocates a shared buffer and passes it to the kernel as an additional<br>
> > argument (which is hidden from the user). The device code is modified<br>
> > so that the DSM object is replaced with the new kernel<br>
> > 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 already<br>
> > emits device code as LLVM bitcode in RDC mode (-fgpu-rdc), so during<br>
> > linking we do receive the device code as LLVM bitcode where to apply<br>
> > these expansions with full view of the device code.<br>
> ><br>
> > The current plan for implementing this is to make the HIPSPVToolChain<br>
> > to build a linker that uses llvm-link for linking device code, opt for<br>
> > running the IR passes needed and the external llvm-spirv tool (llc in<br>
> > the future when the SPIR-V BE lands) for emitting the SPIR-V<br>
> > binary. We load the passes from a path the user provides<br>
> > via --hip-link-pass-path (name pending) or automatically from HIP<br>
> > runtime�s installation location by using the search logic provided by<br>
> > 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, as<br>
> > they are not yet completed and are subject to rapid changes. 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 for<br>
> > tool chain passes?<br>
> ><br>
> ><br>
> > Testing<br>
> > =======<br>
> ><br>
> > We will provide llvm-lit tests for our toolchain in the upstream. We<br>
> > also want to add tests to make sure clang who will run the HIPCL/-LZ<br>
> > runtime passes get run at device code link time. For this we need a<br>
> > 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]: <a href="https://rocmdocs.amd.com/en/latest/Programming_Guides/Programming-Guides.html">
https://rocmdocs.amd.com/en/latest/Programming_Guides/Programming-Guides.html</a><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]: <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>
</div>
</span></font></div>
</body>
</html>