eomii · aaronmondal · May 9, 2023
diff --git a/MODULE.bazel b/MODULE.bazel
@@ -40,6 +40,7 @@ llvm_project_overlay.configure(
         "@rules_ll//patches:rules_ll_overlay_patch.diff",
         "@rules_ll//patches:clang_new_offload_driver.diff",
         "@rules_ll//patches:llvm_use_zlib-ng.diff",
+        "@rules_ll//patches:libcxx_combined_init_patch.diff",
     ],
 )
 
@@ -62,6 +63,7 @@ use_repo(
     "comgr",
     "hip",
     "hipamd",
+    "opensycl",
     "rocclr",
     "rocm-device-libs",
     "rocm-opencl-runtime",

diff --git a/bazel-wrapper/default.nix b/bazel-wrapper/default.nix
@@ -60,6 +60,16 @@ LL_AMD_RPATHS=${(pkgs.lib.concatStringsSep ":" [
   "-rpath=${pkgs.libglvnd}/lib"
   "-rpath=${pkgs.xorg.libX11}/lib"
 ])}
+
+LL_SYCL_INCLUDES=-isystem${(pkgs.boost182.override {
+  enableStatic = true;
+  enableShared = false;
+}).dev}/include
+
+LL_SYCL_LIBRARIES=-L${(pkgs.boost182.override {
+  enableStatic = true;
+  enableShared = false;
+})}/lib
 ${pkgs.lib.strings.optionalString unfree ''
 
 # Flags for CUDA dependencies.
@@ -103,6 +113,8 @@ if [[
         --action_env=LL_AMD_INCLUDES=$LL_AMD_INCLUDES \
         --action_env=LL_AMD_LIBRARIES=$LL_AMD_LIBRARIES \
         --action_env=LL_AMD_RPATHS=$LL_AMD_RPATHS \
+        --action_env=LL_SYCL_INCLUDES=$LL_SYCL_INCLUDES \
+        --action_env=LL_SYCL_LIBRARIES=$LL_SYCL_LIBRARIES \
         --action_env=LL_CUDA_TOOLKIT=$LL_CUDA_TOOLKIT \
         --action_env=LL_CUDA_RUNTIME=$LL_CUDA_RUNTIME \
         --action_env=LL_CUDA_DRIVER=$LL_CUDA_DRIVER \

diff --git a/examples/sycl_cpu_example/BUILD.bazel b/examples/sycl_cpu_example/BUILD.bazel
@@ -0,0 +1,18 @@
+load("@rules_ll//ll:defs.bzl", "ll_binary")
+
+ll_binary(
+    name = "amdgpu",
+    srcs = ["main.cpp"],
+    compilation_mode = "sycl_amdgpu",
+    compile_flags = [
+        "-std=c++20",
+        "--offload-arch=gfx1100",
+    ],
+)
+
+ll_binary(
+    name = "cpu",
+    srcs = ["main.cpp"],
+    compilation_mode = "sycl_cpu",
+    compile_flags = ["-std=c++20"],
+)
diff --git a/examples/sycl_cpu_example/main.cpp b/examples/sycl_cpu_example/main.cpp
@@ -0,0 +1,39 @@
+#include <SYCL/sycl.hpp>
+#include <iostream>
+
+class vector_addition;
+
+int main(int, char **) {
+  sycl::float4 a = {1.0, 2.0, 3.0, 4.0};
+  sycl::float4 b = {4.0, 3.0, 1.0, 1.0};
+  sycl::float4 c = {0.0, 0.0, 0.0, 0.0};
+
+  sycl::queue queue(sycl::default_selector{});
+  std::cout << "Running on "
+            << queue.get_device().get_info<sycl::info::device::name>() << "\n";
+
+  {  // start of scope, ensures data copied back to host
+    sycl::buffer<sycl::float4, 1> a_sycl(&a, sycl::range<1>{1});
+    sycl::buffer<sycl::float4, 1> b_sycl(&b, sycl::range<1>{1});
+    sycl::buffer<sycl::float4, 1> c_sycl(&c, sycl::range<1>{1});
+
+    queue.submit([&](sycl::handler &cgh) {
+      auto a_acc = a_sycl.get_access<sycl::access::mode::read>(cgh);
+      auto b_acc = b_sycl.get_access<sycl::access::mode::read>(cgh);
+      auto c_acc = c_sycl.get_access<sycl::access::mode::write>(cgh);
+
+      cgh.single_task<class vector_addition>(
+          [=]() { c_acc[0] = a_acc[0] + b_acc[0]; });
+    });
+  }  // end of scope, ensures data copied back to host
+
+  std::cout << "  A { " << a.x() << ", " << a.y() << ", " << a.z() << ", "
+            << a.w() << " }\n"
+            << "+ B { " << b.x() << ", " << b.y() << ", " << b.z() << ", "
+            << b.w() << " }\n"
+            << "------------------\n"
+            << "= C { " << c.x() << ", " << c.y() << ", " << c.z() << ", "
+            << c.w() << " }" << std::endl;
+
+  return 0;
+}
diff --git a/ll/BUILD.bazel b/ll/BUILD.bazel
@@ -134,10 +134,15 @@ ll_toolchain(
             "@hip//:headers",
             "@hipamd//:headers",
         ],
+        ":sycl_amdgpu": [
+            "@hip//:headers",
+            "@hipamd//:headers",
+        ],
         "//conditions:default": [],
     }),
     hip_runtime = select({
         ":hip_amdgpu": "@hipamd//:libamdhip64",
+        ":sycl_amdgpu": "@hipamd//:libamdhip64",
         "//conditions:default": None,
     }),
     leak_sanitizer = select({
@@ -148,6 +153,7 @@ ll_toolchain(
     }),
     libomp = select({
         "omp_cpu": "@llvm-project//openmp:libomp",
+        "sycl_cpu": "@llvm-project//openmp:libomp",
         "//conditions:default": None,
     }),
     memory_sanitizer = select({
@@ -159,6 +165,7 @@ ll_toolchain(
     }),
     omp_header = select({
         "omp_cpu": "@llvm-project//openmp:omp_header",
+        "sycl_cpu": "@llvm-project//openmp:omp_header",
         "//conditions:default": None,
     }),
     profile = select({
@@ -167,6 +174,31 @@ ll_toolchain(
     }),
     rocm_device_libs = select({
         ":hip_amdgpu": "@rocm-device-libs//:rocm-device-libs",
+        ":sycl_amdgpu": "@rocm-device-libs//:rocm-device-libs",
+        "//conditions:default": None,
+    }),
+    sycl_headers = select({
+        ":sycl_amdgpu": "@opensycl//:sycl_headers",
+        ":sycl_cpu": "@opensycl//:sycl_headers",
+        "//conditions:default": None,
+    }),
+    sycl_hip_backend = select({
+        ":sycl_amdgpu": "@opensycl//:librt-backend-hip",
+        "//conditions:default": None,
+    }),
+    sycl_omp_backend = select({
+        ":sycl_cpu": "@opensycl//:librt-backend-omp",
+        ":sycl_amdgpu": "@opensycl//:librt-backend-omp",
+        "//conditions:default": None,
+    }),
+    sycl_plugin = select({
+        ":sycl_amdgpu": "@opensycl//:opensycl-clang",
+        ":sycl_cpu": "@opensycl//:opensycl-clang",
+        "//conditions:default": None,
+    }),
+    sycl_runtime = select({
+        ":sycl_amdgpu": "@opensycl//:libhipSYCL-rt",
+        ":sycl_cpu": "@opensycl//:libhipSYCL-rt",
         "//conditions:default": None,
     }),
     thread_sanitizer = select({

diff --git a/ll/args.bzl b/ll/args.bzl
@@ -181,6 +181,7 @@ def compile_object_args(
         "cuda_nvptx",
         "hip_amdgpu",
         "hip_nvptx",
+        "sycl_amdgpu",
     ] and ctx.var["COMPILATION_MODE"] != "dbg":
         args.add_all(["-Xarch_device", "-O3"])
 
@@ -207,7 +208,7 @@ def compile_object_args(
     args.add("-fPIC")
 
     # Maybe enable OpenMP.
-    if ctx.attr.compilation_mode == "omp_cpu":
+    if ctx.attr.compilation_mode in ["omp_cpu", "sycl_cpu"]:
         args.add("-fopenmp")
 
         # TODO: This is obviously not the way lol.
@@ -221,6 +222,7 @@ def compile_object_args(
         "cuda_nvptx",
         "hip_amdgpu",
         "hip_nvptx",
+        "sycl_amdgpu",
     ]:
         args.add("--offload-new-driver")
 
@@ -235,7 +237,7 @@ def compile_object_args(
                 ctx.configuration.default_shell_env["LL_CUDA_TOOLKIT"],
                 format = "--cuda-path=%s",
             )
-    if ctx.attr.compilation_mode in ["hip_nvptx", "hip_amdgpu"]:
+    if ctx.attr.compilation_mode in ["hip_nvptx", "hip_amdgpu", "sycl_amdgpu"]:
         args.add_all(
             [
                 Label("@hip").workspace_root,
@@ -244,9 +246,22 @@ def compile_object_args(
             format_each = "-I%s/include",
         )
 
+    if ctx.attr.compilation_mode in ["sycl_amdgpu", "sycl_cpu"]:
+        args.add_all(
+            [
+                Label("@opensycl").workspace_root,
+                paths.join(
+                    ctx.var["GENDIR"],  # For hipSYCL/common/config.hpp
+                    Label("@opensycl").workspace_root,
+                    "opensycl",
+                ),
+            ],
+            format_each = "-I%s/include",
+        )
+
     clang_resource_dir = paths.join(llvm_bindir_path(ctx), "clang/staging")
 
-    if ctx.attr.compilation_mode == "hip_amdgpu":
+    if ctx.attr.compilation_mode in ["hip_amdgpu", "sycl_amdgpu"]:
         args.add("-xhip")
         args.add(toolchain.hip_runtime.path, format = "--rocm-path=%s")
         args.add(clang_resource_dir, format = "-isystem%s")
@@ -300,7 +315,7 @@ def compile_object_args(
             # become system includes.
             format_each = "-isystem%s",
         )
-        for flags in ["LL_CFLAGS", "LL_AMD_INCLUDES"]:
+        for flags in ["LL_CFLAGS", "LL_AMD_INCLUDES", "LL_SYCL_INCLUDES"]:
             if ctx.configuration.default_shell_env.get(flags) != "":
                 args.add_all(
                     ctx.configuration.default_shell_env[flags].split(":"),
@@ -353,6 +368,16 @@ def compile_object_args(
     # re-enable abi-tagging.
     args.add("-D_LIBCPP_NO_ABI_TAG")
 
+    if ctx.attr.compilation_mode == "sycl_amdgpu":
+        args.add(
+            "-D__HIPSYCL_ENABLE_HIP_TARGET__",
+            "-D__HIPSYCL_CLANG__",
+        )
+    if ctx.attr.compilation_mode == "sycl_cpu":
+        args.add(
+            "-D__HIPSYCL_ENABLE_OMPHOST_TARGET__",
+        )
+
     # Additional compile flags.
     args.add_all(ctx.attr.compile_flags)
 
@@ -419,7 +444,12 @@ def link_executable_args(ctx, in_files, out_file, mode):
     if ctx.var["COMPILATION_MODE"] == "dbg":
         args.add("--verbose")
 
-    for flags in ["LL_LDFLAGS", "LL_AMD_LIBRARIES", "LL_AMD_RPATHS"]:
+    for flags in [
+        "LL_LDFLAGS",
+        "LL_AMD_LIBRARIES",
+        "LL_AMD_RPATHS",
+        "LL_SYCL_LIBRARIES",
+    ]:
         if flags in ctx.configuration.default_shell_env.keys():
             args.add_all(
                 ctx.configuration.default_shell_env[flags].split(":"),
@@ -507,7 +537,7 @@ def link_executable_args(ctx, in_files, out_file, mode):
         args.add("-lcuda")
         args.add("-lcudart_static")
         args.add("-lcupti_static")
-    if ctx.attr.compilation_mode == "hip_amdgpu":
+    if ctx.attr.compilation_mode in ["hip_amdgpu", "sycl_amdgpu"]:
         args.add(toolchain.hip_runtime.dirname, format = "-L%s")
         args.add(toolchain.hip_runtime.basename, format = "-l:%s")
         hip_runtime_rpath = paths.join(
@@ -519,6 +549,18 @@ def link_executable_args(ctx, in_files, out_file, mode):
             hip_runtime_rpath,
             format = "-rpath=$ORIGIN/%s",
         )
+    if ctx.attr.compilation_mode in ["sycl_amdgpu", "sycl_cpu"]:
+        args.add(toolchain.sycl_runtime.dirname, format = "-L%s")
+        args.add(toolchain.sycl_runtime.basename, format = "-l:%s")
+        sycl_runtime_rpath = paths.join(
+            "{}.runfiles".format(ctx.label.name),
+            ctx.workspace_name,
+            paths.dirname(toolchain.sycl_runtime.short_path),
+        )
+        args.add(
+            sycl_runtime_rpath,
+            format = "-rpath=$ORIGIN/%s",
+        )
 
     # Additional system libraries.
     args.add("-lm")  # Math.

diff --git a/ll/attributes.bzl b/ll/attributes.bzl
@@ -35,12 +35,14 @@ DEFAULT_ATTRS = {
         default = "cpp",
         # TODO: hip_amd, sycl_amd
         values = [
+            "bootstrap",
             "cpp",
-            "omp_cpu",
             "cuda_nvptx",
             "hip_amdgpu",
             "hip_nvptx",
-            "bootstrap",
+            "omp_cpu",
+            "sycl_cpu",
+            "sycl_amdgpu",
         ],
     ),
     "compile_flags": attr.string_list(
@@ -496,9 +498,33 @@ LL_TOOLCHAIN_ATTRS = {
     ),
     "rocm_device_libs": attr.label(
         doc = "The ROCm-Device-Libs.",
-        # default "@rocm-device-libs//:rocm-device-libs",
+        # default = "@rocm-device-libs//:rocm-device-libs",
         cfg = transition_to_cpp,
     ),
+    "sycl_headers": attr.label(
+        doc = "The SYCL headers.",
+        # default = "@opensycl//:sycl_headers"
+    ),
+    "sycl_hip_backend": attr.label(
+        doc = "The AMDGPU backend SYCL.",
+        # default = "@opensycl//:librt-backend-hip",
+        allow_single_file = True,
+    ),
+    "sycl_omp_backend": attr.label(
+        doc = "The AMDGPU backend OpenMP/CPU.",
+        # default = "@opensycl//:librt-backend-omp",
+        allow_single_file = True,
+    ),
+    "sycl_plugin": attr.label(
+        doc = "The compiler plugin for OpenSYCL.",
+        # default = "@opensycl//:opensycl-clang",
+        allow_single_file = True,
+    ),
+    "sycl_runtime": attr.label(
+        doc = "The SYCL runtime.",
+        # default = "@opensycl//:libhipSYCL-rt",
+        allow_single_file = True,
+    ),
     "symbolizer": attr.label(
         doc = "The `llvm-symbolizer`.",
         cfg = transition_to_bootstrap,

diff --git a/ll/environment.bzl b/ll/environment.bzl
@@ -20,7 +20,7 @@ def compile_object_environment(ctx):
     config = ctx.attr.toolchain_configuration[BuildSettingInfo].value
     toolchain = ctx.toolchains["//ll:toolchain_type"]
 
-    if config in ["cpp", "omp_cpu"]:
+    if config in ["cpp", "omp_cpu", "sycl_cpu"]:
         return {
             "LINK": toolchain.bitcode_linker.path,
             "LLD": toolchain.linker.path,
@@ -30,7 +30,7 @@ def compile_object_environment(ctx):
                 toolchain.linker_executable.dirname,
             ]),
         }
-    elif config in ["cuda_nvptx", "hip_amdgpu", "hip_nvptx"]:
+    elif config in ["cuda_nvptx", "hip_amdgpu", "hip_nvptx", "sycl_amdgpu"]:
         return {
             "CLANG_OFFLOAD_BUNDLER": toolchain.offload_bundler.path,
             "LINK": toolchain.bitcode_linker.path,

diff --git a/ll/init.bzl b/ll/init.bzl
@@ -118,6 +118,16 @@ def _initialize_rules_ll_impl(_):
         patch_args = ["-p1"],
     )
 
+    http_archive(
+        name = "opensycl",
+        build_file = "@rules_ll//third-party-overlays:opensycl.BUILD.bazel",
+        sha256 = "a01f5633ca62664b163f504495855fd009d340eb38bbf05f69bc71bd9a208afc",
+        strip_prefix = "OpenSYCL-12fdcaedfa990ab58ddf8bce304fa8cf917e6182",
+        urls = [
+            "https://github.com/OpenSYCL/OpenSYCL/archive/12fdcaedfa990ab58ddf8bce304fa8cf917e6182.zip",
+        ],
+    )
+
     http_archive(
         name = "rocclr",
         build_file = "@rules_ll//third-party-overlays:rocclr.BUILD.bazel",