[SYCL][CUDA] Initial Implementation of the CUDA backend

Contributors
Alan Forbes <[email protected]>
Alexander Johnston <[email protected]>
Bjoern Knafla <[email protected]>
Daniel Soutar <[email protected]>
David Wood <[email protected]>
Kumudha Narasimhan <[email protected]>
Mehdi Goli <[email protected]>
Przemek Malon <[email protected]>
Ruyman Reyes <[email protected]>
Stuart Adams <[email protected]>
Svetlozar Georgiev <[email protected]>
Steffen Larsen <[email protected]>
Victor Lomuller <[email protected]>

Signed-off-by: Alexander Johnston <[email protected]>

Author: Alexander Johnston

buildbot/configure.py

@@ -11,30 +11,49 @@ def do_configure(args):
     sycl_dir = os.path.join(args.src_dir, "sycl")
     spirv_dir = os.path.join(args.src_dir, "llvm-spirv")
     ocl_header_dir = os.path.join(args.obj_dir, "OpenCL-Headers")
-    icd_loader_lib = ''
+    icd_loader_lib = os.path.join(args.obj_dir, "OpenCL-ICD-Loader", "build")
+    llvm_targets_to_build = 'X86'
+    llvm_enable_projects = 'clang;llvm-spirv;sycl;opencl-aot'
+    libclc_targets_to_build = ''
+    sycl_build_pi_cuda = 'OFF'
+    llvm_enable_assertions = 'OFF'
 
     if platform.system() == 'Linux':
-      icd_loader_lib = os.path.join(args.obj_dir, "OpenCL-ICD-Loader", "build", "libOpenCL.so")
+        icd_loader_lib = os.path.join(icd_loader_lib, "libOpenCL.so")
     else:
-      icd_loader_lib = os.path.join(args.obj_dir, "OpenCL-ICD-Loader", "build", "OpenCL.lib")
+        icd_loader_lib = os.path.join(icd_loader_lib, "OpenCL.lib")
+
+    if args.cuda:
+        llvm_targets_to_build += ';NVPTX'
+        llvm_enable_projects += ';libclc'
+        libclc_targets_to_build = 'nvptx64--;nvptx64--nvidiacl'
+        sycl_build_pi_cuda = 'ON'
+
+    if args.assertions:
+        llvm_enable_assertions = 'ON'
 
     install_dir = os.path.join(args.obj_dir, "install")
 
-    cmake_cmd = ["cmake",
-                 "-G", "Ninja",
-                 "-DCMAKE_BUILD_TYPE={}".format(args.build_type),
-                 "-DLLVM_EXTERNAL_PROJECTS=sycl;llvm-spirv;opencl-aot",
-                 "-DLLVM_EXTERNAL_SYCL_SOURCE_DIR={}".format(sycl_dir),
-                 "-DLLVM_EXTERNAL_LLVM_SPIRV_SOURCE_DIR={}".format(spirv_dir),
-                 "-DLLVM_ENABLE_PROJECTS=clang;sycl;llvm-spirv;opencl-aot",
-                 "-DOpenCL_INCLUDE_DIR={}".format(ocl_header_dir),
-                 "-DOpenCL_LIBRARY={}".format(icd_loader_lib),
-                 "-DLLVM_BUILD_TOOLS=ON",
-                 "-DSYCL_ENABLE_WERROR=ON",
-                 "-DLLVM_ENABLE_ASSERTIONS=ON",
-                 "-DCMAKE_INSTALL_PREFIX={}".format(install_dir),
-                 "-DSYCL_INCLUDE_TESTS=ON", # Explicitly include all kinds of SYCL tests.
-                 llvm_dir]
+    cmake_cmd = [
+        "cmake",
+        "-G", "Ninja",
+        "-DCMAKE_BUILD_TYPE={}".format(args.build_type),
+        "-DLLVM_ENABLE_ASSERTIONS={}".format(llvm_enable_assertions),
+        "-DLLVM_TARGETS_TO_BUILD={}".format(llvm_targets_to_build),
+        "-DLLVM_EXTERNAL_PROJECTS=sycl;llvm-spirv;opencl-aot",
+        "-DLLVM_EXTERNAL_SYCL_SOURCE_DIR={}".format(sycl_dir),
+        "-DLLVM_EXTERNAL_LLVM_SPIRV_SOURCE_DIR={}".format(spirv_dir),
+        "-DLLVM_ENABLE_PROJECTS={}".format(llvm_enable_projects),
+        "-DLIBCLC_TARGETS_TO_BUILD={}".format(libclc_targets_to_build),
+        "-DOpenCL_INCLUDE_DIR={}".format(ocl_header_dir),
+        "-DOpenCL_LIBRARY={}".format(icd_loader_lib),
+        "-DSYCL_BUILD_PI_CUDA={}".format(sycl_build_pi_cuda),
+        "-DLLVM_BUILD_TOOLS=ON",
+        "-DSYCL_ENABLE_WERROR=ON",
+        "-DCMAKE_INSTALL_PREFIX={}".format(install_dir),
+        "-DSYCL_INCLUDE_TESTS=ON", # Explicitly include all kinds of SYCL tests.
+        llvm_dir
+    ]
 
     print(cmake_cmd)
 
@@ -63,6 +82,8 @@ def main():
     parser.add_argument("-o", "--obj-dir", metavar="OBJ_DIR", required=True, help="build directory")
     parser.add_argument("-t", "--build-type",
                         metavar="BUILD_TYPE", required=True, help="build type, debug or release")
+    parser.add_argument("--cuda", action='store_true', help="switch from OpenCL to CUDA")
+    parser.add_argument("--assertions", action='store_true', help="build with assertions")
 
     args = parser.parse_args()
 
@@ -74,4 +95,3 @@ def main():
     ret = main()
     exit_code = 0 if ret else 1
     sys.exit(exit_code)
-

sycl/CMakeLists.txt

@@ -139,6 +139,9 @@ install(DIRECTORY ${OPENCL_INCLUDE}/CL
 
 # Configure SYCL version macro
 set(sycl_inc_dir ${CMAKE_CURRENT_SOURCE_DIR}/include)
+set(sycl_src_dir ${CMAKE_CURRENT_SOURCE_DIR}/source)
+set(sycl_detail_inc_dir ${CMAKE_CURRENT_SOURCE_DIR}/include/CL/sycl/detail)
+set(sycl_detail_src_dir ${CMAKE_CURRENT_SOURCE_DIR}/source/detail)
 string(TIMESTAMP __SYCL_COMPILER_VERSION "%Y%m%d")
 set(version_header "${sycl_inc_dir}/CL/sycl/version.hpp")
 configure_file("${version_header}.in" "${version_header}")
@@ -197,7 +200,6 @@ add_subdirectory( source )
 # SYCL toolchain builds all components: compiler, libraries, headers, etc.
 add_custom_target( sycl-toolchain
   DEPENDS ${SYCL_RT_LIBS}
-          pi_opencl
           clang
           clang-offload-wrapper
           clang-offload-bundler
@@ -256,6 +258,20 @@ set( SYCL_TOOLCHAIN_DEPLOY_COMPONENTS
      pi_opencl
 )
 
+
+if(SYCL_BUILD_PI_CUDA)
+  # Ensure that libclc is enabled.
+  list(FIND LLVM_ENABLE_PROJECTS libclc LIBCLC_FOUND)
+  if( LIBCLC_FOUND EQUAL -1 )
+    message(FATAL_ERROR
+        "CUDA support requires adding \"libclc\" to the CMake argument \"LLVM_ENABLE_PROJECTS\"")
+  endif()
+
+  add_dependencies(sycl-toolchain libspirv-builtins)
+  list(APPEND SYCL_TOOLCHAIN_DEPLOY_COMPONENTS libspirv-builtins)
+endif()
+
+
 # Use it as fake dependency in order to force another command(s) to execute.
 add_custom_command(OUTPUT __force_it
   COMMAND "${CMAKE_COMMAND}" -E echo

sycl/doc/GetStartedWithSYCLCompiler.md

@@ -10,12 +10,14 @@ OpenCL™ API to offload computations to accelerators.
   * [Create SYCL workspace](#create-sycl-workspace)
 * [Build SYCL toolchain](#build-sycl-toolchain)
   * [Build SYCL toolchain with libc++ library](#build-sycl-toolchain-with-libc-library)
+  * [Build SYCL toolchain with support for NVIDIA CUDA](#build-sycl-toolchain-with-support-for-nvidia-cuda)
 * [Use SYCL toolchain](#use-sycl-toolchain)
   * [Install low level runtime](#install-low-level-runtime)
   * [Test SYCL toolchain](#test-sycl-toolchain)
   * [Run simple SYCL application](#run-simple-sycl-application)
 * [C++ standard](#c-standard)
 * [Known Issues and Limitations](#known-issues-and-limitations)
+* [CUDA backend limitations](#cuda-backend-limitations)
 * [Find More](#find-more)
 
 # Prerequisites
@@ -115,6 +117,28 @@ should be used.
 -DSYCL_LIBCXX_LIBRARY_PATH=<path to libc++ and libc++abi libraries>
 ```
 
+## Build SYCL toolchain with support for NVIDIA CUDA
+
+There is experimental support for SYCL for CUDA devices.
+
+To enable support for CUDA devices, the following arguments need to be added to 
+the CMake command when building the SYCL compiler.
+
+```
+-DCUDA_TOOLKIT_ROOT_DIR=/usr/local/cuda/ \
+-DLLVM_ENABLE_PROJECTS="clang;llvm-spirv;sycl;libclc"
+-DSYCL_BUILD_PI_CUDA=ON 
+```
+
+Enabling this flag requires an installation of 
+[CUDA 10.1](https://developer.nvidia.com/cuda-10.1-download-archive-update2) on the system,
+refer to 
+[NVIDIA CUDA Installation Guide for Linux](https://docs.nvidia.com/cuda/cuda-installation-guide-linux/index.html).
+
+Currently, the only combination tested is Ubuntu 18.04 with CUDA 10.2 using
+a Titan RTX GPU (SM 71), but it should work on any GPU compatible with SM 50 or
+above.
+
 # Use SYCL toolchain
 
 ## Install low level runtime
@@ -354,15 +378,32 @@ and run following command:
 clang++ -fsycl simple-sycl-app.cpp -o simple-sycl-app.exe
 ```
 
+When building for CUDA, use the CUDA target triple as follows:
+
+```bash
+clang++ -fsycl -fsycl-targets=nvptx64-nvidia-cuda-sycldevice \
+  simple-sycl-app.cpp -o simple-sycl-app-cuda.exe
+```
+
 This `simple-sycl-app.exe` application doesn't specify SYCL device for
 execution, so SYCL runtime will use `default_selector` logic to select one
 of accelerators available in the system or SYCL host device.
 
+Note: `nvptx64-nvidia-cuda-sycldevice` is usable with `-fsycl-targets`
+if clang was built with the cmake option `SYCL_BUILD_PI_CUDA=ON`.
+
 **Linux & Windows**
 ```bash
 ./simple-sycl-app.exe
 The results are correct!
 ```
+**Note**:
+Currently, when the application has been built with the CUDA target, the CUDA backend
+must be selected at runtime using the `SYCL_BE` environment variable. 
+
+```bash
+SYCL_BE=PI_CUDA ./simple-sycl-app-cuda.exe
+```
 
 NOTE: SYCL developer can specify SYCL device for execution using device
 selectors (e.g. `cl::sycl::cpu_selector`, `cl::sycl::gpu_selector`,
@@ -414,7 +455,28 @@ int main() {
 
 ```
 
+The device selector below selects an NVIDIA device only, and won't
+execute if there is none.
+
+```c++
+class CUDASelector : public cl::sycl::device_selector {
+  public:
+    int operator()(const cl::sycl::device &Device) const override {
+      using namespace cl::sycl::info;
+
+      const std::string DeviceName = Device.get_info<device::name>();
+      const std::string DeviceVendor = Device.get_info<device::vendor>();
+
+      if (Device.is_gpu() && (DeviceName.find("NVIDIA") != std::string::npos)) {
+        return 1;
+      };
+      return -1;
+    }
+};
+```
+
 # C++ standard
+
 - Minimally support C++ standard is c++11 on Linux and c++14 on Windows.
 
 # Known Issues and Limitations
@@ -426,6 +488,15 @@ int main() {
 - SYCL works only with OpenCL implementations supporting out-of-order queues.
 - On Windows linking SYCL applications with `/MTd` flag is known to cause crashes.
 
+## CUDA back-end limitations
+
+- Backend is only supported on Linux 
+- The only combination tested is Ubuntu 18.04 with CUDA 10.2 using
+a Titan RTX GPU (SM 71), but it should work on any GPU compatible with SM 50 or
+above
+- The NVIDIA OpenCL headers conflict with the OpenCL headers required for this project 
+and may cause compilation issues on some platforms
+
 # Find More
 
 SYCL 1.2.1 specification: [www.khronos.org/registry/SYCL/specs/sycl-1.2.1.pdf](https://www.khronos.org/registry/SYCL/specs/sycl-1.2.1.pdf)