[cfe-dev] RFC: Proposing an LLVM subproject for parallelism runtime and support libraries

Chandler Carruth via cfe-dev cfe-dev at lists.llvm.org
Wed Mar 9 14:22:56 PST 2016


It's all good. Also good for the Clang folks to be aware of the discussion
and context so they know to follow along on the llvm-dev thread. might also
be good to drop a line to openmp-dev that this discussion is taking place.

On Wed, Mar 9, 2016 at 11:21 PM Jason Henline <jhen at google.com> wrote:

> Thanks for the heads-up Chandler. I've moved this thread to llvm-dev.
> Let's consider this thread closed and move the discussion there.
>
> On Wed, Mar 9, 2016 at 2:12 PM Chandler Carruth <chandlerc at google.com>
> wrote:
>
>> FWIW, LLVM sub-project stuff probably is best discussed on llvm-dev.
>> Maybe re-send there?
>>
>> On Wed, Mar 9, 2016 at 10:30 PM Jason Henline via cfe-dev <
>> cfe-dev at lists.llvm.org> wrote:
>>
>>> At Google we're doing a lot of work on parallel programming models for
>>> CPUs, GPUs and other platforms. One place where we're investing a lot are
>>> parallel libraries, especially those closely tied to compiler technology
>>> like runtime and math libraries. We would like to develop these in the
>>> open, and the natural place seems to be as a subproject in LLVM if others
>>> in the community are interested.
>>>
>>> Initially, we'd like to open source our StreamExecutor runtime library,
>>> which is used for simplifying the management of data-parallel workflows on
>>> accelerator devices and can also be extended to support other hardware
>>> platforms. We'd like to teach Clang to use StreamExecutor when targeting
>>> CUDA and work on other integrations, but that makes much more sense if it
>>> is part of the LLVM project.
>>>
>>> However, we think the LLVM subproject should be organized as a set of
>>> several libraries with StreamExecutor as just the first instance. As just
>>> one example of how creating a unified parallelism subproject could help
>>> with code sharing, the StreamExecutor library contains some nice wrappers
>>> around the CUDA driver API and OpenCL API that create a unified API for
>>> managing all kinds of GPU devices. This unified GPU wrapper would be
>>> broadly applicable for libraries that need to communicate with GPU devices.
>>>
>>> Of course, there is already an LLVM subproject for a parallel runtime
>>> library: OpenMP! So there is a question of how it would fit into this
>>> picture.  Eventually, it might make sense to pull in the OpenMP project as
>>> a library in this proposed new subproject. In particular, there is a good
>>> chance that OpenMP and StreamExecutor could share code for offloading to
>>> GPUs and managing workloads on those devices. This is discussed at the end
>>> of the StreamExecutor documentation below. However, if it turns out that
>>> the needs of OpenMP are too specialized to fit well in a generic
>>> parallelism project, then it may make sense to leave OpenMP as a separate
>>> LLVM subproject so it can focus on serving the particular needs of OpenMP.
>>>
>>> Documentation for the StreamExecutor library that is being proposed for
>>> open-sourcing is included below to give a sense of what it is, in order to
>>> give context for how it might fit into a general parallelism LLVM
>>> subproject.
>>>
>>> What do folks think? Is there general interest in something like this?
>>> If so, we can start working on getting a project in place and sketching out
>>> a skeleton for how it would be organized, as well as contributing
>>> StreamExecutor to it. We're happy to iterate on the particulars to figure
>>> out what works for the community.
>>>
>>>
>>> =============================================
>>> StreamExecutor Runtime Library Documentation
>>> =============================================
>>>
>>>
>>> What is StreamExecutor?
>>> ========================
>>>
>>> **StreamExecutor** is a unified wrapper around the **CUDA** and
>>> **OpenCL** host-side programming models (runtimes). It lets host code
>>> target either CUDA or OpenCL devices with identically-functioning
>>> data-parallel kernels. StreamExecutor manages the execution of concurrent
>>> work targeting the accelerator similarly to how an Executor_ from the
>>> Google APIs client library manages the execution of concurrent work on the
>>> host.
>>>
>>> .. _Executor:
>>> http://google.github.io/google-api-cpp-client/latest/doxygen/classgoogleapis_1_1thread_1_1Executor.html
>>>
>>> StreamExecutor is currently used as the runtime for the vast majority of
>>> Google's internal GPGPU applications, and a snapshot of it is included in
>>> the open-source TensorFlow_ project, where it serves as the GPGPU runtime.
>>>
>>> .. _TensorFlow: https://www.tensorflow.org
>>>
>>> It is currently proposed that StreamExecutor itself be independently
>>> open-sourced. As part of that proposal, this document describes the basics
>>> of its design and explains why it would fit in well as an LLVM subproject.
>>>
>>>
>>> -------------------
>>> Key points
>>> -------------------
>>>
>>> StreamExecutor:
>>>
>>> * abstracts the underlying accelerator platform (avoids locking you into
>>> a single vendor, and lets you write code without thinking about which
>>> platform you'll be running on).
>>> * provides an open-source alternative to the CUDA runtime library.
>>> * gives users a stream management model whose terminology matches that
>>> of the CUDA programming model.
>>> * makes use of modern C++ to create a safe, efficient, easy-to-use
>>> programming interface.
>>>
>>> StreamExecutor makes it easy to:
>>>
>>> * move data between host and accelerator (and also between peer
>>> accelerators).
>>> * execute data-parallel kernels written in the OpenCL or CUDA kernel
>>> languages.
>>> * inspect the capabilities of a GPU-like device at runtime.
>>> * manage multiple devices.
>>>
>>>
>>> --------------------------------
>>> Example code snippet
>>> --------------------------------
>>>
>>> The StreamExecutor API uses abstractions that will be familiar to those
>>> who have worked with other GPU APIs: **Streams**, **Timers**, and
>>> **Kernels**. Its API is *fluent*, meaning that it allows the user to chain
>>> together a sequence of related operations on a stream, as in the following
>>> code snippet:
>>>
>>> .. code-block:: c++
>>>
>>>   se::Stream stream(executor);
>>>   se::Timer timer(executor);
>>>   stream.InitWithTimer(&timer)
>>>       .ThenStartTimer(&timer)
>>>       .ThenLaunch(se::ThreadDim(dim_block_x, dim_block_y),
>>>                   se::BlockDim(dim_grid_x, dim_grid_y),
>>>                   my_kernel,
>>>                   arg0, arg1, arg2)
>>>       .ThenStopTimer(&timer)
>>>       .BlockHostUntilDone();
>>>
>>> The name of the kernel being launched in the snippet above is
>>> `my_kernel` and the arguments being passed to the kernel are `arg0`,
>>> `arg1`, and `arg2`. Kernels with any number of arguments of any types are
>>> supported, and the number and types of the arguments is checked at compile
>>> time.
>>>
>>> How does it work?
>>> =======================
>>>
>>>
>>> --------------------------------
>>> Detailed example
>>> --------------------------------
>>>
>>> The following example shows how we can use StreamExecutor to create a
>>> `TypedKernel` instance, associate device code with that instance, and then
>>> use that instance to schedule work on an accelerator device.
>>>
>>> .. code-block:: c++
>>>
>>>     #include <cassert>
>>>
>>>     #include "stream_executor.h"
>>>
>>>     namespace se = streamexecutor;
>>>
>>>     // A PTX string defining a CUDA kernel.
>>>     //
>>>     // This PTX string represents a kernel that takes two arguments: an
>>> input value
>>>     // and an output pointer. The input value is a floating point
>>> number. The output
>>>     // value is a pointer to a floating point value in device memory.
>>> The output
>>>     // pointer is where the output from the kernel will be written.
>>>     //
>>>     // The kernel adds a fixed floating point value to the input and
>>> writes the
>>>     // result to the output location.
>>>     static constexpr const char *KERNEL_PTX = R"(
>>>         .version 3.1
>>>         .target sm_20
>>>         .address_size 64
>>>         .visible .entry add_mystery_value(
>>>             .param .f32 float_literal,
>>>             .param .u64 result_loc
>>>             ) {
>>>           .reg .u64 %rl<2>;
>>>           .reg .f32 %f<2>;
>>>           ld.param.f32 %f1, [float_literal];
>>>           ld.param.u64 %rl1, [result_loc];
>>>           add.f32 %f1, %f1, 123.0;
>>>           st.f32 [%rl1], %f1;
>>>           ret;
>>>         }
>>>         )";
>>>
>>>     // The number of arguments expected by the kernel described in
>>>     // KERNEL_PTX_TEMPLATE.
>>>     static constexpr int KERNEL_ARITY = 2;
>>>
>>>     // The name of the kernel described in KERNEL_PTX.
>>>     static constexpr const char *KERNEL_NAME = "add_mystery_value";
>>>
>>>     // The value added to the input in the kernel described in
>>> KERNEL_PTX.
>>>     static constexpr float MYSTERY_VALUE = 123.0f;
>>>
>>>     int main(int argc, char *argv[]) {
>>>       // Get a CUDA Platform object. (Other platforms such as OpenCL are
>>> also
>>>       // supported.)
>>>       se::Platform *platform =
>>>
>>> se::MultiPlatformManager::PlatformWithName("cuda").ValueOrDie();
>>>
>>>       // Get a StreamExecutor for the chosen Platform. Multiple devices
>>> are
>>>       // supported, we indicate here that we want to run on device 0.
>>>       const int device_ordinal = 0;
>>>       se::StreamExecutor *executor =
>>>           platform->ExecutorForDevice(device_ordinal).ValueOrDie();
>>>
>>>       // Create a MultiKernelLoaderSpec, which knows where to find the
>>> code for our
>>>       // kernel. In this case, the code is stored in memory as a PTX
>>> string.
>>>       //
>>>       // Note that the "arity" and name specified here must match
>>>  "arity" and name
>>>       // of the kernel defined in the PTX string.
>>>       se::MultiKernelLoaderSpec kernel_loader_spec(KERNEL_ARITY);
>>>       kernel_loader_spec.AddCudaPtxInMemory(KERNEL_PTX, KERNEL_NAME);
>>>
>>>       // Next create a kernel handle, which we will associate with our
>>> kernel code
>>>       // (i.e., the PTX string).  The type of this handle is a bit
>>> verbose, so we
>>>       // create an alias for it.
>>>       //
>>>       // This specific type represents a kernel that takes two
>>> arguments: a floating
>>>       // point value and a pointer to a floating point value in device
>>> memory.
>>>       //
>>>       // A type like this is nice to have because it enables static type
>>> checking of
>>>       // kernel arguments when we enqueue work on a stream.
>>>       using KernelType = se::TypedKernel<float, se::DeviceMemory<float>
>>> *>;
>>>
>>>       // Now instantiate an object of the specific kernel type we
>>> declared above.
>>>       // The kernel object is not yet connected with the device code
>>> that we want it
>>>       // to run (that happens with the call to GetKernel below), so it
>>> cannot be
>>>       // used to execute work on the device yet.
>>>       //
>>>       // However, the kernel object is not completely empty when it is
>>> created. From
>>>       // the StreamExecutor passed into its constructor it knows which
>>> platform it
>>>       // is targeted for, and it also knows which device it will run on.
>>>       KernelType kernel(executor);
>>>
>>>       // Use the MultiKernelLoaderSpec defined above to load the kernel
>>> code onto
>>>       // the device pointed to by the kernel object and to make that
>>> kernel object a
>>>       // handle to the kernel code loaded on that device.
>>>       //
>>>       // The MultiKernelLoaderSpec may contain code for several
>>> different platforms,
>>>       // but the kernel object has an associated platform, so there is
>>> no confusion
>>>       // about which code should be loaded.
>>>       //
>>>       // After this call the kernel object can be used to launch its
>>> kernel on its
>>>       // device.
>>>       executor->GetKernel(kernel_loader_spec, &kernel);
>>>
>>>       // Allocate memory in the device memory space to hold the result
>>> of the kernel
>>>       // call. This memory will be freed when this object goes out of
>>> scope.
>>>       se::ScopedDeviceMemory<float> result =
>>> executor->AllocateOwnedScalar<float>();
>>>
>>>       // Create a stream on which to schedule device operations.
>>>       se::Stream stream(executor);
>>>
>>>       // Schedule a kernel launch on the new stream and block until the
>>> kernel
>>>       // completes. The kernel call executes asynchronously on the
>>> device, so we
>>>       // could do more work on the host before calling
>>> BlockHostUntilDone.
>>>       const float kernel_input_argument = 42.5f;
>>>       stream.Init()
>>>           .ThenLaunch(se::ThreadDim(), se::BlockDim(), kernel,
>>>                       kernel_input_argument, result.ptr())
>>>           .BlockHostUntilDone();
>>>
>>>       // Copy the result of the kernel call from device back to the host.
>>>       float host_result = 0.0f;
>>>       executor->SynchronousMemcpyD2H(result.cref(), sizeof(host_result),
>>>                                      &host_result);
>>>
>>>       // Verify that the correct result was computed.
>>>       assert((kernel_input_argument + MYSTERY_VALUE) == host_result);
>>>     }
>>>
>>>
>>> --------------------------------
>>> Kernel Loader Specs
>>> --------------------------------
>>>
>>> An instance of the class `MultiKernelLoaderSpec` is used to encapsulate
>>> knowledge of where the device code for a kernel is stored and what format
>>> it is in.  Given a `MultiKernelLoaderSpec` and an uninitialized
>>> `TypedKernel`, calling the `StreamExecutor::GetKernel` method will load the
>>> code onto the device and associate the `TypedKernel` instance with that
>>> loaded code. So, in order to initialize a `TypedKernel` instance, it is
>>> first necessary to create a `MultiKernelLoaderSpec`.
>>>
>>> A `MultiKernelLoaderSpec` supports a different method for adding device
>>> code
>>> for each combination of platform, format, and storage location. The
>>> following
>>> table shows some examples:
>>>
>>> ===========     =======         ===========     =========================
>>> Platform        Format          Location        Setter
>>> ===========     =======         ===========     =========================
>>> CUDA            PTX             disk            `AddCudaPtxOnDisk`
>>> CUDA            PTX             memory          `AddCudaPtxInMemory`
>>> CUDA            cubin           disk            `AddCudaCubinOnDisk`
>>> CUDA            cubin           memory          `AddCudaCubinInMemory`
>>> OpenCL          text            disk            `AddOpenCLTextOnDisk`
>>> OpenCL          text            memory          `AddOpenCLTextInMemory`
>>> OpenCL          binary          disk            `AddOpenCLBinaryOnDisk`
>>> OpenCL          binary          memory          `AddOpenCLBinaryInMemory`
>>> ===========     =======         ===========     =========================
>>>
>>> The specific method used in the example is `AddCudaPtxInMemory`, but all
>>> other methods are used similarly.
>>>
>>>
>>> ------------------------------------
>>> Compiler Support for StreamExecutor
>>> ------------------------------------
>>>
>>>
>>> General strategies
>>> -------------------
>>>
>>> For illustrative purposes, the PTX code in the example is written by
>>> hand and appears as a string literal in the source code file, but it is far
>>> more typical for the kernel code to be expressed in a high level language
>>> like CUDA C++ or OpenCL C and for the device machine code to be generated
>>> by a compiler.
>>>
>>> There are several ways we can load compiled device code using
>>> StreamExecutor.
>>>
>>> One possibility is that the build system could write the compiled device
>>> code to a file on disk. This can then be added to a `MultiKernelLoaderSpec`
>>> by using one of the `OnDisk` setters.
>>>
>>> Another option is to add a feature to the compiler which embeds the
>>> compiled device code into the host executable and provides some symbol
>>> (probably with a name based on the name of the kernel) that allows the host
>>> code to refer to the embedded code data.
>>>
>>> In fact, as discussed below, in the current use of StreamExecutor inside
>>> Google, the compiler goes even further and generates an instance of
>>> `MultiKernelLoaderSpec` for each kernel. This means the application author
>>> doesn't have to know anything about how or where the compiler decided to
>>> store the compiled device code, but instead gets a pre-made loader object
>>> that handles all those details.
>>>
>>>
>>> Compiler-generated code makes things safe
>>> --------------------------------------------
>>>
>>> Two of the steps in the example above are dangerous because they lack
>>> static safety checks: instantiating the `MultiKernelLoaderSpec` and
>>> specializing the `TypedKernel` class template. This section discusses how
>>> compiler support for StreamExecutor can make these steps safe.
>>>
>>> Instantiating a `MultiKernelLoaderSpec` requires specifying a three
>>> things:
>>>
>>> 1. the kernel *arity* (number of parameters),
>>> 2. the kernel name,
>>> 3. a string containing the device machine code for the kernel (either as
>>> assembly, or some sort of object file).
>>>
>>> The problem with this is that the kernel name and the number of
>>> parameters is already fully determined by the kernel's machine code. In the
>>> best case scenario the *arity* and name arguments passed to the
>>> `MultiKernelLoaderSpec` methods match the information in the machine code
>>> and are simply redundant, but in the worst case these arguments contradict
>>> the information in the machine code and we get a runtime error when we try
>>> to load the kernel..
>>>
>>> The second unsafe operation is specifying the kernel parameter types as
>>> type arguments to the `TypedKernel` class template. The specified types
>>> must match the types defined in the kernel machine code, but again there is
>>> no compile-time checking that these types match. Failure to match these
>>> types will result in a runtime error when the kernel is launched.
>>>
>>> We would like the compiler to perform these checks for the application
>>> author, so as to eliminate this source of runtime errors. In particular, we
>>> want the compiler to create an appropriate `MultiKernelLoaderSpec` instance
>>> and `TypedKernel` specialization for each kernel definition.
>>>
>>> One of the main goals of open-sourcing StreamExecutor is to let us add
>>> this code generation capability to Clang, when the user has chosen to use
>>> StreamExecutor as their runtime for accelerator operations.
>>>
>>> Google has been using an internally developed CUDA compiler based on
>>> Clang called **gpucc** that generates code for StreamExecutor in this way.
>>> The code below shows how the example above would be written using gpucc to
>>> generate the unsafe parts of the code.
>>>
>>> The kernel is defined in a high-level language (CUDA C++ in this
>>> example) in its own file:
>>>
>>> .. code-block:: c++
>>>
>>>     // File: add_mystery_value.cu
>>>
>>>     __global__ void add_mystery_value(float input, float *output) {
>>>       *output = input + 42.0f;
>>>     }
>>>
>>>     The host code is defined in another file:
>>>
>>>     .. code-block:: c++
>>>
>>>     // File: example_host_code.cc
>>>
>>>     #include <cassert>
>>>
>>>     #include "stream_executor.h"
>>>
>>>     // This header is generated by the gpucc compiler and it contains the
>>>     // definitions of gpucc::kernel::AddMysteryValue and
>>>     // gpucc::spec::add_mystery_value().
>>>     //
>>>     // The name of this header file is derived from the name of the file
>>> containing
>>>     // the kernel code. The trailing ".cu" is replaced with ".gpu.h".
>>>     #include "add_mystery_value.gpu.h"
>>>
>>>     namespace se = streamexecutor;
>>>
>>>     int main(int argc, char *argv[]) {
>>>       se::Platform *platform =
>>>
>>> se::MultiPlatformManager::PlatformWithName("cuda").ValueOrDie();
>>>
>>>       const int device_ordinal = 0;
>>>       se::StreamExecutor *executor =
>>>           platform->ExecutorForDevice(device_ordinal).ValueOrDie();
>>>
>>>       // AddMysteryValue is an instance of TypedKernel generated by
>>> gpucc. The
>>>       // template arguments are chosen by the compiler to match the
>>> parameters of
>>>       // the add_mystery_value kernel.
>>>       gpucc::kernel::AddMysteryValue kernel(executor);
>>>
>>>       // gpucc::spec::add_mystery_value() is generated by gpucc. It
>>> returns a
>>>       // MultiKernelLoaderSpec that knows how to find  the compiled code
>>> for the
>>>       // add_mystery_value kernel.
>>>       executor->GetKernel(gpucc::spec::add_mystery_value(), &kernel);
>>>
>>>       se::ScopedDeviceMemory<float> result =
>>> executor->AllocateOwnedScalar<float>();
>>>       se::Stream stream(executor);
>>>
>>>       const float kernel_input_argument = 42.5f;
>>>
>>>       stream.Init()
>>>           .ThenLaunch(se::ThreadDim(), se::BlockDim(), kernel,
>>>                       kernel_input_argument, result.ptr())
>>>           .BlockHostUntilDone();
>>>
>>>       float host_result = 0.0f;
>>>       executor->SynchronousMemcpyD2H(result.cref(), sizeof(host_result),
>>>                                      &host_result);
>>>
>>>       assert((kernel_input_argument + 42.0f) == host_result);
>>>     }
>>>
>>> This support from the compiler makes the use of StreamExecutor safe and
>>> easy.
>>>
>>>
>>> Compiler support for triple angle bracket kernel launches
>>> ----------------------------------------------------------
>>>
>>> For even greater ease of use, Google's gpucc CUDA compiler also supports
>>> an integrated mode that looks like NVIDIA's `CUDA programming model`_,which
>>> uses triple angle brackets (`<<<>>>`) to launch kernels.
>>>
>>> .. _CUDA programming model:
>>> http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#kernels
>>>
>>> .. code-block:: c++
>>>     :emphasize-lines: 22
>>>
>>>     #include <cassert>
>>>
>>>     #include "stream_executor.h"
>>>
>>>     namespace se = streamexecutor;
>>>
>>>     __global__ void add_mystery_value(float input, float *output) {
>>>       *output = input + 42.0f;
>>>     }
>>>
>>>     int main(int argc, char *argv[]) {
>>>       se::Platform *platform =
>>>
>>> se::MultiPlatformManager::PlatformWithName("cuda").ValueOrDie();
>>>
>>>       const int device_ordinal = 0;
>>>       se::StreamExecutor *executor =
>>>           platform->ExecutorForDevice(device_ordinal).ValueOrDie();
>>>
>>>       se::ScopedDeviceMemory<float> result =
>>> executor->AllocateOwnedScalar<float>();
>>>
>>>       const float kernel_input_argument = 42.5f;
>>>       add_mystery_value<<<1, 1>>>(kernel_input_argument, *result.ptr());
>>>
>>>       float host_result = 0.0f;
>>>       executor->SynchronousMemcpyD2H(result.cref(), sizeof(host_result),
>>>                                      &host_result);
>>>
>>>       assert((kernel_input_argument + 42.0f) == host_result);
>>>     }
>>>
>>> Under the hood, gpucc converts the triple angle bracket kernel call into
>>> a series of calls to the StreamExecutor library similar to the calls seen
>>> in the previous examples.
>>>
>>> Clang currently supports the triple angle bracket kernel call syntax for
>>> CUDA compilation by replacing a triple angle bracket call with calls to the
>>> NVIDIA CUDA runtime library, but it would be easy to add a compiler flag to
>>> tell Clang to emit calls to the StreamExecutor library instead. There are
>>> several benefits to supporting this mode of compilation in Clang:
>>>
>>> .. _benefits-of-streamexecutor:
>>>
>>> * StreamExecutor is a high-level, modern C++ API, so is easier to use
>>> and less prone to error than the NVIDIA CUDA runtime and the OpenCL runtime.
>>> * StreamExecutor will be open-source software, so GPU code will not have
>>> to depend on opaque binary blobs like the NVIDIA CUDA runtime library.
>>> * Using StreamExecutor as the runtime would allow for easy extension of
>>> the triple angle bracket kernel launch syntax to support different
>>> accelerator programming models.
>>>
>>>
>>> Supporting other platforms
>>> ===========================
>>>
>>> StreamExecutor currently supports CUDA and OpenCL platforms
>>> out-of-the-box, but it uses a platform plugin architecture that makes it
>>> easy to add new platforms at any time. The CUDA and OpenCL platforms are
>>> both implemented as platform plugins in this way, so they serve as good
>>> examples for future platform developers of how to write these kinds of
>>> plugins.
>>>
>>>
>>> Canned operations
>>> ==================
>>>
>>> StreamExecutor provides several predefined kernels for common
>>> data-parallel operations. The supported classes of operations are:
>>>
>>> * BLAS: basic linear algebra subprograms,
>>> * DNN: deep neural networks,
>>> * FFT: fast Fourier transforms, and
>>> * RNG: random number generation.
>>>
>>> Here is an example of using a canned operation to perform random number
>>> generation:
>>>
>>> .. code-block:: c++
>>>     :emphasize-lines: 12-13,17,34-35
>>>
>>>     #include <array>
>>>
>>>     #include "cuda/cuda_rng.h"
>>>     #include "stream_executor.h"
>>>
>>>     namespace se = streamexecutor;
>>>
>>>     int main(int argc, char *argv[]) {
>>>       se::Platform *platform =
>>>
>>> se::MultiPlatformManager::PlatformWithName("cuda").ValueOrDie();
>>>
>>>       se::PluginConfig plugin_config;
>>>       plugin_config.SetRng(se::cuda::kCuRandPlugin);
>>>
>>>       const int device_ordinal = 0;
>>>       se::StreamExecutor *executor =
>>>           platform->ExecutorForDeviceWithPluginConfig(device_ordinal,
>>> plugin_config)
>>>               .ValueOrDie();
>>>
>>>       const uint8 seed[] = {0x0, 0x1, 0x2, 0x3, 0x4, 0x5, 0x6, 0x7,
>>>                             0x8, 0x9, 0xa, 0xb, 0xc, 0xd, 0xe, 0xf};
>>>       constexpr uint64 random_element_count = 1024;
>>>
>>>       using HostArray = std::array<float, random_element_count>;
>>>
>>>       HostArray host_memory;
>>>       const size_t data_size = host_memory.size() *
>>> sizeof(HostArray::value_type);
>>>
>>>       se::ScopedDeviceMemory<float> device_memory =
>>>           executor->AllocateOwnedArray<float>(random_element_count);
>>>
>>>       se::Stream stream(executor);
>>>       stream.Init()
>>>           .ThenSetRngSeed(seed, sizeof(seed))
>>>           .ThenPopulateRandUniform(device_memory.ptr())
>>>           .BlockHostUntilDone();
>>>
>>>       executor->SynchronousMemcpyD2H(*device_memory.ptr(), data_size,
>>>                                      host_memory.data());
>>>     }
>>>
>>> Each platform plugin can define its own canned operation plugins for
>>> these operations or choose to leave any of them unimplemented.
>>>
>>>
>>> Contrast with OpenMP
>>> =====================
>>>
>>> Recent versions of OpenMP also provide a high-level, easy-to-use
>>> interface for running data-parallel workloads on an accelerator device. One
>>> big difference between OpenMP's approach and that of StreamExecutor is that
>>> OpenMP generates both the kernel code that runs on the device and the
>>> host-side code needed to launch the kernel, whereas StreamExecutor only
>>> generates the host-side code. While the OpenMP model provides the
>>> convenience of allowing the author to write their kernel code in standard
>>> C/C++, the StreamExecutor model allows for the use of any kernel language
>>> (e.g. CUDA C++ or OpenCL C). This lets authors use  platform-specific
>>> features that are only present in platform-specific kernel definition
>>> languages.
>>>
>>> The philosophy of StreamExecutor is that performance is critical on the
>>> device, but less so on the host.  As a result, no attempt is made to use a
>>> high-level device abstraction during device code generation. Instead, the
>>> high-level abstraction provided by StreamExecutor is used only for the
>>> host-side code that moves data and launches kernels.  This host-side work
>>> is tedious and is not performance critical, so it benefits from being
>>> wrapped in a high-level library that can support a wide range of platforms
>>> in an easily extensible manner.
>>>
>>>
>>> Cooperation with OpenMP
>>> ========================
>>>
>>> The Clang OpenMP community is currently in the process of `designing
>>> their implementation`_ of offloading support. They will want the compiler
>>> to convert the various standardized target-oriented OpenMP pragmas into
>>> device code to execute on an accelerator and host code to load and run that
>>> device code. StreamExecutor may provide a convenient API for OpenMP to use
>>> to generate their host-side code.
>>>
>>> .. _designing their implementation:
>>> https://drive.google.com/a/google.com/file/d/0B-jX56_FbGKRM21sYlNYVnB4eFk/view
>>>
>>> In addition to the :ref:`benefits<benefits-of-streamexecutor>` that all
>>> users of StreamExecutor enjoy over the alternative host-side runtime
>>> libraries, OpenMP and StreamExecutor may mutually benefit by sharing work
>>> to support new platforms. If OpenMP makes use of StreamExecutor, then it
>>> should be simple for OpenMP to add support for any new platforms that
>>> StreamExecutor supports in the future. Similarly, for any platforms OpenMP
>>> would like to target, they may add that support in StreamExecutor and take
>>> advantage of the knowledge of platform support in the StreamExecutor
>>> community. The resulting new platform support would then be available not
>>> just within OpenMP, but also to any user of StreamExecutor.
>>>
>>> Although OpenMP and StreamExecutor support different programming models,
>>> some of the work they perform under the hood will likely be very similar.
>>> By sharing code and domain expertise, both projects will be improved and
>>> strengthened as their capabilities are expanded. The StreamExecutor
>>> community looks forward to much collaboration and discussion with OpenMP
>>> about the best places and ways to cooperate.
>>>
>>> _______________________________________________
>>> cfe-dev mailing list
>>> cfe-dev at lists.llvm.org
>>> http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-dev
>>>
>>
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.llvm.org/pipermail/cfe-dev/attachments/20160309/5144033a/attachment-0001.html>


More information about the cfe-dev mailing list