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

[SYCL] Fix launch config in the max_num_work_groups e2e test #15359

Merged
merged 1 commit into from
Sep 13, 2024
Merged

[SYCL] Fix launch config in the max_num_work_groups e2e test #15359

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
14 changes: 6 additions & 8 deletions sycl/test-e2e/Basic/launch_queries/max_num_work_groups.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -71,7 +71,8 @@ int test_max_num_work_groups(sycl::queue &q, const sycl::device &dev) {

const size_t maxWorkGroupSize =
kernel.template get_info<work_group_size>(dev);
const size_t NumWorkItems = maxWorkGroupSize * maxWorkGroupSize;
// Will try to launch 2 work groups.
const size_t NumWorkItems = maxWorkGroupSize * 2;

size_t workGroupSize = 32;
size_t localMemorySizeInBytes{0};
Expand Down Expand Up @@ -100,8 +101,6 @@ int test_max_num_work_groups(sycl::queue &q, const sycl::device &dev) {
// ===================== //
// We must have at least one active group if we are below resource limits.
assert(maxWGs > 0 && "max_num_work_groups query failed");
if (maxWGs == 0)
return 1;

// Run the kernel
auto launch_range = sycl::nd_range<1>{sycl::range<1>{NumWorkItems},
Expand All @@ -121,7 +120,6 @@ int test_max_num_work_groups(sycl::queue &q, const sycl::device &dev) {
// ========================== //
// Test 3 - use max resources //
// ========================== //
// A little over the maximum work-group size for the purpose of exceeding.
workGroupSize = maxWorkGroupSize;
workGroupRange[0] = workGroupSize;
size_t localSize =
Expand All @@ -133,9 +131,8 @@ int test_max_num_work_groups(sycl::queue &q, const sycl::device &dev) {
syclex::info::kernel_queue_specific::max_num_work_groups>(
q, workGroupRange, localMemorySizeInBytes);

// We must have at least one active group if we are at resource limits.
assert(maxWGs > 0 && "max_num_work_groups query failed");
if (maxWGs == 0)
return 1;

launch_range = sycl::nd_range<1>{sycl::range<1>{NumWorkItems},
sycl::range<1>{workGroupSize}};
Expand All @@ -155,6 +152,7 @@ int test_max_num_work_groups(sycl::queue &q, const sycl::device &dev) {
// =============================== //
// Test 4 - exceed resource limits //
// =============================== //
// A little over the maximum work-group size for the purpose of exceeding.
workGroupSize = maxWorkGroupSize + 32;
workGroupRange[0] = workGroupSize;
maxWGs = kernel.template ext_oneapi_get_info<
Expand All @@ -163,10 +161,10 @@ int test_max_num_work_groups(sycl::queue &q, const sycl::device &dev) {
// It cannot be possible to launch a kernel successfully with a configuration
// that exceeds the available resources as in the above defined workGroupSize.
// workGroupSize is larger than maxWorkGroupSize, hence maxWGs must equal 0.
// Note: Level-Zero currently always returns a non-zero value. While other
// backends (i.e., OpenCL, HIP) always return 1 in their implementations.
if (dev.get_backend() == sycl::backend::ext_oneapi_cuda) {
assert(maxWGs == 0 && "max_num_work_groups query failed");
if (maxWGs > 0)
return 1;
}

// As we ensured that the 'max_num_work_groups' query correctly
Expand Down
Loading