Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

[FEA] Support for half-float mixed precise in brute-force #2382

Merged
merged 20 commits into from
Aug 22, 2024
Merged

[FEA] Support for half-float mixed precise in brute-force #2382

Show file tree
Hide file tree
Changes from 1 commit
Commits
Show all changes
20 commits
Select commit Hold shift + click to select a range
0780518
[FEA] Support for half-float mixed precise in brute-force
rhdong Jul 17, 2024
0e9e70b
Merge branch 'branch-24.08' into half_knn
rhdong Jul 17, 2024
3dd5fdf
recover the test cases
rhdong Jul 17, 2024
eb3b8a6
try to fix sddmm error in conda-cpp-tests
rhdong Jul 18, 2024
60ff874
check sparse
rhdong Jul 19, 2024
6103fd1
fix the Floating-point exception
rhdong Jul 19, 2024
d744d59
Merge remote-tracking branch 'origin/branch-24.08' into half_knn
rhdong Jul 19, 2024
efec888
conditionally skip sddmm test cases
rhdong Jul 19, 2024
a55f100
conditionally skip sddmm test cases(2nd try)
rhdong Jul 22, 2024
920997f
half2float -> to_float
rhdong Jul 22, 2024
f2f21c4
Merge branch 'branch-24.08' into half_knn
rhdong Jul 23, 2024
43ee640
CI fix caused by half2float -> to_float
rhdong Jul 23, 2024
11a9988
Merge branch 'branch-24.10' into half_knn
rhdong Jul 26, 2024
aa6da25
Merge branch 'branch-24.10' into half_knn
rhdong Aug 1, 2024
3d8b7ca
update masked_matmul to support mixed precision
rhdong Aug 1, 2024
34dbde8
fix CI fail by skipping low version cuSparse
rhdong Aug 1, 2024
10e5204
fix: illegal instruction by `masked_matmul` on V100 with a dim not mu…
rhdong Aug 9, 2024
4046bf8
Merge branch 'branch-24.10' into half_knn
rhdong Aug 13, 2024
1ed4801
[Fix] bug of `masked_matmul` on non-2-power of `num_cols`
rhdong Aug 13, 2024
4d55c57
[Fix] compilation error of `gemm` & `transpose` under cuml
rhdong Aug 21, 2024
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
Prev Previous commit
Next Next commit
fix: illegal instruction by masked_matmul on V100 with a dim not mu… 
…ltiple of 32
  • Loading branch information
@rhdong
rhdong committed Aug 9, 2024
commit 10e5204cc31a06a77faa3623e9469fc9a8c086c0
1 change: 1 addition & 0 deletions cpp/include/raft/sparse/distance/detail/utils.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -84,6 +84,7 @@ RAFT_KERNEL faster_dot_on_csr_kernel(dot_t* __restrict__ dot,
l_dot_ += s_A[k] * __ldcg(B_col + k);
}
}
__syncthreads();
l_dot_ += __shfl_down_sync(0xffffffff, l_dot_, 16);
l_dot_ += __shfl_down_sync(0xffff, l_dot_, 8);
l_dot_ += __shfl_down_sync(0xff, l_dot_, 4);
Expand Down
18 changes: 12 additions & 6 deletions cpp/test/sparse/masked_matmul.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -363,33 +363,39 @@ using MaskedMatmulTestH = MaskedMatmulTest<half, float, int>;
TEST_P(MaskedMatmulTestH, Result) { Run(); }

const std::vector<MaskedMatmulInputs<float, float, int>> sddmm_inputs_f = {
{0.0001f, 1, 255, 131, 0.01, 1234ULL},
{0.0001f, 9, 255, 13100, 0.009, 1234ULL},
{0.0001f, 10, 5, 32, 0.1, 1234ULL},
{0.0001f, 1024, 32, 1024, 0.1, 1234ULL},
{0.0003f, 32, 1024, 1024, 0.2, 1234ULL},
{0.001f, 1024, 1024, 1024, 0.19, 1234ULL},
{0.001f, 1023, 1023, 1023, 0.19, 1234ULL},
{0.0001f, 1024, 1024, 32, 0.3, 1234ULL},
{0.0001f, 1024, 32, 1024, 0.4, 1234ULL},
{0.0003f, 32, 1024, 1024, 0.19, 1234ULL},
{0.0003f, 31, 1025, 1025, 0.19, 1234ULL},
{0.001f, 1024, 1024, 1024, 0.1, 1234ULL}};

const std::vector<MaskedMatmulInputs<double, double, int>> sddmm_inputs_d = {
{0.0001f, 1, 255, 131, 0.01, 1234ULL},
{0.0001f, 9, 255, 13100, 0.009, 1234ULL},
{0.0001f, 10, 5, 32, 0.01, 1234ULL},
{0.0001f, 1024, 32, 1024, 0.1, 1234ULL},
{0.0001f, 32, 1024, 1024, 0.2, 1234ULL},
{0.0001f, 1024, 1024, 1024, 0.19, 1234ULL},
{0.0001f, 1023, 1023, 1023, 0.19, 1234ULL},
{0.0001f, 1024, 1024, 32, 0.3, 1234ULL},
{0.0001f, 1024, 32, 1024, 0.4, 1234ULL},
{0.0001f, 32, 1024, 1024, 0.19, 1234ULL},
{0.0001f, 31, 1025, 1025, 0.19, 1234ULL},
{0.0001f, 1024, 1024, 1024, 0.1, 1234ULL}};

const std::vector<MaskedMatmulInputs<half, float, int>> sddmm_inputs_h = {
{0.0001f, 1, 255, 131, 0.01, 1234ULL},
{0.0001f, 9, 255, 13101, 0.009, 1234ULL},
{0.0001f, 10, 5, 32, 0.1, 1234ULL},
{0.0001f, 1024, 32, 1024, 0.1, 1234ULL},
{0.0003f, 32, 1024, 1024, 0.2, 1234ULL},
{0.001f, 1024, 1024, 1024, 0.19, 1234ULL},
{0.001f, 1023, 1023, 1023, 0.19, 1234ULL},
{0.0001f, 1024, 1024, 32, 0.3, 1234ULL},
{0.0001f, 1024, 32, 1024, 0.4, 1234ULL},
{0.0003f, 32, 1024, 1024, 0.19, 1234ULL},
{0.0003f, 31, 1025, 1025, 0.19, 1234ULL},
{0.001f, 1024, 1024, 1024, 0.1, 1234ULL}};

INSTANTIATE_TEST_CASE_P(MaskedMatmulTest, MaskedMatmulTestF, ::testing::ValuesIn(sddmm_inputs_f));
Expand Down
Loading