From 82768227b23ce98c074f75bf80a6320cbedda5bd Mon Sep 17 00:00:00 2001 From: "jiashuai.shi@qq.com" Date: Tue, 26 May 2026 02:47:40 +0800 Subject: [PATCH] style(core): apply clang-format to source files Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> --- include/spmv/common.h | 4 ++-- include/spmv/cuda_buffer.h | 3 +-- include/spmv/cuda_compat.h | 6 ++---- include/spmv/spmv.h | 4 ++-- src/bandwidth.cpp | 1 + src/csr_matrix.cpp | 11 ++++++++--- src/ell_matrix.cpp | 11 ++++++++--- src/internal/csr_device.cpp | 10 ++++++---- src/internal/ell_device.cpp | 4 ++-- src/internal/kernel_selector.cpp | 3 +-- src/internal/kernel_selector.h | 3 +-- src/spmv_context.cpp | 4 ++-- src/spmv_cpu.cpp | 3 ++- src/spmv_kernels.cu | 17 +++++++++-------- 14 files changed, 47 insertions(+), 37 deletions(-) diff --git a/include/spmv/common.h b/include/spmv/common.h index 6920454..4f01285 100644 --- a/include/spmv/common.h +++ b/include/spmv/common.h @@ -1,13 +1,13 @@ #ifndef SPMV_COMMON_H #define SPMV_COMMON_H -#include "cuda_compat.h" - #include #include #include #include +#include "cuda_compat.h" + namespace spmv { /** diff --git a/include/spmv/cuda_buffer.h b/include/spmv/cuda_buffer.h index 38efb25..ff82d2a 100644 --- a/include/spmv/cuda_buffer.h +++ b/include/spmv/cuda_buffer.h @@ -1,13 +1,12 @@ #ifndef SPMV_CUDA_BUFFER_H #define SPMV_CUDA_BUFFER_H -#include "cuda_compat.h" - #include #include #include #include "common.h" +#include "cuda_compat.h" namespace spmv { diff --git a/include/spmv/cuda_compat.h b/include/spmv/cuda_compat.h index d0d92ad..a294f64 100644 --- a/include/spmv/cuda_compat.h +++ b/include/spmv/cuda_compat.h @@ -121,10 +121,8 @@ inline cudaError_t cudaMemset(void* dst, int value, size_t count) { return cudaSuccess; } -inline cudaError_t cudaCreateTextureObject(cudaTextureObject_t* tex, - const cudaResourceDesc*, - const cudaTextureDesc*, - const void*) { +inline cudaError_t cudaCreateTextureObject(cudaTextureObject_t* tex, const cudaResourceDesc*, + const cudaTextureDesc*, const void*) { static cudaTextureObject_t next_texture = 1; if (!tex) { return cudaErrorInvalidValue; diff --git a/include/spmv/spmv.h b/include/spmv/spmv.h index 627afb0..9c01777 100644 --- a/include/spmv/spmv.h +++ b/include/spmv/spmv.h @@ -130,8 +130,8 @@ class SpMVExecutionContext { bool is_texture_bound() const { return tex_x_ != 0; } private: - friend int spmv_prepare_texture(SpMVExecutionContext* context, const float* d_x, size_t x_length, - bool requested, cudaTextureObject_t* tex_out, + friend int spmv_prepare_texture(SpMVExecutionContext* context, const float* d_x, + size_t x_length, bool requested, cudaTextureObject_t* tex_out, bool* use_texture_out); cudaTextureObject_t tex_x_ = 0; diff --git a/src/bandwidth.cpp b/src/bandwidth.cpp index 51dc95d..41f4d86 100644 --- a/src/bandwidth.cpp +++ b/src/bandwidth.cpp @@ -1,4 +1,5 @@ #include "spmv/bandwidth.h" + #include "spmv/cuda_compat.h" #include diff --git a/src/csr_matrix.cpp b/src/csr_matrix.cpp index 11b8a1b..7c96eea 100644 --- a/src/csr_matrix.cpp +++ b/src/csr_matrix.cpp @@ -1,4 +1,3 @@ -#include "internal/csr_device.h" #include "spmv/csr_matrix.h" #include @@ -8,6 +7,8 @@ #include #include +#include "internal/csr_device.h" + namespace spmv { // Public API ----------------------------------------------------------------- @@ -132,9 +133,13 @@ float csr_get_element(const CSRMatrix* mat, int row, int col) { return 0.0f; } -int csr_to_gpu(CSRMatrix* mat) { return csr_upload_device_data(mat); } +int csr_to_gpu(CSRMatrix* mat) { + return csr_upload_device_data(mat); +} -int csr_from_gpu(CSRMatrix* mat) { return csr_download_device_data(mat); } +int csr_from_gpu(CSRMatrix* mat) { + return csr_download_device_data(mat); +} int csr_serialize(const CSRMatrix* mat, const char* filename) { if (!mat || !filename) { diff --git a/src/ell_matrix.cpp b/src/ell_matrix.cpp index 58b9673..e7c9ad8 100644 --- a/src/ell_matrix.cpp +++ b/src/ell_matrix.cpp @@ -1,4 +1,3 @@ -#include "internal/ell_device.h" #include "spmv/ell_matrix.h" #include @@ -7,6 +6,8 @@ #include #include +#include "internal/ell_device.h" + namespace spmv { // Public API ----------------------------------------------------------------- @@ -200,9 +201,13 @@ float ell_get_element(const ELLMatrix* mat, int row, int col) { return 0.0f; } -int ell_to_gpu(ELLMatrix* mat) { return ell_upload_device_data(mat); } +int ell_to_gpu(ELLMatrix* mat) { + return ell_upload_device_data(mat); +} -int ell_from_gpu(ELLMatrix* mat) { return ell_download_device_data(mat); } +int ell_from_gpu(ELLMatrix* mat) { + return ell_download_device_data(mat); +} int ell_serialize(const ELLMatrix* mat, const char* filename) { if (!mat || !filename) { diff --git a/src/internal/csr_device.cpp b/src/internal/csr_device.cpp index f92e5d2..8ca6a9c 100644 --- a/src/internal/csr_device.cpp +++ b/src/internal/csr_device.cpp @@ -124,7 +124,8 @@ int csr_upload_device_data(CSRMatrix* mat) { }; if (mat->nnz > 0) { - cudaError_t err = cudaMalloc(reinterpret_cast(&new_d_values), mat->nnz * sizeof(float)); + cudaError_t err = + cudaMalloc(reinterpret_cast(&new_d_values), mat->nnz * sizeof(float)); if (err != cudaSuccess) { cleanup(); return static_cast(SpMVError::CUDA_MALLOC); @@ -137,15 +138,16 @@ int csr_upload_device_data(CSRMatrix* mat) { } } - cudaError_t err = cudaMalloc(reinterpret_cast(&new_d_row_ptrs), - (mat->num_rows + 1) * sizeof(int)); + cudaError_t err = + cudaMalloc(reinterpret_cast(&new_d_row_ptrs), (mat->num_rows + 1) * sizeof(int)); if (err != cudaSuccess) { cleanup(); return static_cast(SpMVError::CUDA_MALLOC); } if (mat->nnz > 0) { - err = cudaMemcpy(new_d_values, mat->values, mat->nnz * sizeof(float), cudaMemcpyHostToDevice); + err = + cudaMemcpy(new_d_values, mat->values, mat->nnz * sizeof(float), cudaMemcpyHostToDevice); if (err != cudaSuccess) { cleanup(); return static_cast(SpMVError::CUDA_MEMCPY); diff --git a/src/internal/ell_device.cpp b/src/internal/ell_device.cpp index 1ce7aa6..496d4c0 100644 --- a/src/internal/ell_device.cpp +++ b/src/internal/ell_device.cpp @@ -148,8 +148,8 @@ int ell_download_device_data(ELLMatrix* mat) { if (!mat->values || !mat->col_indices) { return static_cast(SpMVError::INVALID_ARGUMENT); } - CUDA_CHECK_MEMCPY( - cudaMemcpy(mat->values, internal->d_values, size * sizeof(float), cudaMemcpyDeviceToHost)); + CUDA_CHECK_MEMCPY(cudaMemcpy(mat->values, internal->d_values, size * sizeof(float), + cudaMemcpyDeviceToHost)); CUDA_CHECK_MEMCPY(cudaMemcpy(mat->col_indices, internal->d_col_indices, size * sizeof(int), cudaMemcpyDeviceToHost)); } diff --git a/src/internal/kernel_selector.cpp b/src/internal/kernel_selector.cpp index 6e1f53b..acfba59 100644 --- a/src/internal/kernel_selector.cpp +++ b/src/internal/kernel_selector.cpp @@ -2,8 +2,7 @@ namespace spmv { -SpMVConfig select_kernel(const CSRStats& stats, int num_cols, - const SpMVThresholds& thresholds) { +SpMVConfig select_kernel(const CSRStats& stats, int num_cols, const SpMVThresholds& thresholds) { SpMVConfig config(SpMVConfig::SCALAR_CSR, DEFAULT_BLOCK_SIZE, false); config.use_texture = (num_cols > thresholds.texture_cols_threshold); diff --git a/src/internal/kernel_selector.h b/src/internal/kernel_selector.h index a2dbc4c..4b44ad8 100644 --- a/src/internal/kernel_selector.h +++ b/src/internal/kernel_selector.h @@ -12,8 +12,7 @@ namespace spmv { * immutable inputs (stats, dimensions, thresholds) and returns a deterministic * SpMVConfig. No global state, no matrix internals. */ -SpMVConfig select_kernel(const CSRStats& stats, int num_cols, - const SpMVThresholds& thresholds); +SpMVConfig select_kernel(const CSRStats& stats, int num_cols, const SpMVThresholds& thresholds); } // namespace spmv diff --git a/src/spmv_context.cpp b/src/spmv_context.cpp index fee90d6..a96343e 100644 --- a/src/spmv_context.cpp +++ b/src/spmv_context.cpp @@ -22,8 +22,8 @@ int spmv_prepare_texture(SpMVExecutionContext* context, const float* d_x, size_t return static_cast(SpMVError::INVALID_ARGUMENT); } - bool needs_rebuild = !context->texture_enabled_ || context->tex_x_ == 0 || context->cached_x_ != d_x || - context->cached_x_length_ != x_length; + bool needs_rebuild = !context->texture_enabled_ || context->tex_x_ == 0 || + context->cached_x_ != d_x || context->cached_x_length_ != x_length; if (needs_rebuild) { context->reset(); diff --git a/src/spmv_cpu.cpp b/src/spmv_cpu.cpp index faa3bd0..6b38e8e 100644 --- a/src/spmv_cpu.cpp +++ b/src/spmv_cpu.cpp @@ -1,8 +1,9 @@ -#include "internal/kernel_selector.h" #include "spmv/spmv.h" #include +#include "internal/kernel_selector.h" + namespace spmv { // Global thresholds for kernel selection (can be tuned per-GPU architecture) diff --git a/src/spmv_kernels.cu b/src/spmv_kernels.cu index d0ab615..716c95f 100644 --- a/src/spmv_kernels.cu +++ b/src/spmv_kernels.cu @@ -1,6 +1,3 @@ -#include "internal/csr_device.h" -#include "internal/ell_device.h" -#include "internal/texture_cache.h" #include "spmv/bandwidth.h" #include "spmv/spmv.h" @@ -8,6 +5,10 @@ #include +#include "internal/csr_device.h" +#include "internal/ell_device.h" +#include "internal/texture_cache.h" + namespace spmv { // ---------- RAII helpers ---------- @@ -400,16 +401,16 @@ SpMVResult spmv_csr(const CSRMatrix* A, const float* d_x, float* d_y, const SpMV return result; } spmv_csr_merge_path_kernel<<>>( - A->num_rows, A->nnz, csr_d_row_ptrs(A), csr_d_col_indices(A), csr_d_values(A), - d_x, tex_x, use_texture, d_y); + A->num_rows, A->nnz, csr_d_row_ptrs(A), csr_d_col_indices(A), csr_d_values(A), d_x, + tex_x, use_texture, d_y); break; } case SpMVConfig::VECTOR_CSR: { int warps_per_block = block_size / 32; int num_warps = (A->num_rows + warps_per_block - 1) / warps_per_block; - spmv_csr_vector_kernel<<>>( - A->num_rows, csr_d_row_ptrs(A), csr_d_col_indices(A), csr_d_values(A), d_x, tex_x, - use_texture, d_y); + spmv_csr_vector_kernel<<>>(A->num_rows, csr_d_row_ptrs(A), + csr_d_col_indices(A), csr_d_values(A), + d_x, tex_x, use_texture, d_y); break; } case SpMVConfig::SCALAR_CSR: