From 45ee32a2c6e8ace872cf669a61a1316122eae924 Mon Sep 17 00:00:00 2001 From: Hugo Date: Tue, 14 Oct 2025 10:23:59 +0800 Subject: [PATCH 01/15] feat: support bazel build (#1) 1. Install bazel with [bazelisk](https://github.com/bazelbuild/bazelisk) 2. [Set alias for bazel `alias bazel='bazelisk'` in `.bashrc` or other shell config file on linux system] 3. RUN following commands for test (haven't test on windows) >NOTE : change `--cuda_archs` to your own case ``` bash bazel run //test:muda_unit_test --cuda_archs=compute_89:compute_89,sm_89 -c opt bazel run //test:muda_linear_sysytem_test --cuda_archs=compute_89:compute_89,sm_89 -c opt bazel run //test:muda_eigen_test --cuda_archs=compute_89:compute_89,sm_89 -c opt bazel run //example:examples --cuda_archs=compute_89:compute_89,sm_89 -c opt ``` To enable `muda_check`, `muda_compute_graph`, `muda_nvtx3`, add flags `--//:muda_check` , `--//:muda_compute_graph` or `--//:muda_nvtx3` repectively when build ```bash bazel build --//:muda_check --//:muda_compute_graph --//:muda_nvtx3 //... ``` ## Changes 1. Add MODULE.bazel, BUILD.bazel etc files for bazel build 2. Change dir `external` to `third_party` as `external` is reserved by Bazel and needed for hedron_compile_commands --- .bazelrc | 30 +++++++ .bazelversion | 1 + .gitignore | 14 ++++ BUILD.bazel | 74 ++++++++++++++++ MODULE.bazel | 36 ++++++++ example/BUILD.bazel | 23 +++++ test/BUILD.bazel | 84 +++++++++++++++++++ third_party/BUILD.bazel | 11 +++ {external => third_party}/catch2/catch.hpp | 0 .../catch2/catch_reporter_automake.hpp | 0 .../catch2/catch_reporter_sonarqube.hpp | 0 .../catch2/catch_reporter_tap.hpp | 0 .../catch2/catch_reporter_teamcity.hpp | 0 13 files changed, 273 insertions(+) create mode 100644 .bazelrc create mode 100644 .bazelversion create mode 100644 BUILD.bazel create mode 100644 MODULE.bazel create mode 100644 example/BUILD.bazel create mode 100644 test/BUILD.bazel create mode 100644 third_party/BUILD.bazel rename {external => third_party}/catch2/catch.hpp (100%) rename {external => third_party}/catch2/catch_reporter_automake.hpp (100%) rename {external => third_party}/catch2/catch_reporter_sonarqube.hpp (100%) rename {external => third_party}/catch2/catch_reporter_tap.hpp (100%) rename {external => third_party}/catch2/catch_reporter_teamcity.hpp (100%) diff --git a/.bazelrc b/.bazelrc new file mode 100644 index 00000000..201e7c7e --- /dev/null +++ b/.bazelrc @@ -0,0 +1,30 @@ +# +------------------------------------------------------------+ +# | Build Configurations | +# +------------------------------------------------------------+ +# Do not show warnings from external dependencies. +build --output_filter="^//" + +build --show_timestamps + +build --flag_alias=cuda_archs=@rules_cuda//cuda:archs + +# Work around the sandbox issue. +build --spawn_strategy=local + + +# Consider warning as error. +# inference is not included because tensorflow uses the same log macros with glog and it will +# cause re-definition warning. +build --per_file_copt="-\\.$@-Wall,-Wextra,-Werror" +build --per_file_copt="-\\.$,external/.*@-w" +build --per_file_copt="-\\.$,third_party/.*@-w" + +# If a command fails, print out the full command line. +build --verbose_failures + +# Enable C++17 +build --cxxopt="-std=c++17" +# Workaround as cpp toolchain: stop passing -std=c++0x per default · Issue #18181 · bazelbuild/bazel +build --host_cxxopt="-std=c++17" +# Enable colorful output of GCC +build --cxxopt="-fdiagnostics-color=always" diff --git a/.bazelversion b/.bazelversion new file mode 100644 index 00000000..f4131375 --- /dev/null +++ b/.bazelversion @@ -0,0 +1 @@ +8.4.1 \ No newline at end of file diff --git a/.gitignore b/.gitignore index a7004ca8..31c50d46 100644 --- a/.gitignore +++ b/.gitignore @@ -10,3 +10,17 @@ vsxmake2022 # CMake Cache CMakeBuild/ + +# Bazel build +bazel-* +MODULE.bazel.lock + +# JetBrains IDEs +.idea/ + +.cache + +compile_commands.json + +# //external is reserved by Bazel and needed for hedron_compile_commands +external/ \ No newline at end of file diff --git a/BUILD.bazel b/BUILD.bazel new file mode 100644 index 00000000..2a154ac9 --- /dev/null +++ b/BUILD.bazel @@ -0,0 +1,74 @@ +load("@bazel_skylib//rules:common_settings.bzl", "bool_flag") +load("@rules_cuda//cuda:defs.bzl", "cuda_library") +load("@rules_license//rules:license.bzl", "license") + +package( + default_applicable_licenses = [":license"], + default_visibility = ["//visibility:public"], +) + +license( + name = "license", + package_name = "muda", + package_url = "https://github.com/MuGdxy/muda", + license_text = "LICENSE", + license_kinds = ["@rules_license//licenses/spdx:Apache-2.0"], +) + +bool_flag( + name = "muda_check", + build_setting_default = False, +) + +config_setting( + name = "muda_check_on", + flag_values = {":muda_check": "True"}, +) + +bool_flag( + name = "muda_compute_graph", + build_setting_default = False, +) + +config_setting( + name = "muda_compute_graph_on", + flag_values = {":muda_compute_graph": "True"}, +) + +bool_flag( + name = "muda_nvtx3", + build_setting_default = False, +) + +config_setting( + name = "muda_nvtx3_on", + flag_values = {":muda_nvtx3": "True"}, +) + +cuda_library( + name = "muda", + hdrs = glob( + [ + "src/**/*.h", + "src/**/*.hpp", + "src/**/*.inl", + "src/**/*.cuh", + ], + allow_empty = True, + ), + copts = [ + "--expt-relaxed-constexpr", + "--extended-lambda", + ], + defines = select({ + ":muda_check_on": ["MUDA_CHECK_ON"], + "//conditions:default": ["MUDA_CHECK_ON=0"], + }) + select({ + ":muda_compute_graph_on": ["MUDA_COMPUTE_GRAPH_ON"], + "//conditions:default": [], + }) + select({ + ":muda_nvtx3_on": ["MUDA_NVTX3_ON"], + "//conditions:default": [], + }), + includes = ["src"], +) diff --git a/MODULE.bazel b/MODULE.bazel new file mode 100644 index 00000000..96e611f0 --- /dev/null +++ b/MODULE.bazel @@ -0,0 +1,36 @@ +"""muda""" + +module( + name = "muda", + version = "0.1.0", + compatibility_level = 1, +) + +bazel_dep(name = "rules_cc", version = "0.2.8") + +bazel_dep(name = "rules_license", version = "1.0.0") + +# rules_cuda is required for CUDA toolchain support +bazel_dep(name = "rules_cuda", version = "0.2.4") + +bazel_dep(name = "eigen", version = "3.4.1") + +bazel_dep(name = "bazel_skylib", version = "1.8.2") + +# Optional for debug: Hedron's Compile Commands Extractor for Bazel +# https://github.com/hedronvision/bazel-compile-commands-extractor +bazel_dep(name = "hedron_compile_commands", dev_dependency = True) +git_override( + module_name = "hedron_compile_commands", + remote = "https://github.com/hedronvision/bazel-compile-commands-extractor.git", + commit = "abb61a688167623088f8768cc9264798df6a9d10", + # Replace the commit hash (above) with the latest (https://github.com/hedronvision/bazel-compile-commands-extractor/commits/main). + # Even better, set up Renovate and let it do the work for you (see "Suggestion: Updates" in the README). +) + +cuda = use_extension("@rules_cuda//cuda:extensions.bzl", "toolchain") +cuda.local_toolchain( + name = "local_cuda", + toolkit_path = "/usr/local/cuda", +) +use_repo(cuda, "local_cuda") \ No newline at end of file diff --git a/example/BUILD.bazel b/example/BUILD.bazel new file mode 100644 index 00000000..18214895 --- /dev/null +++ b/example/BUILD.bazel @@ -0,0 +1,23 @@ +load("@rules_cuda//cuda:defs.bzl", "cuda_binary") + +cuda_binary( + name = "examples", + srcs = glob([ + "**/*.cu", + "**/*.cpp", + ]), + hdrs = [ + "example_common.h", + ], + copts = [ + "-std=c++17", + ], + includes = ["."], + rdc = True, + visibility = ["//:__pkg__"], + deps = [ + "//:muda", + "//third_party:catch2", + "@eigen", + ], +) diff --git a/test/BUILD.bazel b/test/BUILD.bazel new file mode 100644 index 00000000..a421e592 --- /dev/null +++ b/test/BUILD.bazel @@ -0,0 +1,84 @@ +load("@rules_cuda//cuda:defs.bzl", "cuda_binary") + +cuda_binary( + name = "muda_unit_test", + srcs = glob([ + "unit_test/**/*.cu", + "unit_test/**/*.cpp", + ]), + copts = [ + "--expt-relaxed-constexpr", + "--extended-lambda", + "-std=c++20", + ] + select({ + "@rules_cc//cc/compiler:msvc-cl": ["/wd4819"], + "//conditions:default": [], + }), + rdc = True, + deps = [ + "//:muda", + "//third_party:catch2", + "@eigen", + ], +) + +cuda_binary( + name = "muda_eigen_test", + srcs = glob([ + "eigen_test/**/*.cu", + "eigen_test/**/*.cpp", + ]), + hdrs = glob( + [ + "eigen_test/**/*.h", + "eigen_test/**/*.hpp", + ], + allow_empty = True, + ), + copts = [ + "--expt-relaxed-constexpr", + "--extended-lambda", + "-std=c++20", + ] + select({ + "@rules_cc//cc/compiler:msvc-cl": ["/wd4819"], + "//conditions:default": [], + }), + rdc = True, + deps = [ + "//:muda", + "//third_party:catch2", + "@eigen", + ], +) + +cuda_binary( + name = "muda_linear_sysytem_test", + srcs = glob([ + "linear_system_test/**/*.cu", + "linear_system_test/**/*.cpp", + ]), + hdrs = glob( + [ + "linear_system_test/**/*.h", + "linear_system_test/**/*.hpp", + ], + allow_empty = True, + ), + copts = [ + "--expt-relaxed-constexpr", + "--extended-lambda", + "-std=c++20", + ] + select({ + "@rules_cc//cc/compiler:msvc-cl": ["/wd4819"], + "//conditions:default": [], + }), + rdc = True, + deps = [ + "//:muda", + "//third_party:catch2", + "@eigen", + "@local_cuda//:cublas", + "@local_cuda//:cusolver", + "@local_cuda//:cusparse", + ], +) diff --git a/third_party/BUILD.bazel b/third_party/BUILD.bazel new file mode 100644 index 00000000..b57b88fd --- /dev/null +++ b/third_party/BUILD.bazel @@ -0,0 +1,11 @@ +load("@rules_cc//cc:defs.bzl", "cc_library") + +package(default_visibility = ["//visibility:public"]) + +cc_library( + name = "catch2", + hdrs = glob([ + "catch2/**/*.hpp", + ]), + includes = ["."], +) diff --git a/external/catch2/catch.hpp b/third_party/catch2/catch.hpp similarity index 100% rename from external/catch2/catch.hpp rename to third_party/catch2/catch.hpp diff --git a/external/catch2/catch_reporter_automake.hpp b/third_party/catch2/catch_reporter_automake.hpp similarity index 100% rename from external/catch2/catch_reporter_automake.hpp rename to third_party/catch2/catch_reporter_automake.hpp diff --git a/external/catch2/catch_reporter_sonarqube.hpp b/third_party/catch2/catch_reporter_sonarqube.hpp similarity index 100% rename from external/catch2/catch_reporter_sonarqube.hpp rename to third_party/catch2/catch_reporter_sonarqube.hpp diff --git a/external/catch2/catch_reporter_tap.hpp b/third_party/catch2/catch_reporter_tap.hpp similarity index 100% rename from external/catch2/catch_reporter_tap.hpp rename to third_party/catch2/catch_reporter_tap.hpp diff --git a/external/catch2/catch_reporter_teamcity.hpp b/third_party/catch2/catch_reporter_teamcity.hpp similarity index 100% rename from external/catch2/catch_reporter_teamcity.hpp rename to third_party/catch2/catch_reporter_teamcity.hpp From 4051545d4712478c8d7d2f2f74007059d6c5b3ee Mon Sep 17 00:00:00 2001 From: zhiguo Date: Wed, 15 Oct 2025 15:55:06 +0800 Subject: [PATCH 02/15] use cuda_test and overwrite the rules_cuda by using the latest commit --- MODULE.bazel | 40 +++++++++++++++++++++++----------------- test/BUILD.bazel | 16 ++++++++-------- 2 files changed, 31 insertions(+), 25 deletions(-) diff --git a/MODULE.bazel b/MODULE.bazel index 96e611f0..aaad2761 100644 --- a/MODULE.bazel +++ b/MODULE.bazel @@ -10,27 +10,33 @@ bazel_dep(name = "rules_cc", version = "0.2.8") bazel_dep(name = "rules_license", version = "1.0.0") -# rules_cuda is required for CUDA toolchain support -bazel_dep(name = "rules_cuda", version = "0.2.4") - -bazel_dep(name = "eigen", version = "3.4.1") - +bazel_dep(name = "eigen", version = "3.4.1", dev_dependency = True) bazel_dep(name = "bazel_skylib", version = "1.8.2") -# Optional for debug: Hedron's Compile Commands Extractor for Bazel -# https://github.com/hedronvision/bazel-compile-commands-extractor -bazel_dep(name = "hedron_compile_commands", dev_dependency = True) +# rules_cuda is required for CUDA toolchain support +bazel_dep(name = "rules_cuda", version = "0.2.4") git_override( - module_name = "hedron_compile_commands", - remote = "https://github.com/hedronvision/bazel-compile-commands-extractor.git", - commit = "abb61a688167623088f8768cc9264798df6a9d10", - # Replace the commit hash (above) with the latest (https://github.com/hedronvision/bazel-compile-commands-extractor/commits/main). - # Even better, set up Renovate and let it do the work for you (see "Suggestion: Updates" in the README). + module_name = "rules_cuda", + remote = "https://github.com/bazel-contrib/rules_cuda.git", + commit = "1bf21ef1eee98042c0b2c5f39365ac39a4217c53", ) + + cuda = use_extension("@rules_cuda//cuda:extensions.bzl", "toolchain") -cuda.local_toolchain( - name = "local_cuda", - toolkit_path = "/usr/local/cuda", +cuda.toolkit( + name = "cuda", + toolkit_path = "/usr/local/cuda/", ) -use_repo(cuda, "local_cuda") \ No newline at end of file +use_repo(cuda, "cuda") + +# Optional for debug: Hedron's Compile Commands Extractor for Bazel +# https://github.com/hedronvision/bazel-compile-commands-extractor +# bazel_dep(name = "hedron_compile_commands", dev_dependency = True) +# git_override( +# module_name = "hedron_compile_commands", +# remote = "https://github.com/hedronvision/bazel-compile-commands-extractor.git", +# commit = "abb61a688167623088f8768cc9264798df6a9d10", +# # Replace the commit hash (above) with the latest (https://github.com/hedronvision/bazel-compile-commands-extractor/commits/main). +# # Even better, set up Renovate and let it do the work for you (see "Suggestion: Updates" in the README). +# ) \ No newline at end of file diff --git a/test/BUILD.bazel b/test/BUILD.bazel index a421e592..ad2541f3 100644 --- a/test/BUILD.bazel +++ b/test/BUILD.bazel @@ -1,6 +1,6 @@ -load("@rules_cuda//cuda:defs.bzl", "cuda_binary") +load("@rules_cuda//cuda:defs.bzl", "cuda_test") -cuda_binary( +cuda_test( name = "muda_unit_test", srcs = glob([ "unit_test/**/*.cu", @@ -22,7 +22,7 @@ cuda_binary( ], ) -cuda_binary( +cuda_test( name = "muda_eigen_test", srcs = glob([ "eigen_test/**/*.cu", @@ -51,8 +51,8 @@ cuda_binary( ], ) -cuda_binary( - name = "muda_linear_sysytem_test", +cuda_test( + name = "muda_linear_system_test", srcs = glob([ "linear_system_test/**/*.cu", "linear_system_test/**/*.cpp", @@ -77,8 +77,8 @@ cuda_binary( "//:muda", "//third_party:catch2", "@eigen", - "@local_cuda//:cublas", - "@local_cuda//:cusolver", - "@local_cuda//:cusparse", + "@cuda//:cublas", + "@cuda//:cusolver", + "@cuda//:cusparse", ], ) From cb1c9d0e70cea2b18b978783917bfb688d2dd0f3 Mon Sep 17 00:00:00 2001 From: zhiguo Date: Wed, 15 Oct 2025 16:04:12 +0800 Subject: [PATCH 03/15] use version 1.0.0-beta --- MODULE.bazel | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/MODULE.bazel b/MODULE.bazel index aaad2761..22a46b58 100644 --- a/MODULE.bazel +++ b/MODULE.bazel @@ -2,7 +2,7 @@ module( name = "muda", - version = "0.1.0", + version = "1.0.0-beta", compatibility_level = 1, ) From 923fb02774ec6fe0056553503783277bcb50624a Mon Sep 17 00:00:00 2001 From: zhiguo Date: Thu, 16 Oct 2025 10:59:56 +0800 Subject: [PATCH 04/15] update module name and usd cc_library for header only --- .gitignore | 4 +++- BUILD.bazel | 8 ++++---- MODULE.bazel | 4 ++-- 3 files changed, 9 insertions(+), 7 deletions(-) diff --git a/.gitignore b/.gitignore index 31c50d46..c2bec551 100644 --- a/.gitignore +++ b/.gitignore @@ -23,4 +23,6 @@ MODULE.bazel.lock compile_commands.json # //external is reserved by Bazel and needed for hedron_compile_commands -external/ \ No newline at end of file +external/ + +presubmit.yml \ No newline at end of file diff --git a/BUILD.bazel b/BUILD.bazel index 2a154ac9..b7272b15 100644 --- a/BUILD.bazel +++ b/BUILD.bazel @@ -1,5 +1,5 @@ load("@bazel_skylib//rules:common_settings.bzl", "bool_flag") -load("@rules_cuda//cuda:defs.bzl", "cuda_library") +load("@rules_cc//cc:defs.bzl", "cc_library") load("@rules_license//rules:license.bzl", "license") package( @@ -9,8 +9,8 @@ package( license( name = "license", - package_name = "muda", - package_url = "https://github.com/MuGdxy/muda", + package_name = "hugooole_muda", + package_url = "https://github.com/hugooole/muda", license_text = "LICENSE", license_kinds = ["@rules_license//licenses/spdx:Apache-2.0"], ) @@ -45,7 +45,7 @@ config_setting( flag_values = {":muda_nvtx3": "True"}, ) -cuda_library( +cc_library( name = "muda", hdrs = glob( [ diff --git a/MODULE.bazel b/MODULE.bazel index 22a46b58..9a5b5602 100644 --- a/MODULE.bazel +++ b/MODULE.bazel @@ -1,7 +1,7 @@ """muda""" module( - name = "muda", + name = "hugooole_muda", version = "1.0.0-beta", compatibility_level = 1, ) @@ -26,7 +26,7 @@ git_override( cuda = use_extension("@rules_cuda//cuda:extensions.bzl", "toolchain") cuda.toolkit( name = "cuda", - toolkit_path = "/usr/local/cuda/", + toolkit_path = "", ) use_repo(cuda, "cuda") From 3a960400945231ffaad3045f181ec5a1c270c9d8 Mon Sep 17 00:00:00 2001 From: zhiguo Date: Thu, 16 Oct 2025 11:39:10 +0800 Subject: [PATCH 05/15] fix build on windows 11 with bazel --- .gitignore | 4 +++- example/BUILD.bazel | 2 +- test/BUILD.bazel | 15 +++------------ 3 files changed, 7 insertions(+), 14 deletions(-) diff --git a/.gitignore b/.gitignore index c2bec551..a28581a2 100644 --- a/.gitignore +++ b/.gitignore @@ -18,6 +18,8 @@ MODULE.bazel.lock # JetBrains IDEs .idea/ +.vs/ + .cache compile_commands.json @@ -25,4 +27,4 @@ compile_commands.json # //external is reserved by Bazel and needed for hedron_compile_commands external/ -presubmit.yml \ No newline at end of file +presubmit.yml diff --git a/example/BUILD.bazel b/example/BUILD.bazel index 18214895..4b6f6c6a 100644 --- a/example/BUILD.bazel +++ b/example/BUILD.bazel @@ -10,7 +10,7 @@ cuda_binary( "example_common.h", ], copts = [ - "-std=c++17", + "-std=c++20", ], includes = ["."], rdc = True, diff --git a/test/BUILD.bazel b/test/BUILD.bazel index ad2541f3..3f0a34e9 100644 --- a/test/BUILD.bazel +++ b/test/BUILD.bazel @@ -10,10 +10,7 @@ cuda_test( "--expt-relaxed-constexpr", "--extended-lambda", "-std=c++20", - ] + select({ - "@rules_cc//cc/compiler:msvc-cl": ["/wd4819"], - "//conditions:default": [], - }), + ], rdc = True, deps = [ "//:muda", @@ -39,10 +36,7 @@ cuda_test( "--expt-relaxed-constexpr", "--extended-lambda", "-std=c++20", - ] + select({ - "@rules_cc//cc/compiler:msvc-cl": ["/wd4819"], - "//conditions:default": [], - }), + ], rdc = True, deps = [ "//:muda", @@ -68,10 +62,7 @@ cuda_test( "--expt-relaxed-constexpr", "--extended-lambda", "-std=c++20", - ] + select({ - "@rules_cc//cc/compiler:msvc-cl": ["/wd4819"], - "//conditions:default": [], - }), + ], rdc = True, deps = [ "//:muda", From c8234a27c576030c45a3f06f6a31c29d8b35903f Mon Sep 17 00:00:00 2001 From: zhiguo Date: Thu, 16 Oct 2025 12:32:15 +0800 Subject: [PATCH 06/15] set muda_check on by default --- BUILD.bazel | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/BUILD.bazel b/BUILD.bazel index b7272b15..2bad480b 100644 --- a/BUILD.bazel +++ b/BUILD.bazel @@ -17,7 +17,7 @@ license( bool_flag( name = "muda_check", - build_setting_default = False, + build_setting_default = True, ) config_setting( From d1612c88c6c44c507eb876472378107b68d1fc9a Mon Sep 17 00:00:00 2001 From: zhiguo Date: Fri, 19 Dec 2025 18:19:14 +0800 Subject: [PATCH 07/15] update --- .github/workflows/cmake.yml | 61 ++++++++++++++++++++++++ CMakeLists.txt | 1 + CMakePresets.json | 48 +++++++++++++++++++ src/muda/launch/details/parallel_for.inl | 2 +- src/muda/launch/parallel_for.h | 5 +- test/eigen_test/CMakeLists.txt | 2 + test/linear_system_test/CMakeLists.txt | 16 ++++--- test/spgrid_test/CMakeLists.txt | 4 +- test/unit_test/CMakeLists.txt | 4 +- 9 files changed, 131 insertions(+), 12 deletions(-) create mode 100644 .github/workflows/cmake.yml create mode 100644 CMakePresets.json diff --git a/.github/workflows/cmake.yml b/.github/workflows/cmake.yml new file mode 100644 index 00000000..da3ad918 --- /dev/null +++ b/.github/workflows/cmake.yml @@ -0,0 +1,61 @@ +name: cmake + +on: + pull_request: + branches: + - "mini20" + paths-ignore: + - ".clang-format" + - ".clangd" + - ".gitignore" + - "LICENSE" + - "README.md" + - "docs/**" + - ".github/workflows/xmake.yml" + - "**/xmake.lua" + - "**/*.bazel" + - "**/*.bzl" + - "**/*.bazelrc" + +jobs: + build_and_test: + strategy: + matrix: + os: [ubuntu-22.04] + + runs-on: ${{ matrix.os }} + + concurrency: + group: ${{ github.ref }}-${{ github.base_ref }}-${{ github.head_ref }}-cmake-${{ matrix.os }} + cancel-in-progress: true + + steps: + - name: Maximize build space + if: matrix.os == 'ubuntu-22.04' + uses: AdityaGarg8/remove-unwanted-software@v4.1 + with: + remove-android: "true" + remove-haskell: "true" + remove-codeql: "true" + + - name: Setup cuda + uses: Jimver/cuda-toolkit@v0.2.28 + id: cuda-toolkit + with: + cuda: "12.8.1" + + - name: Setup ninja and Eigen3 on Ubuntu + if: matrix.os == 'ubuntu-22.04' + run: | + sudo apt update + sudo apt install ninja-build libeigen3-dev -y + + - name: Checkout repository + uses: actions/checkout@v4 + with: + submodules: true + + - name: Build MUDA + run: | + cmake --preset ci-release + cmake --build --preset ci-release --parallel 4 diff --git a/CMakeLists.txt b/CMakeLists.txt index e9992f67..5f255052 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -109,5 +109,6 @@ if(MUDA_BUILD_EXAMPLE) endif() if(MUDA_BUILD_TEST) + enable_testing() add_subdirectory("test") endif() \ No newline at end of file diff --git a/CMakePresets.json b/CMakePresets.json new file mode 100644 index 00000000..bc16b324 --- /dev/null +++ b/CMakePresets.json @@ -0,0 +1,48 @@ +{ + "version": 6, + "cmakeMinimumRequired": { + "major": 3, + "minor": 20, + "patch": 0 + }, + "configurePresets": [ + { + "name": "config-base", + "displayName": "base Configuration", + "hidden": true, + "description": "Default build using Ninja generator", + "generator": "Ninja", + "binaryDir": "${sourceDir}/build", + "cacheVariables": { + "MUDA_FORCE_CHECK": "ON", + "MUDA_WITH_CHECK": "ON", + "MUDA_WITH_COMPUTE_GRAPH": "ON", + "MUDA_WITH_NVTX3": "ON" + } + }, + { + "name": "ci-release", + "displayName": "CI-Release Configuration", + "inherits": ["config-base"], + "description": "Configuration for CI builds", + "cacheVariables": { + "MUDA_BUILD_TEST": "ON", + "MUDA_BUILD_EXAMPLE": "OFF", + "CMAKE_CUDA_ARCHITECTURES": "89", + "CMAKE_BUILD_TYPE": "Release" + } + } + ], + "buildPresets": [ + { + "name": "ci-release", + "configurePreset": "ci-release" + } + ], + "testPresets": [ + { + "name": "ci-release", + "configurePreset": "ci-release" + } + ] +} diff --git a/src/muda/launch/details/parallel_for.inl b/src/muda/launch/details/parallel_for.inl index 528afb50..98685cfd 100644 --- a/src/muda/launch/details/parallel_for.inl +++ b/src/muda/launch/details/parallel_for.inl @@ -127,7 +127,7 @@ MUDA_HOST MUDA_NODISCARD auto ParallelFor::as_node_parms(int count, F&& f) { using CallableType = raw_type_t; - check_input(count); + // check_input(count); auto parms = std::make_shared>(std::forward(f), count); if(m_grid_dim <= 0) // dynamic grid dim diff --git a/src/muda/launch/parallel_for.h b/src/muda/launch/parallel_for.h index 87007cf2..05fbe122 100644 --- a/src/muda/launch/parallel_for.h +++ b/src/muda/launch/parallel_for.h @@ -82,10 +82,10 @@ class ParallelForDetails private: template - friend MUDA_GLOBAL void details::parallel_for_kernel(ParallelForCallable f); + friend MUDA_GLOBAL void details::parallel_for_kernel(details::ParallelForCallable f); template - friend MUDA_GLOBAL void details::grid_stride_loop_kernel(ParallelForCallable f); + friend MUDA_GLOBAL void details::grid_stride_loop_kernel(details::ParallelForCallable f); MUDA_DEVICE ParallelForDetails(ParallelForType type, int i, int total_num) MUDA_NOEXCEPT : m_type(type), @@ -169,6 +169,7 @@ class ParallelFor : public LaunchBase m_block_dim(blockDim), m_shared_mem_size(shared_mem_size) { + std::cout << "[DEBUG]" << "blockDim: " << blockDim << std::endl; } diff --git a/test/eigen_test/CMakeLists.txt b/test/eigen_test/CMakeLists.txt index 80822078..3f68f249 100644 --- a/test/eigen_test/CMakeLists.txt +++ b/test/eigen_test/CMakeLists.txt @@ -11,3 +11,5 @@ target_include_directories(muda_eigen_test PRIVATE target_link_libraries(muda_eigen_test PRIVATE muda Eigen3::Eigen) source_group(TREE "${PROJECT_SOURCE_DIR}/test" PREFIX "test" FILES ${SOURCE}) source_group(TREE "${PROJECT_SOURCE_DIR}/src" PREFIX "src" FILES ${MUDA_HEADER_FILES}) + +add_test(NAME muda_eigen_test COMMAND muda_eigen_test) diff --git a/test/linear_system_test/CMakeLists.txt b/test/linear_system_test/CMakeLists.txt index db29b783..f7c8e64f 100644 --- a/test/linear_system_test/CMakeLists.txt +++ b/test/linear_system_test/CMakeLists.txt @@ -4,13 +4,15 @@ file(GLOB_RECURSE SOURCE "${CMAKE_CURRENT_SOURCE_DIR}/*.cpp" "${CMAKE_CURRENT_SOURCE_DIR}/*.cu" "${CMAKE_CURRENT_SOURCE_DIR}/*.h") -add_executable(muda_linear_sysytem_test) -target_sources(muda_linear_sysytem_test PRIVATE ${SOURCE} ${MUDA_HEADER_FILES}) -target_compile_features(muda_linear_sysytem_test PRIVATE cxx_std_20) -set_target_properties(muda_linear_sysytem_test PROPERTIES CUDA_SEPARABLE_COMPILATION ON) -target_include_directories(muda_linear_sysytem_test PRIVATE +add_executable(muda_linear_system_test) +target_sources(muda_linear_system_test PRIVATE ${SOURCE} ${MUDA_HEADER_FILES}) +target_compile_features(muda_linear_system_test PRIVATE cxx_std_20) +set_target_properties(muda_linear_system_test PROPERTIES CUDA_SEPARABLE_COMPILATION ON) +target_include_directories(muda_linear_system_test PRIVATE "${PROJECT_SOURCE_DIR}/test" "${PROJECT_SOURCE_DIR}/external") -target_link_libraries(muda_linear_sysytem_test PRIVATE muda cusparse cublas cusolver Eigen3::Eigen) +target_link_libraries(muda_linear_system_test PRIVATE muda cusparse cublas cusolver Eigen3::Eigen) source_group(TREE "${PROJECT_SOURCE_DIR}/test" PREFIX "test" FILES ${SOURCE}) -source_group(TREE "${PROJECT_SOURCE_DIR}/src" PREFIX "src" FILES ${MUDA_HEADER_FILES}) \ No newline at end of file +source_group(TREE "${PROJECT_SOURCE_DIR}/src" PREFIX "src" FILES ${MUDA_HEADER_FILES}) + +add_test(NAME muda_linear_system_test COMMAND muda_linear_system_test) diff --git a/test/spgrid_test/CMakeLists.txt b/test/spgrid_test/CMakeLists.txt index ae271781..3048418a 100644 --- a/test/spgrid_test/CMakeLists.txt +++ b/test/spgrid_test/CMakeLists.txt @@ -12,4 +12,6 @@ set_target_properties(muda_spgrid_test PROPERTIES CUDA_SEPARABLE_COMPILATION ON) target_include_directories(muda_spgrid_test PRIVATE "${PROJECT_SOURCE_DIR}/test" "${PROJECT_SOURCE_DIR}/external") - target_link_libraries(muda_spgrid_test PRIVATE muda Eigen3::Eigen) \ No newline at end of file +target_link_libraries(muda_spgrid_test PRIVATE muda Eigen3::Eigen) + +add_test(NAME muda_spgrid_test COMMAND muda_spgrid_test) diff --git a/test/unit_test/CMakeLists.txt b/test/unit_test/CMakeLists.txt index 87043c98..6f1b86bc 100644 --- a/test/unit_test/CMakeLists.txt +++ b/test/unit_test/CMakeLists.txt @@ -14,4 +14,6 @@ target_include_directories(muda_unit_test PRIVATE target_link_libraries(muda_unit_test PRIVATE muda cusparse cublas cusolver Eigen3::Eigen) target_compile_definitions(muda_unit_test PRIVATE "-DMUDA_TEST_DATA_DIR=R\"(${PROJECT_SOURCE_DIR}/test/data)\"") source_group(TREE "${PROJECT_SOURCE_DIR}/test" PREFIX "test" FILES ${SOURCE}) -source_group(TREE "${PROJECT_SOURCE_DIR}/src" PREFIX "src" FILES ${MUDA_HEADER_FILES}) \ No newline at end of file +source_group(TREE "${PROJECT_SOURCE_DIR}/src" PREFIX "src" FILES ${MUDA_HEADER_FILES}) + +add_test(NAME muda_unit_test COMMAND muda_unit_test) From 2a9fc1ba31835edc6bdf1c2e9d47baf058d8bfdb Mon Sep 17 00:00:00 2001 From: zhiguo Date: Thu, 30 Oct 2025 10:54:51 +0800 Subject: [PATCH 08/15] update --- src/muda/launch/details/memory.inl | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/muda/launch/details/memory.inl b/src/muda/launch/details/memory.inl index 25b2c954..60c579bc 100644 --- a/src/muda/launch/details/memory.inl +++ b/src/muda/launch/details/memory.inl @@ -233,7 +233,7 @@ MUDA_INLINE MUDA_HOST Memory& Memory::alloc(cudaPitchedPtr* pitched_ptr, return alloc_3d(pitched_ptr, extent, async); } -MUDA_INLINE MUDA_HOST Memory& muda::Memory::free(cudaPitchedPtr pitched_ptr, bool async) +MUDA_INLINE MUDA_HOST Memory& Memory::free(cudaPitchedPtr pitched_ptr, bool async) { return free(pitched_ptr.ptr, async); } From 8f04432bfcaea6c596bb3239902bd74335cc5fa9 Mon Sep 17 00:00:00 2001 From: Ligo <1569731402@qq.com> Date: Tue, 18 Nov 2025 10:09:48 +0800 Subject: [PATCH 09/15] Comment out debug output in parallel_for.h Commented out debug output for blockDim. --- src/muda/launch/parallel_for.h | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/src/muda/launch/parallel_for.h b/src/muda/launch/parallel_for.h index 05fbe122..83161f39 100644 --- a/src/muda/launch/parallel_for.h +++ b/src/muda/launch/parallel_for.h @@ -169,7 +169,7 @@ class ParallelFor : public LaunchBase m_block_dim(blockDim), m_shared_mem_size(shared_mem_size) { - std::cout << "[DEBUG]" << "blockDim: " << blockDim << std::endl; + // std::cout << "[DEBUG]" << "blockDim: " << blockDim << std::endl; } @@ -235,4 +235,4 @@ class ParallelFor : public LaunchBase }; } // namespace muda -#include "details/parallel_for.inl" \ No newline at end of file +#include "details/parallel_for.inl" From 3fdf37aab78097a4c089b8b962418a6f6dfc6c81 Mon Sep 17 00:00:00 2001 From: MuGdxy Date: Mon, 24 Nov 2025 02:43:19 +0800 Subject: [PATCH 10/15] make no wait on device-device memory transfer --- src/muda/buffer/details/buffer_2d_view.inl | 33 +++++++------- src/muda/buffer/details/buffer_3d_view.inl | 38 ++++++++-------- src/muda/buffer/details/buffer_view.inl | 39 ++++++++-------- src/muda/buffer/details/device_buffer.inl | 44 +++++-------------- src/muda/buffer/details/device_buffer_2d.inl | 38 ++++++---------- src/muda/buffer/details/device_buffer_3d.inl | 36 +++++---------- .../matrix_format_converter_impl.h | 30 ++++++------- 7 files changed, 107 insertions(+), 151 deletions(-) diff --git a/src/muda/buffer/details/buffer_2d_view.inl b/src/muda/buffer/details/buffer_2d_view.inl index bd6ad3c2..ac06d127 100644 --- a/src/muda/buffer/details/buffer_2d_view.inl +++ b/src/muda/buffer/details/buffer_2d_view.inl @@ -42,14 +42,14 @@ MUDA_GENERIC Buffer2DViewT::Buffer2DViewT(auto_const_t* data, } template -MUDA_GENERIC auto Buffer2DViewT::as_const() const MUDA_NOEXCEPT->ConstView +MUDA_GENERIC auto Buffer2DViewT::as_const() const MUDA_NOEXCEPT -> ConstView { return ConstView{*this}; } template -MUDA_GENERIC auto Buffer2DViewT::data(size_t x, - size_t y) const MUDA_NOEXCEPT->auto_const_t* +MUDA_GENERIC auto Buffer2DViewT::data(size_t x, size_t y) const MUDA_NOEXCEPT + -> auto_const_t* { x += m_offset.offset_in_height(); y += m_offset.offset_in_width(); @@ -60,8 +60,8 @@ MUDA_GENERIC auto Buffer2DViewT::data(size_t x, } template -MUDA_GENERIC auto Buffer2DViewT::data(size_t flatten_i) const - MUDA_NOEXCEPT->auto_const_t* +MUDA_GENERIC auto Buffer2DViewT::data(size_t flatten_i) const MUDA_NOEXCEPT + -> auto_const_t* { auto x = flatten_i / m_extent.width(); auto y = flatten_i % m_extent.width(); @@ -69,15 +69,16 @@ MUDA_GENERIC auto Buffer2DViewT::data(size_t flatten_i) const } template -MUDA_GENERIC auto Buffer2DViewT::origin_data() const MUDA_NOEXCEPT->auto_const_t* +MUDA_GENERIC auto Buffer2DViewT::origin_data() const MUDA_NOEXCEPT + -> auto_const_t* { return m_data; } template -MUDA_GENERIC auto Buffer2DViewT::subview( - Offset2D offset, Extent2D extent) const MUDA_NOEXCEPT->ThisView +MUDA_GENERIC auto Buffer2DViewT::subview(Offset2D offset, Extent2D extent) const MUDA_NOEXCEPT + -> ThisView { #ifndef __CUDA_ARCH__ if(ComputeGraphBuilder::is_topo_building()) @@ -97,7 +98,7 @@ MUDA_GENERIC auto Buffer2DViewT::subview( } template -MUDA_GENERIC auto Buffer2DViewT::viewer() const MUDA_NOEXCEPT->ThisViewer +MUDA_GENERIC auto Buffer2DViewT::viewer() const MUDA_NOEXCEPT -> ThisViewer { return ThisViewer{m_data, make_int2((int)m_offset.offset_in_height(), @@ -117,30 +118,30 @@ MUDA_GENERIC cudaPitchedPtr Buffer2DViewT::cuda_pitched_ptr() const } template -MUDA_GENERIC auto Buffer2DViewT::extent() const MUDA_NOEXCEPT->Extent2D +MUDA_GENERIC auto Buffer2DViewT::extent() const MUDA_NOEXCEPT -> Extent2D { return m_extent; } template -MUDA_GENERIC auto Buffer2DViewT::pitch_bytes() const MUDA_NOEXCEPT->size_t +MUDA_GENERIC auto Buffer2DViewT::pitch_bytes() const MUDA_NOEXCEPT -> size_t { return m_pitch_bytes; } template -MUDA_GENERIC auto Buffer2DViewT::offset() const MUDA_NOEXCEPT->Offset2D +MUDA_GENERIC auto Buffer2DViewT::offset() const MUDA_NOEXCEPT -> Offset2D { return m_offset; } template -MUDA_GENERIC auto Buffer2DViewT::total_size() const MUDA_NOEXCEPT->size_t +MUDA_GENERIC auto Buffer2DViewT::total_size() const MUDA_NOEXCEPT -> size_t { return m_extent.width() * m_extent.height(); } template -MUDA_GENERIC auto Buffer2DViewT::cviewer() const MUDA_NOEXCEPT->CViewer +MUDA_GENERIC auto Buffer2DViewT::cviewer() const MUDA_NOEXCEPT -> CViewer { return viewer(); } @@ -154,14 +155,14 @@ MUDA_HOST void Buffer2DViewT::copy_to(T* host) const template MUDA_HOST void Buffer2DViewT::fill(const T& val) MUDA_REQUIRES(!IsConst) { - BufferLaunch().template fill(*this, val).wait(); + BufferLaunch().template fill(*this, val); } template MUDA_HOST void Buffer2DViewT::copy_from(const Buffer2DViewT& other) MUDA_REQUIRES(!IsConst) { - BufferLaunch().template copy(*this, other).wait(); + BufferLaunch().template copy(*this, other); } template diff --git a/src/muda/buffer/details/buffer_3d_view.inl b/src/muda/buffer/details/buffer_3d_view.inl index 617c142a..32931666 100644 --- a/src/muda/buffer/details/buffer_3d_view.inl +++ b/src/muda/buffer/details/buffer_3d_view.inl @@ -45,14 +45,14 @@ MUDA_GENERIC Buffer3DViewT::Buffer3DViewT(T* data, } template -MUDA_GENERIC auto Buffer3DViewT::as_const() const MUDA_NOEXCEPT->ConstView +MUDA_GENERIC auto Buffer3DViewT::as_const() const MUDA_NOEXCEPT -> ConstView { return ConstView{*this}; } template -MUDA_GENERIC auto Buffer3DViewT::data(size_t x, size_t y, size_t z) const - MUDA_NOEXCEPT->auto_const_t* +MUDA_GENERIC auto Buffer3DViewT::data(size_t x, size_t y, size_t z) const MUDA_NOEXCEPT + -> auto_const_t* { x += m_offset.offset_in_depth(); y += m_offset.offset_in_height(); @@ -64,8 +64,8 @@ MUDA_GENERIC auto Buffer3DViewT::data(size_t x, size_t y, size_t z) } template -MUDA_GENERIC auto Buffer3DViewT::data(size_t flatten_i) const - MUDA_NOEXCEPT->auto_const_t* +MUDA_GENERIC auto Buffer3DViewT::data(size_t flatten_i) const MUDA_NOEXCEPT + -> auto_const_t* { auto area = m_extent.width() * m_extent.height(); auto x = flatten_i / area; @@ -77,39 +77,39 @@ MUDA_GENERIC auto Buffer3DViewT::data(size_t flatten_i) const } template -MUDA_GENERIC auto Buffer3DViewT::origin_data() const MUDA_NOEXCEPT->auto_const_t* +MUDA_GENERIC auto Buffer3DViewT::origin_data() const MUDA_NOEXCEPT + -> auto_const_t* { return m_data; } template -MUDA_GENERIC auto Buffer3DViewT::extent() const MUDA_NOEXCEPT->Extent3D +MUDA_GENERIC auto Buffer3DViewT::extent() const MUDA_NOEXCEPT -> Extent3D { return m_extent; } template -MUDA_GENERIC auto Buffer3DViewT::offset() const MUDA_NOEXCEPT->Offset3D +MUDA_GENERIC auto Buffer3DViewT::offset() const MUDA_NOEXCEPT -> Offset3D { return m_offset; } template -MUDA_GENERIC auto Buffer3DViewT::pitch_bytes() const MUDA_NOEXCEPT->size_t +MUDA_GENERIC auto Buffer3DViewT::pitch_bytes() const MUDA_NOEXCEPT -> size_t { return m_pitch_bytes; } template -MUDA_GENERIC auto Buffer3DViewT::pitch_bytes_area() const - MUDA_NOEXCEPT->size_t +MUDA_GENERIC auto Buffer3DViewT::pitch_bytes_area() const MUDA_NOEXCEPT -> size_t { return m_pitch_bytes_area; } template -MUDA_GENERIC auto Buffer3DViewT::cuda_pitched_ptr() const - MUDA_NOEXCEPT->cudaPitchedPtr +MUDA_GENERIC auto Buffer3DViewT::cuda_pitched_ptr() const MUDA_NOEXCEPT + -> cudaPitchedPtr { // don't use make_cudaPitchedPtr (__host__ only function) return cudaPitchedPtr{.ptr = remove_const(m_data), @@ -125,8 +125,8 @@ MUDA_GENERIC size_t Buffer3DViewT::total_size() const MUDA_NOEXCEPT } template -MUDA_GENERIC auto Buffer3DViewT::subview( - Offset3D offset, Extent3D extent) const MUDA_NOEXCEPT->ThisView +MUDA_GENERIC auto Buffer3DViewT::subview(Offset3D offset, Extent3D extent) const MUDA_NOEXCEPT + -> ThisView { #ifndef __CUDA_ARCH__ if(ComputeGraphBuilder::is_topo_building()) @@ -148,7 +148,7 @@ MUDA_GENERIC auto Buffer3DViewT::subview( } template -MUDA_GENERIC auto Buffer3DViewT::viewer() const MUDA_NOEXCEPT->ThisViewer +MUDA_GENERIC auto Buffer3DViewT::viewer() const MUDA_NOEXCEPT -> ThisViewer { return ThisViewer{m_data, make_int3(m_offset.offset_in_depth(), @@ -160,7 +160,7 @@ MUDA_GENERIC auto Buffer3DViewT::viewer() const MUDA_NOEXCEPT->ThisV } template -MUDA_GENERIC auto Buffer3DViewT::cviewer() const MUDA_NOEXCEPT->CViewer +MUDA_GENERIC auto Buffer3DViewT::cviewer() const MUDA_NOEXCEPT -> CViewer { return viewer(); } @@ -169,7 +169,7 @@ template MUDA_HOST void Buffer3DViewT::fill(const T& v) const MUDA_REQUIRES(!IsConst) { static_assert(!IsConst, "This must be non-const buffer."); - BufferLaunch().template fill(*this, v).wait(); + BufferLaunch().template fill(*this, v); } template @@ -177,7 +177,7 @@ MUDA_HOST void Buffer3DViewT::copy_from(const Buffer3DViewT MUDA_REQUIRES(!IsConst) { static_assert(!IsConst, "This must be non-const buffer."); - BufferLaunch().template copy(*this, other).wait(); + BufferLaunch().template copy(*this, other); } template diff --git a/src/muda/buffer/details/buffer_view.inl b/src/muda/buffer/details/buffer_view.inl index 9b250bb8..57b5cb0a 100644 --- a/src/muda/buffer/details/buffer_view.inl +++ b/src/muda/buffer/details/buffer_view.inl @@ -31,41 +31,45 @@ MUDA_GENERIC BufferViewT::BufferViewT(const BufferViewT -MUDA_GENERIC auto BufferViewT::as_const() const MUDA_NOEXCEPT->ConstView +MUDA_GENERIC auto BufferViewT::as_const() const MUDA_NOEXCEPT -> ConstView { return ConstView{*this}; } template -MUDA_GENERIC auto BufferViewT::data() const MUDA_NOEXCEPT->auto_const_t* +MUDA_GENERIC auto BufferViewT::data() const MUDA_NOEXCEPT + -> auto_const_t* { return m_data + m_offset; } template -MUDA_GENERIC auto BufferViewT::data(size_t i) const MUDA_NOEXCEPT->auto_const_t* +MUDA_GENERIC auto BufferViewT::data(size_t i) const MUDA_NOEXCEPT + -> auto_const_t* { i += m_offset; return m_data + i; } template -MUDA_GENERIC auto BufferViewT::origin_data() const MUDA_NOEXCEPT->auto_const_t* +MUDA_GENERIC auto BufferViewT::origin_data() const MUDA_NOEXCEPT + -> auto_const_t* { return m_data; } template -MUDA_GENERIC auto BufferViewT::operator[](size_t i) const - MUDA_NOEXCEPT->auto_const_t& +MUDA_GENERIC auto BufferViewT::operator[](size_t i) const MUDA_NOEXCEPT + -> auto_const_t& { return *data(i); } template -MUDA_GENERIC auto BufferViewT::subview(size_t offset, size_t size) const - MUDA_NOEXCEPT->ThisView +MUDA_GENERIC auto BufferViewT::subview(size_t offset, size_t size) const MUDA_NOEXCEPT + -> ThisView { #ifndef __CUDA_ARCH__ if(ComputeGraphBuilder::is_topo_building()) @@ -81,13 +85,13 @@ MUDA_GENERIC auto BufferViewT::subview(size_t offset, size_t size) c } template -MUDA_GENERIC auto BufferViewT::viewer() const MUDA_NOEXCEPT->ThisViewer +MUDA_GENERIC auto BufferViewT::viewer() const MUDA_NOEXCEPT -> ThisViewer { return ThisViewer{data(), static_cast(m_size)}; } template -MUDA_GENERIC auto BufferViewT::cviewer() const MUDA_NOEXCEPT->CViewer +MUDA_GENERIC auto BufferViewT::cviewer() const MUDA_NOEXCEPT -> CViewer { return CViewer{data(), static_cast(m_size)}; } @@ -97,9 +101,7 @@ MUDA_HOST void BufferViewT::fill(const T& v) const MUDA_REQUIRES(!Is { static_assert(!IsConst, "This must be non-const"); - BufferLaunch() - .template fill(*this, v) // - .wait(); + BufferLaunch().template fill(*this, v); } template @@ -108,9 +110,7 @@ MUDA_HOST void BufferViewT::copy_from(const BufferViewT& ot { static_assert(!IsConst, "This must be non-const"); - BufferLaunch() - .template copy(*this, other) // - .wait(); + BufferLaunch().template copy(*this, other); } template @@ -133,19 +133,20 @@ MUDA_HOST void BufferViewT::copy_to(T* host) const } template -MUDA_GENERIC auto BufferViewT::operator+(int i) const MUDA_NOEXCEPT->ThisView +MUDA_GENERIC auto BufferViewT::operator+(int i) const MUDA_NOEXCEPT -> ThisView { return ThisView{m_data, m_offset + i, m_size - i}; } template -MUDA_GENERIC auto BufferViewT::operator*() const MUDA_NOEXCEPT->reference +MUDA_GENERIC auto BufferViewT::operator*() const MUDA_NOEXCEPT -> reference { return *data(0); } template -MUDA_GENERIC auto BufferViewT::operator[](int i) const MUDA_NOEXCEPT->auto_const_t& +MUDA_GENERIC auto BufferViewT::operator[](int i) const MUDA_NOEXCEPT + -> auto_const_t& { return *data(i); } diff --git a/src/muda/buffer/details/device_buffer.inl b/src/muda/buffer/details/device_buffer.inl index f57886ac..34b8256e 100644 --- a/src/muda/buffer/details/device_buffer.inl +++ b/src/muda/buffer/details/device_buffer.inl @@ -6,9 +6,7 @@ namespace muda template DeviceBuffer::DeviceBuffer(size_t n) { - BufferLaunch() - .alloc(*this, n) // - .wait(); + BufferLaunch().alloc(*this, n); } template @@ -24,8 +22,7 @@ DeviceBuffer::DeviceBuffer(const DeviceBuffer& other) { BufferLaunch() .alloc(*this, other.size()) // - .copy(view(), other.view()) // - .wait(); + .copy(view(), other.view()); } template @@ -47,8 +44,7 @@ DeviceBuffer& DeviceBuffer::operator=(const DeviceBuffer& other) BufferLaunch() .resize(*this, other.size()) // - .copy(view(), other.view()) // - .wait(); + .copy(view(), other.view()); return *this; } @@ -61,9 +57,7 @@ DeviceBuffer& DeviceBuffer::operator=(DeviceBuffer&& other) if(m_data) { - BufferLaunch() - .free(*this) // - .wait(); + BufferLaunch().free(*this); } m_data = other.m_data; @@ -82,8 +76,7 @@ DeviceBuffer::DeviceBuffer(CBufferView other) { BufferLaunch() .alloc(*this, other.size()) // - .copy(view(), other) // - .wait(); + .copy(view(), other); } template @@ -100,8 +93,7 @@ DeviceBuffer& DeviceBuffer::operator=(CBufferView other) { BufferLaunch() .resize(*this, other.size()) // - .copy(view(), other) // - .wait(); + .copy(view(), other); return *this; } @@ -134,41 +126,31 @@ void DeviceBuffer::copy_from(const std::vector& host) template void DeviceBuffer::resize(size_t new_size) { - BufferLaunch() - .resize(*this, new_size) // - .wait(); + BufferLaunch().resize(*this, new_size); } template void DeviceBuffer::resize(size_t new_size, const value_type& value) { - BufferLaunch() - .resize(*this, new_size, value) // - .wait(); + BufferLaunch().resize(*this, new_size, value); } template void DeviceBuffer::reserve(size_t new_capacity) { - BufferLaunch() - .reserve(*this, new_capacity) // - .wait(); + BufferLaunch().reserve(*this, new_capacity); } template void DeviceBuffer::clear() { - BufferLaunch() - .clear(*this) // - .wait(); + BufferLaunch().clear(*this); } template void DeviceBuffer::shrink_to_fit() { - BufferLaunch() - .shrink_to_fit(*this) // - .wait(); + BufferLaunch().shrink_to_fit(*this); } template @@ -218,9 +200,7 @@ DeviceBuffer::~DeviceBuffer() { if(m_data) { - BufferLaunch() - .free(*this) // - .wait(); + BufferLaunch().free(*this); } } } // namespace muda \ No newline at end of file diff --git a/src/muda/buffer/details/device_buffer_2d.inl b/src/muda/buffer/details/device_buffer_2d.inl index 12cad8de..fdce8d18 100644 --- a/src/muda/buffer/details/device_buffer_2d.inl +++ b/src/muda/buffer/details/device_buffer_2d.inl @@ -5,9 +5,7 @@ namespace muda template DeviceBuffer2D::DeviceBuffer2D(const Extent2D& n) { - BufferLaunch() - .resize(*this, n) // - .wait(); + BufferLaunch().resize(*this, n); } template @@ -24,8 +22,7 @@ DeviceBuffer2D::DeviceBuffer2D(const DeviceBuffer2D& other) { BufferLaunch() .resize(*this, other.extent()) // - .copy(view(), other.view()) // - .wait(); + .copy(view(), other.view()); } template @@ -49,8 +46,7 @@ DeviceBuffer2D& DeviceBuffer2D::operator=(const DeviceBuffer2D& other) BufferLaunch() .resize(*this, other.extent()) // - .copy(view(), other.view()) // - .wait(); + .copy(view(), other.view()); return *this; } @@ -62,7 +58,7 @@ DeviceBuffer2D& DeviceBuffer2D::operator=(DeviceBuffer2D&& other) return *this; if(m_data) - BufferLaunch().free(*this).wait(); + BufferLaunch().free(*this); m_data = other.m_data; m_pitch_bytes = other.m_pitch_bytes; @@ -82,8 +78,7 @@ DeviceBuffer2D::DeviceBuffer2D(CBuffer2DView other) { BufferLaunch() .alloc(*this, other.extent()) // - .copy(view(), other) // - .wait(); + .copy(view(), other); } template @@ -91,8 +86,7 @@ DeviceBuffer2D& DeviceBuffer2D::operator=(CBuffer2DView other) { BufferLaunch() .resize(*this, other.extent()) // - .copy(view(), other) // - .wait(); + .copy(view(), other); return *this; } @@ -122,49 +116,43 @@ void DeviceBuffer2D::copy_from(const std::vector& host) template void DeviceBuffer2D::resize(Extent2D new_extent) { - BufferLaunch() - .resize(*this, new_extent) // - .wait(); + BufferLaunch().resize(*this, new_extent); } template void DeviceBuffer2D::resize(Extent2D new_extent, const T& value) { - BufferLaunch() - .resize(*this, new_extent, value) // - .wait(); + BufferLaunch().resize(*this, new_extent, value); } template void DeviceBuffer2D::reserve(Extent2D new_capacity) { - BufferLaunch() - .reserve(*this, new_capacity) // - .wait(); + BufferLaunch().reserve(*this, new_capacity); } template void DeviceBuffer2D::clear() { - BufferLaunch().clear(*this).wait(); + BufferLaunch().clear(*this); } template void DeviceBuffer2D::shrink_to_fit() { - BufferLaunch().shrink_to_fit(*this).wait(); + BufferLaunch().shrink_to_fit(*this); } template void DeviceBuffer2D::fill(const T& v) { - BufferLaunch().fill(view(), v).wait(); + BufferLaunch().fill(view(), v); } template DeviceBuffer2D::~DeviceBuffer2D() { if(m_data) - BufferLaunch().free(*this).wait(); + BufferLaunch().free(*this); } } // namespace muda \ No newline at end of file diff --git a/src/muda/buffer/details/device_buffer_3d.inl b/src/muda/buffer/details/device_buffer_3d.inl index 11763684..3dd26602 100644 --- a/src/muda/buffer/details/device_buffer_3d.inl +++ b/src/muda/buffer/details/device_buffer_3d.inl @@ -5,9 +5,7 @@ namespace muda template DeviceBuffer3D::DeviceBuffer3D(const Extent3D& n) { - BufferLaunch() - .resize(*this, n) // - .wait(); + BufferLaunch().resize(*this, n); } template @@ -25,8 +23,7 @@ DeviceBuffer3D::DeviceBuffer3D(const DeviceBuffer3D& other) { BufferLaunch() .resize(*this, other.extent()) // - .copy(view(), other.view()) // - .wait(); + .copy(view(), other.view()); } template @@ -52,8 +49,7 @@ DeviceBuffer3D& DeviceBuffer3D::operator=(const DeviceBuffer3D& other) BufferLaunch() .resize(*this, other.extent()) // - .copy(view(), other.view()) // - .wait(); + .copy(view(), other.view()); return *this; } @@ -87,8 +83,7 @@ DeviceBuffer3D::DeviceBuffer3D(CBuffer3DView other) { BufferLaunch() .alloc(*this, other.extent()) // - .copy(view(), other) // - .wait(); + .copy(view(), other); } template @@ -96,8 +91,7 @@ DeviceBuffer3D& DeviceBuffer3D::operator=(CBuffer3DView other) { BufferLaunch() .resize(*this, other.extent()) // - .copy(view(), other) // - .wait(); + .copy(view(), other); return *this; } @@ -125,49 +119,43 @@ void DeviceBuffer3D::copy_from(const std::vector& host) template void DeviceBuffer3D::resize(Extent3D new_extent) { - BufferLaunch() - .resize(*this, new_extent) // - .wait(); + BufferLaunch().resize(*this, new_extent); } template void DeviceBuffer3D::resize(Extent3D new_extent, const T& value) { - BufferLaunch() - .resize(*this, new_extent, value) // - .wait(); + BufferLaunch().resize(*this, new_extent, value); } template void DeviceBuffer3D::reserve(Extent3D new_capacity) { - BufferLaunch() - .reserve(*this, new_capacity) // - .wait(); + BufferLaunch().reserve(*this, new_capacity); } template void DeviceBuffer3D::clear() { - BufferLaunch().clear(*this).wait(); + BufferLaunch().clear(*this); } template void DeviceBuffer3D::shrink_to_fit() { - BufferLaunch().shrink_to_fit(*this).wait(); + BufferLaunch().shrink_to_fit(*this); } template void DeviceBuffer3D::fill(const T& v) { - BufferLaunch().fill(view(), v).wait(); + BufferLaunch().fill(view(), v); } template DeviceBuffer3D::~DeviceBuffer3D() { if(m_data) - BufferLaunch().free(*this).wait(); + BufferLaunch().free(*this); } } // namespace muda \ No newline at end of file diff --git a/src/muda/ext/linear_system/matrix_format_converter_impl.h b/src/muda/ext/linear_system/matrix_format_converter_impl.h index 1ef10cfe..f67c8aec 100644 --- a/src/muda/ext/linear_system/matrix_format_converter_impl.h +++ b/src/muda/ext/linear_system/matrix_format_converter_impl.h @@ -21,14 +21,14 @@ #include +namespace muda +{ // for encode run length usage MUDA_GENERIC constexpr bool operator==(const int2& a, const int2& b) { return a.x == b.x && a.y == b.y; } -namespace muda -{ namespace details { class MatrixFormatConverterBase @@ -292,13 +292,12 @@ namespace details [sort_index = sort_index.viewer().name("sort_index")] __device__( int i) mutable { sort_index(i) = i; }); - DeviceMergeSort().SortPairs(ij_pairs.data(), - sort_index.data(), - ij_pairs.size(), - [] __device__(const int2& a, const int2& b) { - return a.x < b.x - || (a.x == b.x && a.y < b.y); - }); + DeviceMergeSort().SortPairs( + ij_pairs.data(), + sort_index.data(), + ij_pairs.size(), + [] __device__(const int2& a, const int2& b) + { return a.x < b.x || (a.x == b.x && a.y < b.y); }); // set ij_pairs back to row_indices and col_indices @@ -527,13 +526,12 @@ namespace details ij_pairs(i).y = col_indices(i); }); - DeviceMergeSort().SortPairs(ij_pairs.data(), - to.m_values.data(), - ij_pairs.size(), - [] __device__(const int2& a, const int2& b) { - return a.x < b.x - || (a.x == b.x && a.y < b.y); - }); + DeviceMergeSort().SortPairs( + ij_pairs.data(), + to.m_values.data(), + ij_pairs.size(), + [] __device__(const int2& a, const int2& b) + { return a.x < b.x || (a.x == b.x && a.y < b.y); }); // set ij_pairs back to row_indices and col_indices From ac177254b2ad3bb9bef49f74eff315675ded138a Mon Sep 17 00:00:00 2001 From: MuGdxy Date: Wed, 3 Dec 2025 22:37:56 +0800 Subject: [PATCH 11/15] use muda::details::IntPair instead of int2 to prevent equal_to<> conflict --- .../matrix_format_converter_impl.h | 30 ++++++++++++++----- 1 file changed, 22 insertions(+), 8 deletions(-) diff --git a/src/muda/ext/linear_system/matrix_format_converter_impl.h b/src/muda/ext/linear_system/matrix_format_converter_impl.h index f67c8aec..d08f4d9a 100644 --- a/src/muda/ext/linear_system/matrix_format_converter_impl.h +++ b/src/muda/ext/linear_system/matrix_format_converter_impl.h @@ -21,14 +21,28 @@ #include + namespace muda { -// for encode run length usage -MUDA_GENERIC constexpr bool operator==(const int2& a, const int2& b) +namespace details { - return a.x == b.x && a.y == b.y; -} + struct IntPair + { + int x; + int y; + }; + // A internal muda pair type to avoid int2 equality redefinition conflict + constexpr bool operator==(const IntPair& l, const IntPair& r) + { + return l.x == r.x && l.y == r.y; + } +} // namespace details +} // namespace muda + + +namespace muda +{ namespace details { class MatrixFormatConverterBase @@ -113,10 +127,10 @@ namespace details DeviceVar count; - DeviceBuffer ij_pairs; + DeviceBuffer ij_pairs; DeviceBuffer ij_hash; DeviceBuffer ij_hash_input; - DeviceBuffer unique_ij_pairs; + DeviceBuffer unique_ij_pairs; muda::DeviceBuffer blocks_sorted; DeviceBuffer unique_blocks; @@ -296,7 +310,7 @@ namespace details ij_pairs.data(), sort_index.data(), ij_pairs.size(), - [] __device__(const int2& a, const int2& b) + [] __device__(const IntPair& a, const IntPair& b) { return a.x < b.x || (a.x == b.x && a.y < b.y); }); @@ -530,7 +544,7 @@ namespace details ij_pairs.data(), to.m_values.data(), ij_pairs.size(), - [] __device__(const int2& a, const int2& b) + [] __device__(const IntPair& a, const IntPair& b) { return a.x < b.x || (a.x == b.x && a.y < b.y); }); // set ij_pairs back to row_indices and col_indices From fbafe50f29ff555c171c0dbe8fe679eb17bf5c6c Mon Sep 17 00:00:00 2001 From: MuGdxy Date: Thu, 4 Dec 2025 01:05:43 +0800 Subject: [PATCH 12/15] add missing cuda_type_mapper header --- src/muda/ext/linear_system/device_bcoo_vector.h | 1 + 1 file changed, 1 insertion(+) diff --git a/src/muda/ext/linear_system/device_bcoo_vector.h b/src/muda/ext/linear_system/device_bcoo_vector.h index 91abbf6c..611d1f29 100644 --- a/src/muda/ext/linear_system/device_bcoo_vector.h +++ b/src/muda/ext/linear_system/device_bcoo_vector.h @@ -1,5 +1,6 @@ #pragma once #include +#include #include #include From cee2a36bced926268d6beb333cb91db6c63f7362 Mon Sep 17 00:00:00 2001 From: zhiguo Date: Fri, 19 Dec 2025 17:55:15 +0800 Subject: [PATCH 13/15] update --- src/muda/launch/stream.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/muda/launch/stream.h b/src/muda/launch/stream.h index ffc23d45..d23c2435 100644 --- a/src/muda/launch/stream.h +++ b/src/muda/launch/stream.h @@ -77,7 +77,7 @@ class Stream std::byte* workspace(size_t byte_size); private: - Stream(nullptr_t) + Stream(std::nullptr_t) : m_handle(nullptr) { } From a8e7787c746cc15a1ed350182dff200484817cc8 Mon Sep 17 00:00:00 2001 From: zhiguo Date: Mon, 22 Dec 2025 15:53:37 +0800 Subject: [PATCH 14/15] fix cmake to use third party --- example/CMakeLists.txt | 2 +- test/spgrid_test/CMakeLists.txt | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/example/CMakeLists.txt b/example/CMakeLists.txt index 767f4fad..9f21716d 100644 --- a/example/CMakeLists.txt +++ b/example/CMakeLists.txt @@ -19,7 +19,7 @@ set_target_properties(muda_example PROPERTIES CUDA_SEPARABLE_COMPILATION ON) # target_link_libraries(muda_example PRIVATE fmt::fmt-header-only) target_include_directories(muda_example PRIVATE "${PROJECT_SOURCE_DIR}/example" - "${PROJECT_SOURCE_DIR}/external") + "${PROJECT_SOURCE_DIR}/third_party") target_link_libraries(muda_example PRIVATE muda Eigen3::Eigen) source_group(TREE "${PROJECT_SOURCE_DIR}/example" PREFIX "example" FILES ${SOURCE_FILES}) source_group(TREE "${PROJECT_SOURCE_DIR}/src" PREFIX "src" FILES ${MUDA_HEADER_FILES}) diff --git a/test/spgrid_test/CMakeLists.txt b/test/spgrid_test/CMakeLists.txt index 3048418a..61e36ccb 100644 --- a/test/spgrid_test/CMakeLists.txt +++ b/test/spgrid_test/CMakeLists.txt @@ -11,7 +11,7 @@ target_compile_features(muda_spgrid_test PRIVATE cxx_std_20) set_target_properties(muda_spgrid_test PROPERTIES CUDA_SEPARABLE_COMPILATION ON) target_include_directories(muda_spgrid_test PRIVATE "${PROJECT_SOURCE_DIR}/test" - "${PROJECT_SOURCE_DIR}/external") + "${PROJECT_SOURCE_DIR}/third_party") target_link_libraries(muda_spgrid_test PRIVATE muda Eigen3::Eigen) add_test(NAME muda_spgrid_test COMMAND muda_spgrid_test) From d6e8c60c6af4b77068eb5e8b23c8c9501c15f2a8 Mon Sep 17 00:00:00 2001 From: zhiguo Date: Mon, 22 Dec 2025 15:57:44 +0800 Subject: [PATCH 15/15] update --- test/eigen_test/CMakeLists.txt | 2 +- test/linear_system_test/CMakeLists.txt | 2 +- test/unit_test/CMakeLists.txt | 2 +- 3 files changed, 3 insertions(+), 3 deletions(-) diff --git a/test/eigen_test/CMakeLists.txt b/test/eigen_test/CMakeLists.txt index 3f68f249..ea278c2c 100644 --- a/test/eigen_test/CMakeLists.txt +++ b/test/eigen_test/CMakeLists.txt @@ -7,7 +7,7 @@ target_compile_features(muda_eigen_test PRIVATE cxx_std_20) set_target_properties(muda_eigen_test PROPERTIES CUDA_SEPARABLE_COMPILATION ON) target_include_directories(muda_eigen_test PRIVATE "${PROJECT_SOURCE_DIR}/test" - "${PROJECT_SOURCE_DIR}/external") + "${PROJECT_SOURCE_DIR}/third_party") target_link_libraries(muda_eigen_test PRIVATE muda Eigen3::Eigen) source_group(TREE "${PROJECT_SOURCE_DIR}/test" PREFIX "test" FILES ${SOURCE}) source_group(TREE "${PROJECT_SOURCE_DIR}/src" PREFIX "src" FILES ${MUDA_HEADER_FILES}) diff --git a/test/linear_system_test/CMakeLists.txt b/test/linear_system_test/CMakeLists.txt index f7c8e64f..bbe5fd7b 100644 --- a/test/linear_system_test/CMakeLists.txt +++ b/test/linear_system_test/CMakeLists.txt @@ -10,7 +10,7 @@ target_compile_features(muda_linear_system_test PRIVATE cxx_std_20) set_target_properties(muda_linear_system_test PROPERTIES CUDA_SEPARABLE_COMPILATION ON) target_include_directories(muda_linear_system_test PRIVATE "${PROJECT_SOURCE_DIR}/test" - "${PROJECT_SOURCE_DIR}/external") + "${PROJECT_SOURCE_DIR}/third_party") target_link_libraries(muda_linear_system_test PRIVATE muda cusparse cublas cusolver Eigen3::Eigen) source_group(TREE "${PROJECT_SOURCE_DIR}/test" PREFIX "test" FILES ${SOURCE}) source_group(TREE "${PROJECT_SOURCE_DIR}/src" PREFIX "src" FILES ${MUDA_HEADER_FILES}) diff --git a/test/unit_test/CMakeLists.txt b/test/unit_test/CMakeLists.txt index 6f1b86bc..79fb23a6 100644 --- a/test/unit_test/CMakeLists.txt +++ b/test/unit_test/CMakeLists.txt @@ -10,7 +10,7 @@ target_compile_features(muda_unit_test PRIVATE cxx_std_20) set_target_properties(muda_unit_test PROPERTIES CUDA_SEPARABLE_COMPILATION ON) target_include_directories(muda_unit_test PRIVATE "${PROJECT_SOURCE_DIR}/test" - "${PROJECT_SOURCE_DIR}/external") + "${PROJECT_SOURCE_DIR}/third_party") target_link_libraries(muda_unit_test PRIVATE muda cusparse cublas cusolver Eigen3::Eigen) target_compile_definitions(muda_unit_test PRIVATE "-DMUDA_TEST_DATA_DIR=R\"(${PROJECT_SOURCE_DIR}/test/data)\"") source_group(TREE "${PROJECT_SOURCE_DIR}/test" PREFIX "test" FILES ${SOURCE})