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

Adding fused_l2_nn_argmin wrapper to Pylibraft #924

Merged
merged 7 commits into from
Oct 19, 2022
Merged

Adding fused_l2_nn_argmin wrapper to Pylibraft #924

Changes from 1 commit
Commits
Show all changes
7 commits
Select commit Hold shift + click to select a range
0e15a2d
Adding fused l2 nn argmin primitive for pylibraft
cjnolet Oct 18, 2022
aaf75bb
Adding a few more test caseds.
cjnolet Oct 18, 2022
b7a5835
Cleaning up style and prints
cjnolet Oct 18, 2022
e6731dd
CLeaning up
cjnolet Oct 18, 2022
e9a86ec
Style
cjnolet Oct 18, 2022
7c3f50f
Removing accidental checkin
cjnolet Oct 18, 2022
6d2b9d0
Fixing flake8 style
cjnolet Oct 19, 2022
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
Style
  • Loading branch information
@cjnolet
cjnolet committed Oct 18, 2022
commit e9a86ec8a6a3b5da4707f8774fc83f4128661b22
130 changes: 65 additions & 65 deletions cpp/test/util/fast_int_div.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -15,100 +15,100 @@
*/

#include "../test_utils.h"
#include <raft/util/fast_int_div.cuh>
#include <raft/util/cudart_utils.hpp>
#include <raft/util/fast_int_div.cuh>

#include <rmm/device_uvector.hpp>

#include <gtest/gtest.h>

namespace raft::util {

TEST(FastIntDiv, CpuTest)
TEST(FastIntDiv, CpuTest)
{
for (int i = 0; i < 100; ++i) {
for (int i = 0; i < 100; ++i) {
// get a positive divisor
int divisor;
do {
divisor = rand();
} while (divisor <= 0);
FastIntDiv fid(divisor);
// run it against a few random numbers and compare the outputs
for (int i = 0; i < 10000; ++i) {
auto num = rand();
auto correct = num / divisor;
auto computed = num / fid;
ASSERT_EQ(correct, computed) << " divisor=" << divisor << " num=" << num;
num = rand();
correct = num % divisor;
computed = num % fid;
ASSERT_EQ(correct, computed) << " divisor=" << divisor << " num=" << num;
num = -num;
correct = num / divisor;
computed = num / fid;
ASSERT_EQ(correct, computed) << " divisor=" << divisor << " num=" << num;
num = rand();
correct = num % divisor;
computed = num % fid;
ASSERT_EQ(correct, computed) << " divisor=" << divisor << " num=" << num;
}
}
divisor = rand();
} while (divisor <= 0);
FastIntDiv fid(divisor);
// run it against a few random numbers and compare the outputs
for (int i = 0; i < 10000; ++i) {
auto num = rand();
auto correct = num / divisor;
auto computed = num / fid;
ASSERT_EQ(correct, computed) << " divisor=" << divisor << " num=" << num;
num = rand();
correct = num % divisor;
computed = num % fid;
ASSERT_EQ(correct, computed) << " divisor=" << divisor << " num=" << num;
num = -num;
correct = num / divisor;
computed = num / fid;
ASSERT_EQ(correct, computed) << " divisor=" << divisor << " num=" << num;
num = rand();
correct = num % divisor;
computed = num % fid;
ASSERT_EQ(correct, computed) << " divisor=" << divisor << " num=" << num;
}
}
}

__global__ void fastIntDivTestKernel(
int* computed, int* correct, const int* in, FastIntDiv fid, int divisor, int len)
int* computed, int* correct, const int* in, FastIntDiv fid, int divisor, int len)
{
auto tid = threadIdx.x + blockIdx.x * blockDim.x;
if (tid < len) {
computed[tid] = in[tid] % fid;
correct[tid] = in[tid] % divisor;
computed[len + tid] = -in[tid] % fid;
correct[len + tid] = -in[tid] % divisor;
}
auto tid = threadIdx.x + blockIdx.x * blockDim.x;
if (tid < len) {
computed[tid] = in[tid] % fid;
correct[tid] = in[tid] % divisor;
computed[len + tid] = -in[tid] % fid;
correct[len + tid] = -in[tid] % divisor;
}
}

TEST(FastIntDiv, GpuTest)
{
cudaStream_t stream = 0;
RAFT_CUDA_TRY(cudaStreamCreate(&stream));
cudaStream_t stream = 0;
RAFT_CUDA_TRY(cudaStreamCreate(&stream));

static const int len = 100000;
static const int TPB = 128;
rmm::device_uvector<int> computed(len * 2, stream);
rmm::device_uvector<int> correct(len * 2, stream);
rmm::device_uvector<int> in(len, stream);
for (int i = 0; i < 100; ++i) {
// get a positive divisor
int divisor;
do {
divisor = rand();
} while (divisor <= 0);
FastIntDiv fid(divisor);
// run it against a few random numbers and compare the outputs
std::vector<int> h_in(len);
for (int i = 0; i < len; ++i) {
h_in[i] = rand();
}
raft::update_device(in.data(), h_in.data(), len, stream);
int nblks = raft::ceildiv(len, TPB);
fastIntDivTestKernel<<<nblks, TPB, 0, 0>>>(
computed.data(), correct.data(), in.data(), fid, divisor, len);
RAFT_CUDA_TRY(cudaStreamSynchronize(0));
ASSERT_TRUE(devArrMatch(correct.data(), computed.data(), len * 2, raft::Compare<int>()))
<< " divisor=" << divisor;
}
static const int len = 100000;
static const int TPB = 128;
rmm::device_uvector<int> computed(len * 2, stream);
rmm::device_uvector<int> correct(len * 2, stream);
rmm::device_uvector<int> in(len, stream);
for (int i = 0; i < 100; ++i) {
// get a positive divisor
int divisor;
do {
divisor = rand();
} while (divisor <= 0);
FastIntDiv fid(divisor);
// run it against a few random numbers and compare the outputs
std::vector<int> h_in(len);
for (int i = 0; i < len; ++i) {
h_in[i] = rand();
}
raft::update_device(in.data(), h_in.data(), len, stream);
int nblks = raft::ceildiv(len, TPB);
fastIntDivTestKernel<<<nblks, TPB, 0, 0>>>(
computed.data(), correct.data(), in.data(), fid, divisor, len);
RAFT_CUDA_TRY(cudaStreamSynchronize(0));
ASSERT_TRUE(devArrMatch(correct.data(), computed.data(), len * 2, raft::Compare<int>()))
<< " divisor=" << divisor;
}
}

FastIntDiv dummyFunc(int num)
{
FastIntDiv fd(num);
return fd;
FastIntDiv fd(num);
return fd;
}

TEST(FastIntDiv, IncorrectUsage)
{
ASSERT_THROW(dummyFunc(-1), raft::exception);
ASSERT_THROW(dummyFunc(0), raft::exception);
ASSERT_THROW(dummyFunc(-1), raft::exception);
ASSERT_THROW(dummyFunc(0), raft::exception);
}

} // namespace raft::util