Skip to content

Commit

Permalink
revert the sparsity value
Browse files Browse the repository at this point in the history
  • Loading branch information
rhdong committed Dec 4, 2024
1 parent d9a42d2 commit c42cbf7
Show file tree
Hide file tree
Showing 2 changed files with 30 additions and 29 deletions.
7 changes: 4 additions & 3 deletions cpp/include/raft/sparse/convert/detail/bitmap_to_csr.cuh
Original file line number Diff line number Diff line change
Expand Up @@ -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<bitmap_t>;
using BlockReduce = cub::BlockReduce<index_t, bitmap_to_csr_tpb>;

__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;
Expand Down Expand Up @@ -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<bitmap_t>;
using BlockScan = cub::BlockScan<int, bitmap_to_csr_tpb>;

using BlockScan = cub::BlockScan<int, bitmap_to_csr_tpb>;
__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.
Expand Down
52 changes: 26 additions & 26 deletions cpp/test/sparse/convert_csr.cu
Original file line number Diff line number Diff line change
Expand Up @@ -249,7 +249,7 @@ class BitmapToCSRTest : public ::testing::TestWithParam<BitmapToCSRInputs<index_
index_t create_sparse_matrix(index_t m, index_t n, float sparsity, std::vector<bitmap_t>& bitmap)
{
index_t total = static_cast<index_t>(m * n);
index_t num_ones = static_cast<index_t>((total * 1.0f) * sparsity);
index_t num_ones = static_cast<index_t>((total * 1.0f) * (1.0f - sparsity));
index_t res = num_ones;

for (auto& item : bitmap) {
Expand Down Expand Up @@ -431,35 +431,35 @@ TEST_P(BitmapToCSRTestLOnLargeSize, Result) { Run(); }

template <typename index_t>
const std::vector<BitmapToCSRInputs<index_t>> 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 <typename index_t>
const std::vector<BitmapToCSRInputs<index_t>> 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,
Expand Down

0 comments on commit c42cbf7

Please sign in to comment.