Skip to content

Commit

Permalink
fixes for cl_khr_command_buffer_mutable_dispatch tests
Browse files Browse the repository at this point in the history
mutable_dispatch_global_size was failing because the update_global_size = 3
was not a multiple of the Local WS (original command's LWS = GWS / 64 and PoCL
doesn't support non-uniform WGs). Fixed by increasing the GWS to 256K and LWS
to 16K. The test should really detect the max Local WS of the device and use
multiples of that value - that ensures changing global WS without changing
local WS will still result in an acceptable GWS/LWS combination.

Fixes another test which had an out-of-bounds access because it hardcoded
old value of GWS.
  • Loading branch information
franz committed Oct 28, 2024
1 parent ac736ac commit 83613f9
Show file tree
Hide file tree
Showing 3 changed files with 7 additions and 5 deletions.
Original file line number Diff line number Diff line change
Expand Up @@ -126,7 +126,7 @@ struct BasicMutableCommandBufferTest : BasicCommandBufferTest
clUpdateMutableCommandsKHR_fn clUpdateMutableCommandsKHR = nullptr;

const char* kernelString = "__kernel void empty() {}";
const size_t global_work_size = 4 * 16;
const size_t global_work_size = 256 * 1024;
};

struct InfoMutableCommandBufferTest : BasicMutableCommandBufferTest
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -153,9 +153,9 @@ struct MutableDispatchGlobalSize : public InfoMutableCommandBufferTest
}

size_t info_global_size = 0;
const size_t update_global_size = 3;
const size_t sizeToAllocate = global_work_size;
const size_t num_elements = sizeToAllocate / sizeof(cl_int);
const size_t update_global_size = 16 * 1024;
const size_t sizeToAllocate = global_work_size * 4;
const size_t num_elements = global_work_size;
cl_mutable_command_khr command = nullptr;
};

Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -268,7 +268,9 @@ struct MutableDispatchWorkGroups : public BasicMutableCommandBufferTest
static constexpr size_t test_global_work_size = 64;
static constexpr size_t update_global_size = 16;
const size_t local_work_size = 8;
const size_t sizeToAllocate = 64 * sizeof(cl_int);
const size_t sizeToAllocate = (test_global_work_size > global_work_size
? test_global_work_size
: global_work_size) * sizeof(cl_int);
};

int test_command_buffer_with_no_additional_work_groups(cl_device_id device,
Expand Down

0 comments on commit 83613f9

Please sign in to comment.