[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:12:14 PST 2016
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/d1273a2c/attachment.html>
More information about the cfe-dev
mailing list