issue/158/feat: 支持天数编译

Signed-off-by: YdrMaster <ydrml@hotmail.com>

src/infiniop/devices/cuda/cuda_common.cu

@@ -68,8 +68,10 @@ cudnnDataType_t getCudnnDtype(infiniDtype_t dt) {
         return CUDNN_DATA_INT8;
     case INFINI_DTYPE_I32:
         return CUDNN_DATA_INT32;
+#ifndef ENABLE_ILUVATAR_CUDA_API
     case INFINI_DTYPE_I64:
         return CUDNN_DATA_INT64;
+#endif
     case INFINI_DTYPE_U8:
         return CUDNN_DATA_UINT8;
     default:

src/infiniop/ops/gemm/cuda/gemm_cuda.cu

@@ -43,12 +43,20 @@ infiniStatus_t Descriptor::calculate(
     void *stream) const {
 
     cudaDataType a_type, b_type, c_type;
+#ifdef ENABLE_ILUVATAR_CUDA_API
+    cudaDataType compute_type;
+#else
     cublasComputeType_t compute_type;
+#endif
 
     switch (_dtype) {
     case INFINI_DTYPE_F16:
         a_type = b_type = c_type = CUDA_R_16F;
+#ifdef ENABLE_ILUVATAR_CUDA_API
+        compute_type = CUDA_R_32F;
+#else
         compute_type = CUBLAS_COMPUTE_32F;
+#endif
         break;
     case INFINI_DTYPE_BF16:
         a_type = b_type = c_type = CUDA_R_16BF;
@@ -56,7 +64,9 @@ infiniStatus_t Descriptor::calculate(
         break;
     case INFINI_DTYPE_F32:
         a_type = b_type = c_type = CUDA_R_32F;
-#ifdef ENABLE_SUGON_CUDA_API
+#if defined ENABLE_ILUVATAR_CUDA_API
+        compute_type = CUDA_R_32F;
+#elif defined ENABLE_SUGON_CUDA_API
         compute_type = CUBLAS_COMPUTE_32F;
 #else
         compute_type = CUBLAS_COMPUTE_32F_FAST_TF32;

src/infiniop/ops/gemm/info.h

@@ -8,10 +8,7 @@
 
 namespace op::gemm {
 
-class BlasMatrix {
-    BlasMatrix() = default;
-
-public:
+struct BlasMatrix {
     size_t ndim;
     size_t batch;
     ptrdiff_t stride;

src/infinirt/cuda/infinirt_cuda.cu

@@ -38,8 +38,12 @@ infiniStatus_t streamSynchronize(infinirtStream_t stream) {
 }
 
 infiniStatus_t streamWaitEvent(infinirtStream_t stream, infinirtEvent_t event) {
+#ifndef ENABLE_ILUVATAR_CUDA_API
     CHECK_CUDART(cudaStreamWaitEvent((cudaStream_t)stream, (cudaEvent_t)event));
     return INFINI_STATUS_SUCCESS;
+#else
+    return INFINI_STATUS_NOT_IMPLEMENTED;
+#endif
 }
 
 infiniStatus_t eventCreate(infinirtEvent_t *event_ptr) {

xmake.lua

@@ -46,6 +46,19 @@ if has_config("nv-gpu") then
     includes("xmake/cuda.lua")
 end
 
+-- 天数智芯
+option("iluvatar-gpu")
+    set_default(false)
+    set_showmenu(true)
+    set_description("Whether to complie implementations for Iluvatar GPU")
+option_end()
+
+if has_config("iluvatar-gpu") then
+    add_defines("ENABLE_CUDA_API")
+    add_defines("ENABLE_ILUVATAR_CUDA_API")
+    includes("xmake/iluvatar.lua")
+end
+
 -- 寒武纪
 option("cambricon-mlu")
     set_default(false)
@@ -174,6 +187,9 @@ target("infinirt")
     if has_config("moore-gpu") then
         add_deps("infinirt-moore")
     end
+    if has_config("iluvatar-gpu") then
+        add_deps("infinirt-iluvatar")
+    end
     if has_config("kunlun-xpu") then
         add_deps("infinirt-kunlun")
     end
@@ -193,6 +209,9 @@ target("infiniop")
     if has_config("nv-gpu") then
         add_deps("infiniop-cuda")
     end
+    if has_config("iluvatar-gpu") then
+        add_deps("infiniop-iluvatar")
+    end
     if has_config("sugon-dcu") then
         local builddir = string.format(
             "build/%s/%s/%s",
@@ -248,7 +267,7 @@ target("infiniccl")
     if has_config("metax-gpu") then
         add_deps("infiniccl-metax")
     end
-    
+
     set_languages("cxx17")
 
     add_files("src/infiniccl/*.cc")

xmake/iluvatar.lua

+toolchain("iluvatar.toolchain")
+    set_toolset("cc"  , "clang"  )
+    set_toolset("cxx" , "clang++")
+    set_toolset("cu"  , "clang++")
+    set_toolset("culd", "clang++")
+    set_toolset("cu-ccbin", "$(env CXX)", "$(env CC)")
+toolchain_end()
+
+rule("iluvatar.env")
+    add_deps("cuda.env", {order = true})
+    after_load(function (target)
+        local old = target:get("syslinks")
+        local new = {}
+
+        for _, link in ipairs(old) do
+            if link ~= "cudadevrt" then
+                table.insert(new, link)
+            end
+        end
+
+        if #old > #new then
+            target:set("syslinks", new)
+            local log = "cudadevrt removed, syslinks = { "
+            for _, link in ipairs(new) do
+                log = log .. link .. ", "
+            end
+            log = log:sub(0, -3) .. " }"
+            print(log)
+        end
+    end)
+rule_end()
+
+target("infiniop-iluvatar")
+    set_kind("static")
+    add_deps("infini-utils")
+    on_install(function (target) end)
+
+    set_toolchains("iluvatar.toolchain")
+    add_rules("iluvatar.env")
+    set_values("cuda.rdc", false)
+
+    add_links("cublas", "cudnn")
+
+    set_warnings("all", "error")
+    add_cuflags("-fPIC", "-x", "ivcore", "-std=c++17", {force = true})
+    add_cuflags("-fPIC")
+    add_culdflags("-fPIC")
+    add_cxflags("-fPIC")
+
+    -- set_languages("cxx17") 天数似乎不能用这个配置
+    add_files("../src/infiniop/devices/cuda/*.cu", "../src/infiniop/ops/*/cuda/*.cu")
+target_end()
+
+target("infinirt-iluvatar")
+    set_kind("static")
+    add_deps("infini-utils")
+    on_install(function (target) end)
+
+    set_toolchains("iluvatar.toolchain")
+    add_rules("iluvatar.env")
+    set_values("cuda.rdc", false)
+
+    add_links("cudart")
+
+    set_warnings("all", "error")
+    add_cuflags("-fPIC", "-x", "ivcore", "-std=c++17", {force = true})
+    add_cuflags("-fPIC")
+    add_culdflags("-fPIC")
+    add_cxflags("-fPIC")
+
+    -- set_languages("cxx17") 天数似乎不能用这个配置
+    add_files("../src/infinirt/cuda/*.cu")
+target_end()