Add SSCP MUSA backend

CMakeLists.txt

@@ -102,6 +102,8 @@ endif()
 list(INSERT CMAKE_MODULE_PATH 0 "${PROJECT_SOURCE_DIR}/cmake/")
 find_package(CUDA QUIET)
 find_package(HIP QUIET HINTS ${ROCM_PATH} ${ROCM_PATH}/lib/cmake)
+list(APPEND CMAKE_MODULE_PATH /usr/local/musa/cmake)
+find_package(MUSA QUIET)
 
 # Check for OpenCL unless user has explicitly set WITH_SSCP_COMPILER to False
 if (DEFINED WITH_SSCP_COMPILER)
@@ -145,6 +147,11 @@ if(WITH_ROCM_BACKEND)
     # Let's assume the user knows what he/she is doing.
   endif()
 endif()
+if(WITH_MUSA_BACKEND)
+  if(NOT MUSA_FOUND)
+    message(SEND_ERROR "MUSA was not found")
+  endif()
+endif()
 
 if(WITH_OPENCL_BACKEND)
   if(NOT OpenCL_FOUND)
@@ -154,6 +161,7 @@ endif()
 
 set(WITH_CUDA_BACKEND ${CUDA_FOUND} CACHE BOOL "Build hipSYCL support for NVIDIA GPUs with CUDA")
 set(WITH_ROCM_BACKEND ${ROCM_FOUND} CACHE BOOL "Build hipSYCL support for AMD GPUs with ROCm")
+set(WITH_MUSA_BACKEND ${MUSA_FOUND} CACHE BOOL "Build hipSYCL support for Moore Threads GPUs with MUSA")
 set(WITH_OPENCL_BACKEND ${OpenCL_FOUND} CACHE BOOL "Build hipSYCL support for OpenCL SPIR-V devices supporting USM")
 set(WITH_CUDA_NVCXX_ONLY FALSE CACHE BOOL "Whether to target CUDA exclusively with nvc++")
 
@@ -206,7 +214,7 @@ endif()
 if(WITH_ROCM_BACKEND)
   set(BUILD_CLANG_PLUGIN true)
 endif()
-if(WITH_OPENCL_BACKEND OR WITH_LEVEL_ZERO_BACKEND)
+if(WITH_OPENCL_BACKEND OR WITH_LEVEL_ZERO_BACKEND OR WITH_MUSA_BACKEND)
   set(BUILD_CLANG_PLUGIN true)
   set(WITH_SSCP_COMPILER true)
 endif()

include/hipSYCL/compiler/llvm-to-backend/musa/LLVMToMusa.hpp

@@ -0,0 +1,61 @@
+/*
+ * This file is part of hipSYCL, a SYCL implementation based on CUDA/HIP
+ *
+ * Copyright (c) 2019-2022 Aksel Alpay
+ * All rights reserved.
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions are met:
+ *
+ * 1. Redistributions of source code must retain the above copyright notice, this
+ *    list of conditions and the following disclaimer.
+ * 2. Redistributions in binary form must reproduce the above copyright notice,
+ *    this list of conditions and the following disclaimer in the documentation
+ *    and/or other materials provided with the distribution.
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
+ * ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
+ * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+ * DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR CONTRIBUTORS BE LIABLE FOR
+ * ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
+ * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
+ * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
+ * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
+ * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
+ * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ */
+
+#ifndef HIPSYCL_LLVM_TO_MUSA_HPP
+#define HIPSYCL_LLVM_TO_MUSA_HPP
+
+
+#include "../LLVMToBackend.hpp"
+
+#include <vector>
+#include <string>
+
+namespace hipsycl {
+namespace compiler {
+
+class LLVMToMusaTranslator : public LLVMToBackendTranslator{
+public:
+  LLVMToMusaTranslator(const std::vector<std::string>& KernelNames);
+
+  virtual ~LLVMToMusaTranslator() {}
+
+  virtual bool prepareBackendFlavor(llvm::Module& M) override {return true;}
+  virtual bool toBackendFlavor(llvm::Module &M, PassHandler& PH) override;
+  virtual bool translateToBackendFormat(llvm::Module &FlavoredModule, std::string &out) override;
+protected:
+  virtual bool applyBuildOption(const std::string &Option, const std::string &Value) override;
+  virtual bool isKernelAfterFlavoring(llvm::Function& F) override;
+  virtual AddressSpaceMap getAddressSpaceMap() const override;
+private:
+  std::vector<std::string> KernelNames;
+  unsigned MusaTarget = 10;
+};
+
+}
+}
+
+#endif

include/hipSYCL/compiler/llvm-to-backend/musa/LLVMToMusaFactory.hpp

@@ -0,0 +1,45 @@
+/*
+ * This file is part of hipSYCL, a SYCL implementation based on CUDA/HIP
+ *
+ * Copyright (c) 2019-2022 Aksel Alpay
+ * All rights reserved.
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions are met:
+ *
+ * 1. Redistributions of source code must retain the above copyright notice, this
+ *    list of conditions and the following disclaimer.
+ * 2. Redistributions in binary form must reproduce the above copyright notice,
+ *    this list of conditions and the following disclaimer in the documentation
+ *    and/or other materials provided with the distribution.
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
+ * ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
+ * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+ * DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR CONTRIBUTORS BE LIABLE FOR
+ * ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
+ * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
+ * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
+ * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
+ * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
+ * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ */
+
+#ifndef HIPSYCL_LLVM_TO_MUSA_FACTORY_HPP
+#define HIPSYCL_LLVM_TO_MUSA_FACTORY_HPP
+
+#include <memory>
+#include <vector>
+#include <string>
+#include "../LLVMToBackend.hpp"
+
+namespace hipsycl {
+namespace compiler {
+
+std::unique_ptr<LLVMToBackendTranslator>
+createLLVMToMusaTranslator(const std::vector<std::string> &KernelNames);
+
+}
+}
+
+#endif

include/hipSYCL/glue/backend_interop.hpp

@@ -63,5 +63,6 @@ template <sycl::backend b> struct backend_interop {
 #include "hip/hip_interop.hpp"
 #include "ze/ze_interop.hpp"
 #include "omp/omp_interop.hpp"
+#include "musa/musa_interop.hpp"
 
 #endif

include/hipSYCL/glue/kernel_configuration.hpp

@@ -70,7 +70,9 @@ enum class kernel_build_option : int {
   amdgpu_rocm_device_libs_path,
   amdgpu_rocm_path,
 
-  spirv_dynamic_local_mem_allocation_size
+  spirv_dynamic_local_mem_allocation_size,
+
+  musa_target_device
 };
 
 enum class kernel_build_flag : int {
@@ -97,7 +99,8 @@ class string_build_config_mapper {
       {"amdgpu-target-device", kernel_build_option::amdgpu_target_device},
       {"rocm-device-libs-path", kernel_build_option::amdgpu_rocm_device_libs_path},
       {"rocm-path", kernel_build_option::amdgpu_rocm_path},
-      {"spirv-dynamic-local-mem-allocation-size", kernel_build_option::spirv_dynamic_local_mem_allocation_size}
+      {"spirv-dynamic-local-mem-allocation-size", kernel_build_option::spirv_dynamic_local_mem_allocation_size},
+      {"musa-target-device", kernel_build_option::musa_target_device}
     };
 
     _flags = {

include/hipSYCL/glue/llvm-sscp/s2_ir_constants.hpp

@@ -58,6 +58,7 @@ inline constexpr int spirv = 0;
 inline constexpr int ptx = 1;
 inline constexpr int amdgpu = 2;
 inline constexpr int host = 3;
+inline constexpr int musa = 4;
 
 }
 

include/hipSYCL/glue/musa/musa_interop.hpp

@@ -0,0 +1,135 @@
+/*
+ * This file is part of hipSYCL, a SYCL implementation based on CUDA/HIP
+ *
+ * Copyright (c) 2020 Aksel Alpay
+ * All rights reserved.
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions are met:
+ *
+ * 1. Redistributions of source code must retain the above copyright notice, this
+ *    list of conditions and the following disclaimer.
+ * 2. Redistributions in binary form must reproduce the above copyright notice,
+ *    this list of conditions and the following disclaimer in the documentation
+ *    and/or other materials provided with the distribution.
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
+ * ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
+ * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+ * DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR CONTRIBUTORS BE LIABLE FOR
+ * ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
+ * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
+ * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
+ * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
+ * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
+ * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ */
+
+
+
+#include "hipSYCL/sycl/backend.hpp"
+#include "hipSYCL/sycl/device.hpp"
+
+#ifdef SYCL_EXT_HIPSYCL_BACKEND_MUSA
+#include "hipSYCL/runtime/musa/musa_queue.hpp"
+#include "hipSYCL/runtime/error.hpp"
+
+#ifndef HIPSYCL_GLUE_MUSA_BACKEND_INTEROP_HPP
+#define HIPSYCL_GLUE_MUSA_BACKEND_INTEROP_HPP
+
+struct MUstream_st;
+
+namespace hipsycl {
+namespace glue {
+
+template <> struct backend_interop<sycl::backend::musa> {
+  // Well, there's not a really a native error code type
+  using error_type = int;
+
+  using native_mem_type = void *;
+  using native_device_type = int;
+  using native_queue_type = MUstream_st*;
+
+  template <class Accessor_type>
+  static native_mem_type get_native_mem(const Accessor_type &a) {
+    return a.get_pointer();
+  }
+
+  static native_device_type get_native_device(const sycl::device &d) {
+    return sycl::detail::extract_rt_device(d).get_id();
+  }
+
+  static native_queue_type
+  get_native_queue(void *launcher_params) {
+
+    if (!launcher_params) {
+      rt::register_error(
+          __hipsycl_here(),
+          rt::error_info{"Invalid argument to get_native_queue()"});
+
+      return native_queue_type{};
+    }
+
+    rt::inorder_queue* q = static_cast<rt::inorder_queue*>(launcher_params);
+    return static_cast<native_queue_type>(q->get_native_type());
+  }
+
+  static native_queue_type
+  get_native_queue(rt::device_id dev, rt::backend_executor *executor) {
+    rt::multi_queue_executor *mqe =
+        dynamic_cast<rt::multi_queue_executor *>(executor);
+
+    if (!mqe) {
+      rt::register_error(
+          __hipsycl_here(),
+          rt::error_info{"Invalid argument to get_native_queue()"});
+      return native_queue_type{};
+    }
+
+    rt::inorder_queue *q = nullptr;
+    mqe->for_each_queue(
+        dev, [&](rt::inorder_queue *current_queue) { q = current_queue; });
+    assert(q);
+
+    return static_cast<native_queue_type>(q->get_native_type());
+  }
+
+  static sycl::device make_sycl_device(int device_id) {
+    return sycl::device{
+        rt::device_id{rt::backend_descriptor{rt::hardware_platform::musa,
+                                             rt::api_platform::musa},
+                      device_id}};
+  }
+
+  static constexpr bool can_make_platform = false;
+  static constexpr bool can_make_device = true;
+  static constexpr bool can_make_context = false;
+  static constexpr bool can_make_queue = false;
+  static constexpr bool can_make_event = false;
+  static constexpr bool can_make_buffer = false;
+  static constexpr bool can_make_sampled_image = false;
+  static constexpr bool can_make_image_sampler = false;
+  static constexpr bool can_make_stream = false;
+  static constexpr bool can_make_kernel = false;
+  static constexpr bool can_make_module = false;
+
+  static constexpr bool can_extract_native_platform = false;
+  static constexpr bool can_extract_native_device = true;
+  static constexpr bool can_extract_native_context = false;
+  static constexpr bool can_extract_native_queue = false;
+  static constexpr bool can_extract_native_event = false;
+  static constexpr bool can_extract_native_buffer = false;
+  static constexpr bool can_extract_native_sampled_image = false;
+  static constexpr bool can_extract_native_image_sampler = false;
+  static constexpr bool can_extract_native_stream = false;
+  static constexpr bool can_extract_native_kernel = false;
+  static constexpr bool can_extract_native_module = false;
+  static constexpr bool can_extract_native_device_event = false;
+  static constexpr bool can_extract_native_mem = true;
+};
+
+}
+}
+
+#endif
+#endif

include/hipSYCL/runtime/device_id.hpp

@@ -40,6 +40,7 @@ enum class hardware_platform
   rocm,
   cuda,
   level_zero,
+  musa,
   ocl,
   cpu
 };
@@ -48,6 +49,7 @@ enum class api_platform {
   cuda,
   hip,
   level_zero,
+  musa,
   ocl,
   omp
 };
@@ -56,6 +58,7 @@ enum class backend_id {
   cuda,
   hip,
   level_zero,
+  musa,
   ocl,
   omp
 };
@@ -81,6 +84,9 @@ struct backend_descriptor
     else if(hw_plat == hardware_platform::level_zero &&
             sw_plat == api_platform::level_zero)
       id = backend_id::level_zero;
+    else if (hw_plat == hardware_platform::musa &&
+             sw_plat == api_platform::musa)
+      id = backend_id::musa;
     else if (hw_plat == hardware_platform::ocl && sw_plat == api_platform::ocl)
       id = backend_id::ocl;
     else