apache · moazreyad · Jul 12, 2026
diff --git a/.github/workflows/conda-gpu-build.yaml b/.github/workflows/conda-gpu-build.yaml
@@ -0,0 +1,51 @@
+#
+# Licensed to the Apache Software Foundation (ASF) under one or more
+# contributor license agreements.  See the NOTICE file distributed with
+# this work for additional information regarding copyright ownership.
+# The ASF licenses this file to You under the Apache License, Version 2.0
+# (the "License"); you may not use this file except in compliance with
+# the License.  You may obtain a copy of the License at
+#
+#     http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing, software
+# distributed under the License is distributed on an "AS IS" BASIS,
+# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+# See the License for the specific language governing permissions and
+# limitations under the License.
+#
+
+name: Conda GPU Build
+
+on:
+  push:
+  pull_request:
+  workflow_dispatch:
+
+jobs:
+  build-gpu-conda-package:
+    runs-on: ubuntu-latest
+
+    env:
+      CUDA: "12.2"
+      DIST: "OFF"
+      TEST_COMMAND: python -c "import singa; import singa.singa_wrap; print(singa.__version__)"
+
+    steps:
+      - uses: actions/checkout@v7
+
+      - name: setup-miniconda
+        uses: conda-incubator/setup-miniconda@v4.0.1
+        with:
+          auto-update-conda: true
+          channels: nvidia,conda-forge,nusdbsystem,defaults
+          channel-priority: flexible
+
+      - name: install-conda-build
+        shell: bash -el {0}
+        run: conda install -y conda-build
+
+      - name: build-gpu-package
+        shell: bash -el {0}
+        run: |
+          conda-build tool/conda/singa -c nvidia -c conda-forge -c nusdbsystem -c defaults --python 3.8 --no-anaconda-upload --output-folder conda-bld
diff --git a/CMakeLists.txt b/CMakeLists.txt
@@ -119,7 +119,7 @@ IF (USE_CUDA)
     #LIST(APPEND SINGA_LINKER_LIBS cnmem)
     SET(global_cuda_objs "")
     # add support cuda fp16
-    SET(CUDA_NVCC_FLAGS "${CUDA_NVCC_FLAGS} --gpu-architecture=compute_75")
+    SET(CUDA_NVCC_FLAGS "${CUDA_NVCC_FLAGS} --gpu-architecture=compute_75 --std=c++11")
 ENDIF()
 
 # TODO(wangwei) detect the ev lib

diff --git a/src/core/tensor/math_kernel.h b/src/core/tensor/math_kernel.h
@@ -24,6 +24,7 @@
 #include "singa/singa_config.h"
 #ifdef USE_CUDA
 
+#include <thrust/count.h>
 #include <thrust/execution_policy.h>
 #include <thrust/remove.h>
 #include <thrust/sort.h>

diff --git a/src/model/layer/cudnn_convolution.cc b/src/model/layer/cudnn_convolution.cc
@@ -19,6 +19,7 @@
 #ifdef USE_CUDNN
 #include <cudnn.h>
 #include <chrono>
+#include <limits>
 #include "./cudnn_utils.h"
 #include "singa/utils/logging.h"
 
@@ -88,6 +89,62 @@ void CudnnConvolution::InitCudnn(const Tensor &input) {
                                          channels_, kernel_h_, kernel_w_));
   if (prefer_ == "fastest" || prefer_ == "limited_workspace" ||
       prefer_ == "no_workspace") {
+#if CUDNN_MAJOR >= 8
+    const int max_algos = 16;
+    int num_fp_alg = 0, num_bp_filt_alg = 0, num_bp_data_alg = 0;
+    cudnnConvolutionFwdAlgoPerf_t fp_alg_perf[max_algos];
+    cudnnConvolutionBwdFilterAlgoPerf_t bp_filt_perf[max_algos];
+    cudnnConvolutionBwdDataAlgoPerf_t bp_data_perf[max_algos];
+    const size_t workspace_limit =
+        prefer_ == "fastest" ? std::numeric_limits<size_t>::max()
+                             : (prefer_ == "limited_workspace" ?
+                                workspace_byte_limit_ : 0);
+
+    CUDNN_CHECK(cudnnGetConvolutionForwardAlgorithm_v7(
+                  ctx->cudnn_handle, x_desc_, filter_desc_, conv_desc_, y_desc_,
+                  max_algos, &num_fp_alg, fp_alg_perf));
+    bool found_alg = false;
+    for (int i = 0; i < num_fp_alg; ++i) {
+      if (fp_alg_perf[i].status == CUDNN_STATUS_SUCCESS &&
+          fp_alg_perf[i].memory <= workspace_limit) {
+        fp_alg_ = fp_alg_perf[i].algo;
+        found_alg = true;
+        break;
+      }
+    }
+    CHECK(found_alg) << "No cuDNN forward convolution algorithm matched "
+                     << prefer_ << " workspace preference";
+
+    CUDNN_CHECK(cudnnGetConvolutionBackwardFilterAlgorithm_v7(
+                  ctx->cudnn_handle, x_desc_, y_desc_, conv_desc_, filter_desc_,
+                  max_algos, &num_bp_filt_alg, bp_filt_perf));
+    found_alg = false;
+    for (int i = 0; i < num_bp_filt_alg; ++i) {
+      if (bp_filt_perf[i].status == CUDNN_STATUS_SUCCESS &&
+          bp_filt_perf[i].memory <= workspace_limit) {
+        bp_filter_alg_ = bp_filt_perf[i].algo;
+        found_alg = true;
+        break;
+      }
+    }
+    CHECK(found_alg) << "No cuDNN backward-filter convolution algorithm matched "
+                     << prefer_ << " workspace preference";
+
+    CUDNN_CHECK(cudnnGetConvolutionBackwardDataAlgorithm_v7(
+                  ctx->cudnn_handle, filter_desc_, y_desc_, conv_desc_, x_desc_,
+                  max_algos, &num_bp_data_alg, bp_data_perf));
+    found_alg = false;
+    for (int i = 0; i < num_bp_data_alg; ++i) {
+      if (bp_data_perf[i].status == CUDNN_STATUS_SUCCESS &&
+          bp_data_perf[i].memory <= workspace_limit) {
+        bp_data_alg_ = bp_data_perf[i].algo;
+        found_alg = true;
+        break;
+      }
+    }
+    CHECK(found_alg) << "No cuDNN backward-data convolution algorithm matched "
+                     << prefer_ << " workspace preference";
+#else
     cudnnConvolutionFwdPreference_t fwd_pref;
     cudnnConvolutionBwdFilterPreference_t bwd_filt_pref;
     cudnnConvolutionBwdDataPreference_t bwd_data_pref;
@@ -110,10 +167,10 @@ void CudnnConvolution::InitCudnn(const Tensor &input) {
     CUDNN_CHECK(cudnnGetConvolutionBackwardFilterAlgorithm(
                   ctx->cudnn_handle, x_desc_, y_desc_, conv_desc_, filter_desc_,
                   bwd_filt_pref, workspace_byte_limit_, &bp_filter_alg_));
-    // deprecated in cudnn v7
     CUDNN_CHECK(cudnnGetConvolutionBackwardDataAlgorithm(
                   ctx->cudnn_handle, filter_desc_, y_desc_, conv_desc_, x_desc_,
                   bwd_data_pref, workspace_byte_limit_, &bp_data_alg_));
+#endif  // CUDNN_MAJOR >= 8
   } else if (prefer_ == "autotune") {
     const int topk = 1;
     int num_fp_alg, num_bp_filt_alg, num_bp_data_alg;

diff --git a/src/model/layer/cudnn_rnn.cc b/src/model/layer/cudnn_rnn.cc
@@ -144,7 +144,14 @@ void CudnnRNN::SetRNNDescriptor(shared_ptr<Device> dev) {
     rnn_mode = CUDNN_RNN_TANH;
   else if (rnn_mode_ == "gru")
     rnn_mode = CUDNN_GRU;
-#if CUDNN_MAJOR <= 5
+#if CUDNN_MAJOR >= 8
+  CUDNN_CHECK(cudnnSetRNNDescriptor_v8(
+      rnn_desc_, CUDNN_RNN_ALGO_STANDARD, rnn_mode, CUDNN_RNN_DOUBLE_BIAS,
+      direction, input_mode, dtype_, dtype_, CUDNN_DEFAULT_MATH,
+      static_cast<int32_t>(input_size_), static_cast<int32_t>(hidden_size_),
+      static_cast<int32_t>(hidden_size_), static_cast<int32_t>(num_stacks_),
+      dropout_desc_, 0));
+#elif CUDNN_MAJOR <= 5
   CUDNN_CHECK(cudnnSetRNNDescriptor(rnn_desc_, hidden_size_, num_stacks_,
                                     dropout_desc_, input_mode, direction,
                                     rnn_mode, dtype_));

diff --git a/src/model/operation/convolution.cc b/src/model/operation/convolution.cc
@@ -23,6 +23,7 @@
 #include "convolution.h"
 
 #include <cctype>
+#include <limits>
 
 namespace singa {
 
@@ -502,6 +503,62 @@ CudnnConvHandle::CudnnConvHandle(
     bp_data_alg = CUDNN_CONVOLUTION_BWD_DATA_ALGO_1;
   } else if (prefer == "fastest" || prefer == "limited_workspace" ||
       prefer == "no_workspace") {
+#if CUDNN_MAJOR >= 8
+    const int max_algos = 16;
+    int num_fp_alg = 0, num_bp_filt_alg = 0, num_bp_data_alg = 0;
+    cudnnConvolutionFwdAlgoPerf_t fp_alg_perf[max_algos];
+    cudnnConvolutionBwdFilterAlgoPerf_t bp_filt_perf[max_algos];
+    cudnnConvolutionBwdDataAlgoPerf_t bp_data_perf[max_algos];
+    const size_t workspace_limit =
+        prefer == "fastest" ? std::numeric_limits<size_t>::max()
+                            : (prefer == "limited_workspace" ?
+                               workspace_byte_limit : 0);
+
+    CUDNN_CHECK(cudnnGetConvolutionForwardAlgorithm_v7(
+        ctx->cudnn_handle, x_desc, filter_desc, conv_desc, y_desc, max_algos,
+        &num_fp_alg, fp_alg_perf));
+    bool found_alg = false;
+    for (int i = 0; i < num_fp_alg; ++i) {
+      if (fp_alg_perf[i].status == CUDNN_STATUS_SUCCESS &&
+          fp_alg_perf[i].memory <= workspace_limit) {
+        fp_alg = fp_alg_perf[i].algo;
+        found_alg = true;
+        break;
+      }
+    }
+    CHECK(found_alg) << "No cuDNN forward convolution algorithm matched "
+                     << prefer << " workspace preference";
+
+    CUDNN_CHECK(cudnnGetConvolutionBackwardFilterAlgorithm_v7(
+        ctx->cudnn_handle, x_desc, y_desc, conv_desc, filter_desc, max_algos,
+        &num_bp_filt_alg, bp_filt_perf));
+    found_alg = false;
+    for (int i = 0; i < num_bp_filt_alg; ++i) {
+      if (bp_filt_perf[i].status == CUDNN_STATUS_SUCCESS &&
+          bp_filt_perf[i].memory <= workspace_limit) {
+        bp_filter_alg = bp_filt_perf[i].algo;
+        found_alg = true;
+        break;
+      }
+    }
+    CHECK(found_alg) << "No cuDNN backward-filter convolution algorithm matched "
+                     << prefer << " workspace preference";
+
+    CUDNN_CHECK(cudnnGetConvolutionBackwardDataAlgorithm_v7(
+        ctx->cudnn_handle, filter_desc, y_desc, conv_desc, x_desc, max_algos,
+        &num_bp_data_alg, bp_data_perf));
+    found_alg = false;
+    for (int i = 0; i < num_bp_data_alg; ++i) {
+      if (bp_data_perf[i].status == CUDNN_STATUS_SUCCESS &&
+          bp_data_perf[i].memory <= workspace_limit) {
+        bp_data_alg = bp_data_perf[i].algo;
+        found_alg = true;
+        break;
+      }
+    }
+    CHECK(found_alg) << "No cuDNN backward-data convolution algorithm matched "
+                     << prefer << " workspace preference";
+#else
     cudnnConvolutionFwdPreference_t fwd_pref;
     cudnnConvolutionBwdFilterPreference_t bwd_filt_pref;
     cudnnConvolutionBwdDataPreference_t bwd_data_pref;
@@ -524,10 +581,10 @@ CudnnConvHandle::CudnnConvHandle(
     CUDNN_CHECK(cudnnGetConvolutionBackwardFilterAlgorithm(
         ctx->cudnn_handle, x_desc, y_desc, conv_desc, filter_desc,
         bwd_filt_pref, workspace_byte_limit, &bp_filter_alg));
-    // deprecated in cudnn v7
     CUDNN_CHECK(cudnnGetConvolutionBackwardDataAlgorithm(
         ctx->cudnn_handle, filter_desc, y_desc, conv_desc, x_desc,
         bwd_data_pref, workspace_byte_limit, &bp_data_alg));
+#endif  // CUDNN_MAJOR >= 8
   } else if (prefer == "autotune") {
     const int topk = 1;
     int num_fp_alg, num_bp_filt_alg, num_bp_data_alg;

diff --git a/src/model/operation/rnn.cc b/src/model/operation/rnn.cc
@@ -111,12 +111,23 @@ void CudnnRNNHandle::init_rnn_desc() {
     RNNMode = CUDNN_LSTM;
   else if (mode == 3)
     RNNMode = CUDNN_GRU;
+#if CUDNN_MAJOR >= 8
+  CUDNN_CHECK(cudnnSetRNNDescriptor_v8(
+      rnnDesc, cudnnRNNAlgo, RNNMode,
+      bias ? CUDNN_RNN_DOUBLE_BIAS : CUDNN_RNN_NO_BIAS,
+      bidirectional ? CUDNN_BIDIRECTIONAL : CUDNN_UNIDIRECTIONAL,
+      CUDNN_LINEAR_INPUT, cudnnDataType, cudnnDataType, CUDNN_DEFAULT_MATH,
+      static_cast<int32_t>(feature_size), static_cast<int32_t>(hidden_size),
+      static_cast<int32_t>(hidden_size), static_cast<int32_t>(num_layers),
+      dropoutDesc, 0));
+#else
   CUDNN_CHECK(cudnnSetRNNDescriptor(
       ctx->cudnn_handle, rnnDesc, hidden_size, num_layers, dropoutDesc,
       CUDNN_LINEAR_INPUT,
       bidirectional ? CUDNN_BIDIRECTIONAL : CUDNN_UNIDIRECTIONAL, RNNMode,
       cudnnRNNAlgo,  // CUDNN_RNN_ALGO_STANDARD,
       cudnnDataType));
+#endif  // CUDNN_MAJOR >= 8
 }
 void CudnnRNNHandle::init_dropout_desc() {
   /* drop out */

diff --git a/tool/conda/singa/README.md b/tool/conda/singa/README.md
@@ -23,10 +23,10 @@
 
 ## Environment variables
 
-We export the CUDA version if SINGA is compiled with CUDA enabled. The cuDNN version is fixed by SINGA and cuDNN is installed from [anaconda cloud](https://anaconda.org/anaconda/cudnn).
+We export the CUDA version if SINGA is compiled with CUDA enabled. The value must be in `x.y` form and must match a CUDA variant in `conda_build_config.yaml`. The CUDA and cuDNN packages are installed into the conda-build environment from the configured conda channels.
 
-    # for SINGA with GPU, e.g. cuda9.0-cudnn7.3.1
-    export CUDA=9.0
+    # for SINGA with GPU, e.g. cuda12.2-cudnn8.9.2
+    export CUDA=12.2
 
 Then, we export a flag DIST to indicate if SINGA is compiled with distributed training enabled.
 
@@ -39,14 +39,15 @@ We need to export both CUDA and DIST for GPU version. For CPU-only version, we d
 
 After exporting the environment variables, we need to add the necessary conda channels
 
+    conda config --add channels nvidia
     conda config --add channels conda-forge
     conda config --add channels nusdbsystem
 
 Then, we can execute the following commands to compile SINGA and package it
 
     conda-build .  --python 3.6
 
-You will see the package path from the screen output, e.g., `xx/yy/singa-1.2.0-cpu.tar.bz2` or `xx/yy/singa-1.2.0-cudnn7.3.1_cuda9.0.tar.bz2`.
+You will see the package path from the screen output, e.g., `xx/yy/singa-1.2.0-cpu.tar.bz2` or `xx/yy/singa-1.2.0-cudnn8.9.2_cuda12.2.tar.bz2`.
 
 To clean the cache
 

diff --git a/tool/conda/singa/conda_build_config.yaml b/tool/conda/singa/conda_build_config.yaml
@@ -18,16 +18,25 @@
 #
 
 c_compiler_version:         # [linux]
-    - 5.4                   # [linux]
+    - 11                    # [linux]
 cxx_compiler_version:       # [linux]
-    - 5.4                   # [linux]
+    - 11                    # [linux]
 # https://docs.conda.io/projects/conda-build/en/latest/resources/compiler-tools.html#macos-sdk
 CONDA_BUILD_SYSROOT:
     - "/Applications/Xcode_11.7.app/Contents/Developer/Platforms/MacOSX.platform/Developer/SDKs/MacOSX.sdk" # [osx]
 cudnn:                      # [linux]
+    - "8.9.2.26 cuda12_0"   # [environ.get("CUDA")=="12.2"]
     - "7.6.5 cuda10.2_0"    # [environ.get("CUDA")=="10.2"]
 #    - "7.6.5 cuda10.0_0"    # [environ.get("CUDA")=="10.0"]
 #    - "7.6.5 cuda9.0_0"     # [environ.get("CUDA")=="9.0"]
+cuda_nvcc:                  # [linux]
+    - 12.2.140              # [environ.get("CUDA")=="12.2"]
+cuda_libraries:             # [linux]
+    - 12.2.2                # [environ.get("CUDA")=="12.2"]
+cuda_cudart:                # [linux]
+    - 12.2.140              # [environ.get("CUDA")=="12.2"]
+cuda_cccl:                  # [linux]
+    - 12.2.140              # [environ.get("CUDA")=="12.2"]
 dnnl:
     - 1.1
 python:
@@ -40,10 +49,11 @@ nccl:
 mpich:
     - 3.3.2
 build_str:
-    - "cudnn7.6.5_cuda10.2"   # [environ.get("CUDA")=="10.2"] && [environ.get("DIST")=="OFF"]
-    - "cudnn7.6.5_cuda10.0"   # [environ.get("CUDA")=="10.0"] && [environ.get("DIST")=="OFF"]
-    - "cudnn7.6.5_cuda9.0"    # [environ.get("CUDA")=="9.0"] && [environ.get("DIST")=="OFF"]
+    - "cudnn8.9.2_cuda12.2"  # [environ.get("CUDA")=="12.2" and environ.get("DIST")=="OFF"]
+    - "cudnn7.6.5_cuda10.2"   # [environ.get("CUDA")=="10.2" and environ.get("DIST")=="OFF"]
+    - "cudnn7.6.5_cuda10.0"   # [environ.get("CUDA")=="10.0" and environ.get("DIST")=="OFF"]
+    - "cudnn7.6.5_cuda9.0"    # [environ.get("CUDA")=="9.0" and environ.get("DIST")=="OFF"]
     - "cpu"                   # [environ.get("CUDA", "")== ""]
-    - "cudnn7.6.5_cuda10.2_nccl2.6.4.1_mpich3.3.2"     # [environ.get("CUDA")=="10.2"] && [environ.get("DIST")=="ON"]
-    - "cudnn7.6.5_cuda10.0_nccl2.4.8.1_mpich3.3.2"     # [environ.get("CUDA")=="10.0"] && [environ.get("DIST")=="ON"]
-    - "cudnn7.6.5_cuda9.0_nccl2.4.8.1_mpich3.3.2"      # [environ.get("CUDA")=="9.0"] && [environ.get("DIST")=="ON"]
+    - "cudnn7.6.5_cuda10.2_nccl2.6.4.1_mpich3.3.2"     # [environ.get("CUDA")=="10.2" and environ.get("DIST")=="ON"]
+    - "cudnn7.6.5_cuda10.0_nccl2.4.8.1_mpich3.3.2"     # [environ.get("CUDA")=="10.0" and environ.get("DIST")=="ON"]
+    - "cudnn7.6.5_cuda9.0_nccl2.4.8.1_mpich3.3.2"      # [environ.get("CUDA")=="9.0" and environ.get("DIST")=="ON"]
diff --git a/tool/conda/singa/meta.yaml b/tool/conda/singa/meta.yaml
@@ -41,7 +41,7 @@ requirements:
   build:
     - {{ compiler('cxx') }}
     - {{ compiler('c') }}
-    - cmake >=3.12.2
+    - cmake >=3.12.2,<4
     - make # [unix]
 
   host:
@@ -53,6 +53,11 @@ requirements:
     - numpy >=1.16,<2.0
     - pytest
     - deprecated 1.2.7
+    - cuda-nvcc {{ cuda_nvcc }}       # ['cuda' in str(build_str)]
+    - cuda-cccl {{ cuda_cccl }}       # ['cuda' in str(build_str)]
+    - cuda-cudart {{ cuda_cudart }}       # ['cuda' in str(build_str)]
+    - cuda-cudart-dev {{ cuda_cudart }}       # ['cuda' in str(build_str)]
+    - cuda-libraries-dev {{ cuda_libraries }}       # ['cuda' in str(build_str)]
     - cudnn {{ cudnn }}       # ['cudnn' in str(build_str)]
     - dnnl {{ dnnl }}
     - python {{ python }}
@@ -63,6 +68,8 @@ requirements:
     - {{ pin_compatible('glog', max_pin='x.x') }}
     - {{ pin_compatible('numpy', max_pin='x.x') }}
     - {{ pin_compatible('dnnl', max_pin='x.x') }}
+    - cuda-cudart {{ cuda_cudart }}       # ['cuda' in str(build_str)]
+    - cuda-libraries {{ cuda_libraries }}       # ['cuda' in str(build_str)]
     - cudnn {{ cudnn }}       # ['cudnn' in str(build_str)]
     - python {{ python }}
     - nccl {{ nccl }}         # ['nccl' in str(build_str)]