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

Revert the commit 9e02f53 to make it a standalone PR. #202

Merged
merged 1 commit into from
Aug 21, 2024
Merged

Revert the commit 9e02f53 to make it a standalone PR. #202

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
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
18 changes: 9 additions & 9 deletions include/merlin/core_kernels/group_lock_kernels.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -21,15 +21,15 @@ namespace nv {
namespace merlin {
namespace group_lock {

static __global__ void init_kernel(
__global__ void init_kernel(
cuda::atomic<int, cuda::thread_scope_device>* update_count,
cuda::atomic<int, cuda::thread_scope_device>* read_count,
cuda::atomic<bool, cuda::thread_scope_device>* unique_flag) {
new (update_count) cuda::atomic<int, cuda::thread_scope_device>{0};
new (read_count) cuda::atomic<int, cuda::thread_scope_device>{0};
new (unique_flag) cuda::atomic<bool, cuda::thread_scope_device>{false};
}
static __global__ void lock_read_kernel(
__global__ void lock_read_kernel(
cuda::atomic<int, cuda::thread_scope_device>* update_count,
cuda::atomic<int, cuda::thread_scope_device>* read_count) {
for (;;) {
Expand All @@ -43,12 +43,12 @@ static __global__ void lock_read_kernel(
}
}

static __global__ void unlock_read_kernel(
__global__ void unlock_read_kernel(
cuda::atomic<int, cuda::thread_scope_device>* read_count) {
read_count->fetch_sub(1, cuda::std::memory_order_relaxed);
}

static __global__ void lock_update_kernel(
__global__ void lock_update_kernel(
cuda::atomic<int, cuda::thread_scope_device>* update_count,
cuda::atomic<int, cuda::thread_scope_device>* read_count) {
for (;;) {
Expand All @@ -62,12 +62,12 @@ static __global__ void lock_update_kernel(
}
}

static __global__ void unlock_update_kernel(
__global__ void unlock_update_kernel(
cuda::atomic<int, cuda::thread_scope_device>* update_count) {
update_count->fetch_sub(1, cuda::std::memory_order_relaxed);
}

static __global__ void lock_update_read_kernel(
__global__ void lock_update_read_kernel(
cuda::atomic<int, cuda::thread_scope_device>* update_count,
cuda::atomic<int, cuda::thread_scope_device>* read_count,
cuda::atomic<bool, cuda::thread_scope_device>* unique_flag) {
Expand Down Expand Up @@ -101,7 +101,7 @@ static __global__ void lock_update_read_kernel(
}
}

static __global__ void unlock_update_read_kernel(
__global__ void unlock_update_read_kernel(
cuda::atomic<int, cuda::thread_scope_device>* update_count,
cuda::atomic<int, cuda::thread_scope_device>* read_count,
cuda::atomic<bool, cuda::thread_scope_device>* unique_flag) {
Expand All @@ -110,12 +110,12 @@ static __global__ void unlock_update_read_kernel(
unique_flag->store(false, cuda::std::memory_order_relaxed);
}

static __global__ void update_count_kernel(
__global__ void update_count_kernel(
int* counter, cuda::atomic<int, cuda::thread_scope_device>* update_count) {
*counter = update_count->load(cuda::std::memory_order_relaxed);
}

static __global__ void read_count_kernel(
__global__ void read_count_kernel(
int* counter, cuda::atomic<int, cuda::thread_scope_device>* read_count) {
*counter = read_count->load(cuda::std::memory_order_relaxed);
}
Expand Down
4 changes: 2 additions & 2 deletions include/merlin/types.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -54,7 +54,7 @@ constexpr uint64_t MAX_SCORE = UINT64_C(0xFFFFFFFFFFFFFFFF);
constexpr uint64_t EMPTY_SCORE = UINT64_C(0);
constexpr uint64_t IGNORED_GLOBAL_EPOCH = UINT64_C(0xFFFFFFFFFFFFFFFF);

static uint64_t EMPTY_KEY_CPU = DEFAULT_EMPTY_KEY;
uint64_t EMPTY_KEY_CPU = DEFAULT_EMPTY_KEY;
__constant__ uint64_t EMPTY_KEY = DEFAULT_EMPTY_KEY;
__constant__ uint64_t RECLAIM_KEY = DEFAULT_RECLAIM_KEY;
__constant__ uint64_t LOCKED_KEY = DEFAULT_LOCKED_KEY;
Expand All @@ -76,7 +76,7 @@ __forceinline__ __device__ bool IS_VACANT_KEY(K key) {
return (VACANT_KEY_MASK_1 & key) == VACANT_KEY_MASK_2;
}

static cudaError_t init_reserved_keys(int index) {
cudaError_t init_reserved_keys(int index) {
if (index < 1 || index > MAX_RESERVED_KEY_BIT) {
// index = 0 is the default,
// index = 62 is the maximum index can be set for reserved keys.
Expand Down
3 changes: 1 addition & 2 deletions include/merlin/utils.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -314,8 +314,7 @@ inline void free_pointers(cudaStream_t stream, int n, ...) {
va_end(args);
}

static __global__ void memset64bitKernel(void* devPtr, uint64_t value,
size_t count) {
__global__ void memset64bitKernel(void* devPtr, uint64_t value, size_t count) {
size_t idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < count) {
static_cast<uint64_t*>(devPtr)[idx] = value;
Expand Down
Loading