<div dir="ltr">In response to the latest questions from C Bergström:<br><div><br></div><div><div>Is</div><div>there "CUDA" or OpenCL hidden in the headers and that's where the</div><div>actual offload portion is happening</div></div><div><br></div><div><div> Does StreamExecutor</div><div>wrapper around public or private CUDA/OpenCL runtimes?</div></div><div><br></div><div>Yes, StreamExecutor is a wrapper around the public OpenCL and CUDA userspace driver libraries. The actual offloading is achieved by making calls to those libraries.<br></div><div><br></div><div><div>Is there anything stopping you from exposing "wrapper" interfaces</div><div>which are the same as the NVIDIA runtime?</div></div><div><br></div><div>There is nothing stopping us from doing that. The reason we haven't to this point is because we felt the current StreamExecutor API was nicer to work with.<br></div><div><br></div><div>Where is the StreamExecutor runtime source now?<br></div><div><br></div><div>It is currently housed in Google's internal code repo, where it is being used in production code. There is also a local copy in the open-source TensorFlow project (<a href="https://www.tensorflow.org">https://www.tensorflow.org</a>) which we want to replace with a dependency on a separate open source StreamExecutor project.<br></div><div><br></div><div><div>/*</div><div>I have said this before and I really get uncomfortable with the</div><div>generic term "CUDA" in clang. Until someone from NVIDIA (lawyers) put</div><div>something in writing. CUDA is an NV trademark and clang/llvm project</div><div>can't claim to be "CUDA" and need to make a distinction. Informally</div><div>this is all friendly now, but I do hope it's officially clarified at</div><div>some point. Maybe it's as simple as saying "CUDA compatible" - I don't</div><div>know..</div><div>*/</div></div><div><br></div><div>Good point! I will try to keep that in mind.<br></div><div><br></div><div><div>I think having a nice model that lowers cleanly (high performance) to</div><div>at least some targets is (should be) very important. From my</div><div>experience - if you have complex or perfectly nested loops - how would</div><div>you take this sort of algorithm and map it to StreamExecutor? Getting</div><div>reductions right or wrong can also have a performance impact - If your</div><div>goal is to create a "one wrapper rules them all" approach - I'm hoping</div><div>you can find a common way to also make it easier for basic needs to be</div><div>expressed to the underlying target. (In a target agnostic way)</div></div><div><br></div><div>I'm not quite sure how to answer this in all generality, but here are some thoughts. Any complex or nested looping control flow that happens on data stored in device memory can be handled completely within the kernel definition, and should be independent of StreamExecutor. If the complexity arises instead from coordinating data transfers to device memory with kernel launches, then StreamExecutor proposes to model those dependencies as "streams" where one operation can be forced to wait on another (much in the way CUDA streams work). It would be possible to create new "canned" operations to perform common operations like reductions where the data won't all fit on the device at once, but those canned operations would probably not be optimal for all platforms, and in those cases the user might need to roll their own.<br></div><div><br></div><div><div>Microsoft did a really nice job of documenting C++AMP - Does google</div><div>have a bunch of example codes which show how StreamExecutor can be</div><div>used to implement various algorithms?</div></div><div><br></div><div>We don't currently have any simplified public examples, but I agree that would be something useful to have. I may write up a few in the coming weeks.<br></div><div><br></div><div><div>Does clang/llvm</div><div>accept anything or is there some metric for generally deciding what</div><div>should get a sub-project and what just is too early.</div></div><div><br></div><div>I'm a newcomer to the community myself, so I'll leave this to others to give a better answer than I could.<br></div><div><br></div><div><div>Does Google have a plan to engage and bring other</div><div>stakeholders into supporting this? </div></div><div><br></div><div>We see this unified model as a benifit to all accelerator platforms because we think it will make it easier for programmers to use their systems. We plan to propose this model to these vendors and see if we can get them interested in providing code or advertising this model as a way to program their devices.<br></div><div><br></div><div>I hope all my questions are viewed as positive and meant to be constructive.<br></div><div><br></div><div>Absolutely. I feel that your input has been very constructive, and I appreciate you helping us think through this design.<br></div></div><br><div class="gmail_quote"><div dir="ltr">On Thu, Mar 10, 2016 at 9:54 AM Hal Finkel <<a href="mailto:hfinkel@anl.gov">hfinkel@anl.gov</a>> wrote:<br></div><blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex">----- Original Message -----<br>
> From: "Arpith C Jacob" <<a href="mailto:acjacob@us.ibm.com" target="_blank">acjacob@us.ibm.com</a>><br>
> To: <a href="mailto:llvm-dev@lists.llvm.org" target="_blank">llvm-dev@lists.llvm.org</a><br>
> Cc: <a href="mailto:jhen@google.com" target="_blank">jhen@google.com</a>, "Hal J. Finkel" <<a href="mailto:hfinkel@anl.gov" target="_blank">hfinkel@anl.gov</a>><br>
> Sent: Thursday, March 10, 2016 10:38:46 AM<br>
> Subject: Re: [llvm-dev] RFC: Proposing an LLVM subproject for parallelism runtime and support libraries<br>
><br>
> Hi Jason,<br>
><br>
> I'm trying to better understand your StreamExecutor proposal and how<br>
> it relates to other parallel programming models and runtimes such as<br>
> RAJA [1], KOKKOS [2], or some hypothetical SPARK C++ API.<br>
><br>
> Please correct me if I'm misunderstanding your proposal, but I think<br>
> the essence of what you want from the compiler is type safety for<br>
> accelerator kernel launches, i.e., you would like the frontend to<br>
> parse, check, and codegen for the construct:<br>
> add_mystery_value<<<1, 1>>>(kernel_input_argument, *result.ptr());<br>
><br>
> Is that a correct understanding?<br>
><br>
<br>
Without answering your question, I'll point out that, as I understand it, StreamExecutor completely replaces the CUDA userspace library runtime components and talks directly to the drivers. Jason, please correct me if I'm wrong.<br>
<br>
 -Hal<br>
<br>
> Thanks,<br>
> Arpith<br>
><br>
> [1]<br>
> <a href="http://computation.llnl.gov/projects/raja-managing-application-portability-next-generation-platforms" rel="noreferrer" target="_blank">http://computation.llnl.gov/projects/raja-managing-application-portability-next-generation-platforms</a><br>
> [2] <a href="https://github.com/kokkos/kokkos" rel="noreferrer" target="_blank">https://github.com/kokkos/kokkos</a><br>
><br>
<br>
--<br>
Hal Finkel<br>
Assistant Computational Scientist<br>
Leadership Computing Facility<br>
Argonne National Laboratory<br>
</blockquote></div>