Skip to content
This repository was archived by the owner on Nov 17, 2025. It is now read-only.

Feature: use hipMemset to improve kernel sync performance which can be faster than hipMemcpy #39

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

Feature: use hipMemset to improve kernel sync performance which can be faster than hipMemcpy #39

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
7 changes: 6 additions & 1 deletion src/usparse/usparse_utils.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -21,11 +21,12 @@
}

struct sync_handle {
static constexpr int SYNC_POLICY = 2;
static constexpr int SYNC_POLICY = 3;

static constexpr int SYNC_POLICY_DEVICE_SYNC = 0;
static constexpr int SYNC_POLICY_EVENT_SYNC = 1;
static constexpr int SYNC_POLICY_MEMCOPY_SYNC = 2;
static constexpr int SYNC_POLICY_MEMSET_SYNC = 3;

inline sync_handle() {
if (SYNC_POLICY == SYNC_POLICY_DEVICE_SYNC) {
Expand All @@ -34,6 +35,8 @@ struct sync_handle {
hipEventCreate(&sync_kernel_event);
} else if (SYNC_POLICY == SYNC_POLICY_MEMCOPY_SYNC) {
mem_malloc();
} else if (SYNC_POLICY == SYNC_POLICY_MEMSET_SYNC) {
HIP_CHECK(hipMalloc((int **)&dev_data, 1 * sizeof(int)));
}
}

Expand All @@ -45,6 +48,8 @@ struct sync_handle {
hipEventSynchronize(sync_kernel_event);
} else if (SYNC_POLICY == SYNC_POLICY_MEMCOPY_SYNC) {
mem_copy();
} else if (SYNC_POLICY == SYNC_POLICY_MEMSET_SYNC) {
hipMemset(dev_data, 0, sizeof(int)); // TODO: HIP_CHECK
}
}

Expand Down