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

Jump to bottom

[Fix] a standalone PR to fixup multi-defination issue across multiple… #204

Open
wants to merge 1 commit into
base: master
Choose a base branch
from
Open

[Fix] a standalone PR to fixup multi-defination issue across multiple… #204

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 {

__global__ void init_kernel(
static __global__ void init_kernel(
Copy link
Collaborator

@Lifann Lifann Aug 22, 2024 •
edited
Loading

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Is there relative issue or code snippet to reproduce the multi-definition problem?

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};
}
__global__ void lock_read_kernel(
static __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 @@ __global__ void lock_read_kernel(
}
}

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

__global__ void lock_update_kernel(
static __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 @@ __global__ void lock_update_kernel(
}
}

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

__global__ void lock_update_read_kernel(
static __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 @@ __global__ void lock_update_read_kernel(
}
}

__global__ void unlock_update_read_kernel(
static __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 @@ __global__ void unlock_update_read_kernel(
unique_flag->store(false, cuda::std::memory_order_relaxed);
}

__global__ void update_count_kernel(
static __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);
}

__global__ void read_count_kernel(
static __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);

uint64_t EMPTY_KEY_CPU = DEFAULT_EMPTY_KEY;
static 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;
}

cudaError_t init_reserved_keys(int index) {
static 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: 2 additions & 1 deletion include/merlin/utils.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -314,7 +314,8 @@ inline void free_pointers(cudaStream_t stream, int n, ...) {
va_end(args);
}

__global__ void memset64bitKernel(void* devPtr, uint64_t value, size_t count) {
static __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