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

Remove cuda event deadlocking issues in device mr tests #1097

Merged
merged 6 commits into from
Sep 27, 2022
Merged

Remove cuda event deadlocking issues in device mr tests #1097

Changes from 1 commit
Commits
Show all changes
6 commits
Select commit Hold shift + click to select a range
7309481
Workaround for cuda event deadlocking issues in device mr tests
robertmaynard Sep 22, 2022
f5aa46b
remove unneded cv notify
robertmaynard Sep 22, 2022
2026411
Address style issues brought up in code review
robertmaynard Sep 23, 2022
3145417
Remove testing code
robertmaynard Sep 23, 2022
cff675d
Address style issues brought up in code review
robertmaynard Sep 23, 2022
ce11027
Remove unneeded unlock as destructor does it
robertmaynard Sep 26, 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
21 changes: 16 additions & 5 deletions tests/mr/device/mr_multithreaded_tests.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -179,6 +179,7 @@ void allocate_loop(rmm::mr::device_memory_resource* mr,
std::size_t num_allocations,
std::list<allocation>& allocations,
std::mutex& mtx,
std::condition_variable& cv,
robertmaynard marked this conversation as resolved.
Show resolved Hide resolved
cudaEvent_t& event,
rmm::cuda_stream_view stream)
{
Expand All @@ -195,34 +196,42 @@ void allocate_loop(rmm::mr::device_memory_resource* mr,
RMM_CUDA_TRY(cudaEventRecord(event, stream.value()));
allocations.emplace_back(ptr, size);
}
cv.notify_one();
}
// Work around for threads going away before cudaEvent has finished async processing
cudaEventSynchronize(event);
}

void deallocate_loop(rmm::mr::device_memory_resource* mr,
std::size_t num_allocations,
std::list<allocation>& allocations,
std::mutex& mtx,
std::condition_variable& cv,
cudaEvent_t& event,
rmm::cuda_stream_view stream)
{
for (std::size_t i = 0; i < num_allocations;) {
std::lock_guard<std::mutex> lock(mtx);
if (allocations.empty()) { continue; }
i++;
for (std::size_t i = 0; i < num_allocations; i++) {
std::unique_lock lk(mtx);
robertmaynard marked this conversation as resolved.
Show resolved Hide resolved
cv.wait(lk, [&allocations] { return !allocations.empty(); });
RMM_CUDA_TRY(cudaStreamWaitEvent(stream.value(), event));
allocation alloc = allocations.front();
allocations.pop_front();
mr->deallocate(alloc.ptr, alloc.size, stream);
lk.unlock();
cv.notify_one();
robertmaynard marked this conversation as resolved.
Show resolved Hide resolved
}
}

// Work around for threads going away before cudaEvent has finished async processing
cudaEventSynchronize(event);
}
void test_allocate_free_different_threads(rmm::mr::device_memory_resource* mr,
rmm::cuda_stream_view streamA,
rmm::cuda_stream_view streamB)
{
constexpr std::size_t num_allocations{100};

std::mutex mtx;
std::condition_variable cv;
robertmaynard marked this conversation as resolved.
Show resolved Hide resolved
std::list<allocation> allocations;
cudaEvent_t event;

Expand All @@ -233,6 +242,7 @@ void test_allocate_free_different_threads(rmm::mr::device_memory_resource* mr,
num_allocations,
std::ref(allocations),
std::ref(mtx),
std::ref(cv),
std::ref(event),
streamA);

Expand All @@ -241,6 +251,7 @@ void test_allocate_free_different_threads(rmm::mr::device_memory_resource* mr,
num_allocations,
std::ref(allocations),
std::ref(mtx),
std::ref(cv),
std::ref(event),
streamB);

Expand Down