From c42cbf74159d1629dd986484b1f884d88de49b1b Mon Sep 17 00:00:00 2001 From: rhdong Date: Wed, 4 Dec 2024 08:54:36 -0800 Subject: [PATCH] revert the sparsity value --- .../sparse/convert/detail/bitmap_to_csr.cuh | 7 +-- cpp/test/sparse/convert_csr.cu | 52 +++++++++---------- 2 files changed, 30 insertions(+), 29 deletions(-) diff --git a/cpp/include/raft/sparse/convert/detail/bitmap_to_csr.cuh b/cpp/include/raft/sparse/convert/detail/bitmap_to_csr.cuh index 1b500f30c7..dbe788b15c 100644 --- a/cpp/include/raft/sparse/convert/detail/bitmap_to_csr.cuh +++ b/cpp/include/raft/sparse/convert/detail/bitmap_to_csr.cuh @@ -55,11 +55,12 @@ RAFT_KERNEL __launch_bounds__(bitmap_to_csr_tpb) calc_nnz_by_rows_kernel(const b { using mutable_bitmap_t = typename std::remove_const_t; using BlockReduce = cub::BlockReduce; + __shared__ typename BlockReduce::TempStorage reduce_storage; constexpr index_t BITS_PER_BITMAP = sizeof(bitmap_t) * 8; - auto tid = threadIdx.x; + const auto tid = threadIdx.x; const auto row = blockIdx.x; const auto num_sub_cols = gridDim.y; @@ -150,11 +151,11 @@ RAFT_KERNEL __launch_bounds__(bitmap_to_csr_tpb) constexpr index_t BITS_PER_BITMAP = sizeof(bitmap_t) * 8; using mutable_bitmap_t = typename std::remove_const_t; + using BlockScan = cub::BlockScan; - using BlockScan = cub::BlockScan; __shared__ typename BlockScan::TempStorage scan_storage; - auto tid = threadIdx.x; + const auto tid = threadIdx.x; const auto row = blockIdx.x; // Ensure the HBM allocated for CSR values is sufficient to handle all non-zero bitmap bits. diff --git a/cpp/test/sparse/convert_csr.cu b/cpp/test/sparse/convert_csr.cu index ae08c3cfe2..df6b4352e9 100644 --- a/cpp/test/sparse/convert_csr.cu +++ b/cpp/test/sparse/convert_csr.cu @@ -249,7 +249,7 @@ class BitmapToCSRTest : public ::testing::TestWithParam& bitmap) { index_t total = static_cast(m * n); - index_t num_ones = static_cast((total * 1.0f) * sparsity); + index_t num_ones = static_cast((total * 1.0f) * (1.0f - sparsity)); index_t res = num_ones; for (auto& item : bitmap) { @@ -431,35 +431,35 @@ TEST_P(BitmapToCSRTestLOnLargeSize, Result) { Run(); } template const std::vector> bitmaptocsr_inputs = { - {0, 0, 0.2, false}, - {10, 32, 0.4, false}, - {10, 3, 0.2, false}, - {32, 1024, 0.4, false}, - {1024, 1048576, 0.01, false}, - {1024, 1024, 0.4, false}, - {64 * 1024 + 10, 2, 0.3, false}, // 64K + 10 is slightly over maximum of blockDim.y - {16, 16, 0.3, false}, // No peeling-remainder - {17, 16, 0.3, false}, // Check peeling-remainder - {18, 16, 0.3, false}, // Check peeling-remainder - {32 + 9, 33, 0.2, false}, // Check peeling-remainder - {2, 33, 0.2, false}, // Check peeling-remainder - {0, 0, 0.2, true}, - {10, 32, 0.4, true}, - {10, 3, 0.2, true}, - {32, 1024, 0.4, true}, - {1024, 1048576, 0.01, true}, - {1024, 1024, 0.4, true}, - {64 * 1024 + 10, 2, 0.3, true}, // 64K + 10 is slightly over maximum of blockDim.y - {16, 16, 0.3, true}, // No peeling-remainder - {17, 16, 0.3, true}, // Check peeling-remainder - {18, 16, 0.3, true}, // Check peeling-remainder - {32 + 9, 33, 0.2, true}, // Check peeling-remainder - {2, 33, 0.2, true}, // Check peeling-remainder + {0, 0, 0.8, false}, + {10, 32, 0.6, false}, + {10, 3, 0.8, false}, + {32, 1024, 0.6, false}, + {1024, 1048576, 0.99, false}, + {1024, 1024, 0.6, false}, + {64 * 1024 + 10, 2, 0.7, false}, // 64K + 10 is slightly over maximum of blockDim.y + {16, 16, 0.7, false}, // No peeling-remainder + {17, 16, 0.7, false}, // Check peeling-remainder + {18, 16, 0.7, false}, // Check peeling-remainder + {32 + 9, 33, 0.8, false}, // Check peeling-remainder + {2, 33, 0.8, false}, // Check peeling-remainder + {0, 0, 0.8, true}, + {10, 32, 0.6, true}, + {10, 3, 0.8, true}, + {32, 1024, 0.6, true}, + {1024, 1048576, 0.99, true}, + {1024, 1024, 0.6, true}, + {64 * 1024 + 10, 2, 0.7, true}, // 64K + 10 is slightly over maximum of blockDim.y + {16, 16, 0.7, true}, // No peeling-remainder + {17, 16, 0.7, true}, // Check peeling-remainder + {18, 16, 0.7, true}, // Check peeling-remainder + {32 + 9, 33, 0.8, true}, // Check peeling-remainder + {2, 33, 0.8, true}, // Check peeling-remainder }; template const std::vector> bitmaptocsr_large_inputs = { - {100, 100000000, 0.01, true}, {100, 100000000, 0.05, false}, {100, 100000000 + 17, 0.05, false}}; + {100, 100000000, 0.99, true}, {100, 100000000, 0.95, false}, {100, 100000000 + 17, 0.95, false}}; INSTANTIATE_TEST_CASE_P(SparseConvertCSRTest, BitmapToCSRTestI,