From bb604702a5c34e71697f306b693e812b288828c6 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?K=C3=A9vin=20Petit?= Date: Tue, 11 Mar 2025 22:00:38 +0000 Subject: [PATCH] Migrate workgroups suite to the new test registration framework (#2308) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Contributes to #2181 Signed-off-by: Kévin Petit --- test_conformance/workgroups/main.cpp | 24 +- test_conformance/workgroups/procs.h | 74 ------- test_conformance/workgroups/testBase.h | 11 +- test_conformance/workgroups/test_wg_all.cpp | 15 +- test_conformance/workgroups/test_wg_any.cpp | 15 +- .../workgroups/test_wg_broadcast.cpp | 50 ++--- .../workgroups/test_wg_scan_reduce.cpp | 205 +++++++++--------- 7 files changed, 133 insertions(+), 261 deletions(-) delete mode 100644 test_conformance/workgroups/procs.h diff --git a/test_conformance/workgroups/main.cpp b/test_conformance/workgroups/main.cpp index 11849e1f..bdb489c4 100644 --- a/test_conformance/workgroups/main.cpp +++ b/test_conformance/workgroups/main.cpp @@ -16,32 +16,12 @@ #include "harness/compat.h" #include "harness/testHarness.h" -#include "procs.h" #include #include #if !defined(_WIN32) #include #endif -test_definition test_list[] = { - ADD_TEST_VERSION(work_group_all, Version(2, 0)), - ADD_TEST_VERSION(work_group_any, Version(2, 0)), - ADD_TEST_VERSION(work_group_reduce_add, Version(2, 0)), - ADD_TEST_VERSION(work_group_reduce_min, Version(2, 0)), - ADD_TEST_VERSION(work_group_reduce_max, Version(2, 0)), - ADD_TEST_VERSION(work_group_scan_inclusive_add, Version(2, 0)), - ADD_TEST_VERSION(work_group_scan_inclusive_min, Version(2, 0)), - ADD_TEST_VERSION(work_group_scan_inclusive_max, Version(2, 0)), - ADD_TEST_VERSION(work_group_scan_exclusive_add, Version(2, 0)), - ADD_TEST_VERSION(work_group_scan_exclusive_min, Version(2, 0)), - ADD_TEST_VERSION(work_group_scan_exclusive_max, Version(2, 0)), - ADD_TEST_VERSION(work_group_broadcast_1D, Version(2, 0)), - ADD_TEST_VERSION(work_group_broadcast_2D, Version(2, 0)), - ADD_TEST_VERSION(work_group_broadcast_3D, Version(2, 0)), -}; - -const int test_num = ARRAY_SIZE(test_list); - test_status InitCL(cl_device_id device) { auto version = get_device_cl_version(device); auto expected_min_version = Version(1, 2); @@ -77,6 +57,8 @@ test_status InitCL(cl_device_id device) { } int main(int argc, const char *argv[]) { - return runTestHarnessWithCheck(argc, argv, test_num, test_list, false, 0, InitCL); + return runTestHarnessWithCheck( + argc, argv, test_registry::getInstance().num_tests(), + test_registry::getInstance().definitions(), false, 0, InitCL); } diff --git a/test_conformance/workgroups/procs.h b/test_conformance/workgroups/procs.h deleted file mode 100644 index 0baa9066..00000000 --- a/test_conformance/workgroups/procs.h +++ /dev/null @@ -1,74 +0,0 @@ -// -// Copyright (c) 2017, 2021 The Khronos Group Inc. -// -// Licensed under the Apache License, Version 2.0 (the "License"); -// you may not use this file except in compliance with the License. -// You may obtain a copy of the License at -// -// http://www.apache.org/licenses/LICENSE-2.0 -// -// Unless required by applicable law or agreed to in writing, software -// distributed under the License is distributed on an "AS IS" BASIS, -// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -// See the License for the specific language governing permissions and -// limitations under the License. -// -#include "harness/testHarness.h" -#include "harness/kernelHelpers.h" -#include "harness/errorHelpers.h" -#include "harness/typeWrappers.h" -#include "harness/conversions.h" -#include "harness/mt19937.h" - -extern int create_program_and_kernel(const char *source, - const char *kernel_name, - cl_program *program_ret, - cl_kernel *kernel_ret); - -extern int test_work_group_all(cl_device_id deviceID, cl_context context, - cl_command_queue queue, int num_elements); -extern int test_work_group_any(cl_device_id deviceID, cl_context context, - cl_command_queue queue, int num_elements); -extern int test_work_group_broadcast_1D(cl_device_id deviceID, - cl_context context, - cl_command_queue queue, - int num_elements); -extern int test_work_group_broadcast_2D(cl_device_id deviceID, - cl_context context, - cl_command_queue queue, - int num_elements); -extern int test_work_group_broadcast_3D(cl_device_id deviceID, - cl_context context, - cl_command_queue queue, - int num_elements); -extern int test_work_group_reduce_add(cl_device_id deviceID, cl_context context, - cl_command_queue queue, int num_elements); -extern int test_work_group_reduce_min(cl_device_id deviceID, cl_context context, - cl_command_queue queue, int num_elements); -extern int test_work_group_reduce_max(cl_device_id deviceID, cl_context context, - cl_command_queue queue, int num_elements); - -extern int test_work_group_scan_exclusive_add(cl_device_id deviceID, - cl_context context, - cl_command_queue queue, - int num_elements); -extern int test_work_group_scan_exclusive_min(cl_device_id deviceID, - cl_context context, - cl_command_queue queue, - int num_elements); -extern int test_work_group_scan_exclusive_max(cl_device_id deviceID, - cl_context context, - cl_command_queue queue, - int num_elements); -extern int test_work_group_scan_inclusive_add(cl_device_id deviceID, - cl_context context, - cl_command_queue queue, - int num_elements); -extern int test_work_group_scan_inclusive_min(cl_device_id deviceID, - cl_context context, - cl_command_queue queue, - int num_elements); -extern int test_work_group_scan_inclusive_max(cl_device_id deviceID, - cl_context context, - cl_command_queue queue, - int num_elements); diff --git a/test_conformance/workgroups/testBase.h b/test_conformance/workgroups/testBase.h index 5b49bfd7..0251ad25 100644 --- a/test_conformance/workgroups/testBase.h +++ b/test_conformance/workgroups/testBase.h @@ -16,6 +16,12 @@ #ifndef _testBase_h #define _testBase_h +#include "harness/testHarness.h" +#include "harness/kernelHelpers.h" +#include "harness/errorHelpers.h" +#include "harness/typeWrappers.h" +#include "harness/conversions.h" +#include "harness/mt19937.h" #include "harness/compat.h" #include @@ -23,9 +29,4 @@ #include #include -#include "procs.h" - #endif // _testBase_h - - - diff --git a/test_conformance/workgroups/test_wg_all.cpp b/test_conformance/workgroups/test_wg_all.cpp index 7269f5cd..1eefffa8 100644 --- a/test_conformance/workgroups/test_wg_all.cpp +++ b/test_conformance/workgroups/test_wg_all.cpp @@ -20,8 +20,7 @@ #include #include -#include "procs.h" - +#include "testBase.h" const char *wg_all_kernel_code = "__kernel void test_wg_all(global float *input, global int *output)\n" @@ -65,8 +64,7 @@ verify_wg_all(float *inptr, int *outptr, size_t n, size_t wg_size) return 0; } -int -test_work_group_all(cl_device_id device, cl_context context, cl_command_queue queue, int n_elems) +REGISTER_TEST_VERSION(work_group_all, Version(2, 0)) { cl_mem streams[2]; cl_float *input_ptr[1], *p; @@ -74,8 +72,7 @@ test_work_group_all(cl_device_id device, cl_context context, cl_command_queue qu cl_program program; cl_kernel kernel; size_t threads[1]; - size_t wg_size[1]; - size_t num_elements; + size_t wg_size[1]; int err; MTdata d; @@ -88,8 +85,6 @@ test_work_group_all(cl_device_id device, cl_context context, cl_command_queue qu err = get_max_allowed_1d_work_group_size_on_device(device, kernel, wg_size); test_error(err, "get_max_allowed_1d_work_group_size_on_device failed"); - num_elements = n_elems; - input_ptr[0] = (cl_float*)malloc(sizeof(cl_float) * (num_elements+1)); output_ptr = (cl_int*)malloc(sizeof(cl_int) * (num_elements+1)); streams[0] = @@ -111,7 +106,7 @@ test_work_group_all(cl_device_id device, cl_context context, cl_command_queue qu p = input_ptr[0]; d = init_genrand( gRandomSeed ); - for (size_t i = 0; i < (num_elements + 1); i++) + for (int i = 0; i < (num_elements + 1); i++) { p[i] = get_random_float((float)(-100000.f * M_PI), (float)(100000.f * M_PI) ,d); } @@ -133,7 +128,7 @@ test_work_group_all(cl_device_id device, cl_context context, cl_command_queue qu } // Line below is troublesome... - threads[0] = (size_t)n_elems; + threads[0] = (size_t)num_elements; err = clEnqueueNDRangeKernel( queue, kernel, 1, NULL, threads, wg_size, 0, NULL, NULL ); if (err != CL_SUCCESS) { diff --git a/test_conformance/workgroups/test_wg_any.cpp b/test_conformance/workgroups/test_wg_any.cpp index 66639302..2b2e1b9b 100644 --- a/test_conformance/workgroups/test_wg_any.cpp +++ b/test_conformance/workgroups/test_wg_any.cpp @@ -20,8 +20,7 @@ #include #include -#include "procs.h" - +#include "testBase.h" const char *wg_any_kernel_code = "__kernel void test_wg_any(global float *input, global int *output)\n" @@ -65,8 +64,7 @@ verify_wg_any(float *inptr, int *outptr, size_t n, size_t wg_size) return 0; } -int -test_work_group_any(cl_device_id device, cl_context context, cl_command_queue queue, int n_elems) +REGISTER_TEST_VERSION(work_group_any, Version(2, 0)) { cl_mem streams[2]; cl_float *input_ptr[1], *p; @@ -74,8 +72,7 @@ test_work_group_any(cl_device_id device, cl_context context, cl_command_queue qu cl_program program; cl_kernel kernel; size_t threads[1]; - size_t wg_size[1]; - size_t num_elements; + size_t wg_size[1]; int err; MTdata d; @@ -88,8 +85,6 @@ test_work_group_any(cl_device_id device, cl_context context, cl_command_queue qu err = get_max_allowed_1d_work_group_size_on_device(device, kernel, wg_size); test_error(err, "get_max_allowed_1d_work_group_size_on_device failed"); - num_elements = n_elems; - input_ptr[0] = (cl_float*)malloc(sizeof(cl_float) * (num_elements+1)); output_ptr = (cl_int*)malloc(sizeof(cl_int) * (num_elements+1)); streams[0] = @@ -111,7 +106,7 @@ test_work_group_any(cl_device_id device, cl_context context, cl_command_queue qu p = input_ptr[0]; d = init_genrand( gRandomSeed ); - for (size_t i = 0; i < (num_elements + 1); i++) + for (int i = 0; i < (num_elements + 1); i++) { p[i] = get_random_float((float)(-100000.f * M_PI), (float)(100000.f * M_PI) ,d); } @@ -133,7 +128,7 @@ test_work_group_any(cl_device_id device, cl_context context, cl_command_queue qu } // Line below is troublesome... - threads[0] = (size_t)n_elems; + threads[0] = (size_t)num_elements; err = clEnqueueNDRangeKernel( queue, kernel, 1, NULL, threads, wg_size, 0, NULL, NULL ); if (err != CL_SUCCESS) { diff --git a/test_conformance/workgroups/test_wg_broadcast.cpp b/test_conformance/workgroups/test_wg_broadcast.cpp index 05d30ed8..3c8beacb 100644 --- a/test_conformance/workgroups/test_wg_broadcast.cpp +++ b/test_conformance/workgroups/test_wg_broadcast.cpp @@ -22,8 +22,7 @@ #include -#include "procs.h" - +#include "testBase.h" const char *wg_broadcast_1D_kernel_code = "__kernel void test_wg_broadcast_1D(global float *input, global float *output)\n" @@ -168,8 +167,7 @@ verify_wg_broadcast_3D(float *inptr, float *outptr, size_t nx, size_t ny, size_t } -int -test_work_group_broadcast_1D(cl_device_id device, cl_context context, cl_command_queue queue, int n_elems) +REGISTER_TEST_VERSION(work_group_broadcast_1D, Version(2, 0)) { cl_mem streams[2]; cl_float *input_ptr[1], *p; @@ -177,8 +175,7 @@ test_work_group_broadcast_1D(cl_device_id device, cl_context context, cl_command cl_program program; cl_kernel kernel; size_t globalsize[1]; - size_t wg_size[1]; - size_t num_elements; + size_t wg_size[1]; int err; MTdata d; @@ -192,8 +189,6 @@ test_work_group_broadcast_1D(cl_device_id device, cl_context context, cl_command err = get_max_allowed_1d_work_group_size_on_device(device, kernel, wg_size); test_error(err, "get_max_allowed_1d_work_group_size_on_device failed"); - num_elements = n_elems; - input_ptr[0] = (cl_float*)malloc(sizeof(cl_float) * num_elements); output_ptr = (cl_float*)malloc(sizeof(cl_float) * num_elements); streams[0] = clCreateBuffer(context, CL_MEM_READ_WRITE, @@ -214,7 +209,7 @@ test_work_group_broadcast_1D(cl_device_id device, cl_context context, cl_command p = input_ptr[0]; d = init_genrand( gRandomSeed ); - for (size_t i = 0; i < num_elements; i++) + for (int i = 0; i < num_elements; i++) { p[i] = get_random_float((float)(-100000.f * M_PI), (float)(100000.f * M_PI) ,d); } @@ -236,7 +231,7 @@ test_work_group_broadcast_1D(cl_device_id device, cl_context context, cl_command } // Line below is troublesome... - globalsize[0] = (size_t)n_elems; + globalsize[0] = (size_t)num_elements; err = clEnqueueNDRangeKernel( queue, kernel, 1, NULL, globalsize, wg_size, 0, NULL, NULL ); if (err != CL_SUCCESS) { @@ -271,8 +266,7 @@ test_work_group_broadcast_1D(cl_device_id device, cl_context context, cl_command } -int -test_work_group_broadcast_2D(cl_device_id device, cl_context context, cl_command_queue queue, int n_elems) +REGISTER_TEST_VERSION(work_group_broadcast_2D, Version(2, 0)) { cl_mem streams[2]; cl_float *input_ptr[1], *p; @@ -282,8 +276,7 @@ test_work_group_broadcast_2D(cl_device_id device, cl_context context, cl_command size_t globalsize[2]; size_t localsize[2]; size_t wg_size[1]; - size_t num_workgroups; - size_t num_elements; + size_t num_workgroups; int err; MTdata d; @@ -314,7 +307,7 @@ test_work_group_broadcast_2D(cl_device_id device, cl_context context, cl_command localsize[0] = localsize[1] = 1; } - num_workgroups = std::max(n_elems / wg_size[0], (size_t)16); + num_workgroups = std::max(num_elements / wg_size[0], (size_t)16); globalsize[0] = num_workgroups * localsize[0]; globalsize[1] = num_workgroups * localsize[1]; num_elements = globalsize[0] * globalsize[1]; @@ -339,7 +332,7 @@ test_work_group_broadcast_2D(cl_device_id device, cl_context context, cl_command p = input_ptr[0]; d = init_genrand( gRandomSeed ); - for (size_t i = 0; i < num_elements; i++) + for (int i = 0; i < num_elements; i++) { p[i] = get_random_float((float)(-100000.f * M_PI), (float)(100000.f * M_PI) ,d); } @@ -394,8 +387,7 @@ test_work_group_broadcast_2D(cl_device_id device, cl_context context, cl_command } -int -test_work_group_broadcast_3D(cl_device_id device, cl_context context, cl_command_queue queue, int n_elems) +REGISTER_TEST_VERSION(work_group_broadcast_3D, Version(2, 0)) { cl_mem streams[2]; cl_float *input_ptr[1], *p; @@ -405,8 +397,7 @@ test_work_group_broadcast_3D(cl_device_id device, cl_context context, cl_command size_t globalsize[3]; size_t localsize[3]; size_t wg_size[1]; - size_t num_workgroups; - size_t num_elements; + size_t num_workgroups; int err; MTdata d; @@ -437,7 +428,7 @@ test_work_group_broadcast_3D(cl_device_id device, cl_context context, cl_command localsize[0] = localsize[1] = localsize[2] = 1; } - num_workgroups = std::max(n_elems / wg_size[0], (size_t)8); + num_workgroups = std::max(num_elements / wg_size[0], (size_t)8); globalsize[0] = num_workgroups * localsize[0]; globalsize[1] = num_workgroups * localsize[1]; globalsize[2] = num_workgroups * localsize[2]; @@ -463,7 +454,7 @@ test_work_group_broadcast_3D(cl_device_id device, cl_context context, cl_command p = input_ptr[0]; d = init_genrand( gRandomSeed ); - for (size_t i = 0; i < num_elements; i++) + for (int i = 0; i < num_elements; i++) { p[i] = get_random_float((float)(-100000.f * M_PI), (float)(100000.f * M_PI) ,d); } @@ -516,18 +507,3 @@ test_work_group_broadcast_3D(cl_device_id device, cl_context context, cl_command return err; } - - -int -test_work_group_broadcast(cl_device_id device, cl_context context, cl_command_queue queue, int n_elems) -{ - int err; - - err = test_work_group_broadcast_1D(device, context, queue, n_elems); - if (err) return err; - err = test_work_group_broadcast_2D(device, context, queue, n_elems); - if (err) return err; - return err; -} - - diff --git a/test_conformance/workgroups/test_wg_scan_reduce.cpp b/test_conformance/workgroups/test_wg_scan_reduce.cpp index bf4dc89e..f1f28cee 100644 --- a/test_conformance/workgroups/test_wg_scan_reduce.cpp +++ b/test_conformance/workgroups/test_wg_scan_reduce.cpp @@ -19,7 +19,7 @@ #include #include -#include "procs.h" +#include "testBase.h" static std::string make_kernel_string(const std::string &type, const std::string &kernelName, @@ -272,184 +272,181 @@ static int run_test(cl_device_id device, cl_context context, return TEST_PASS; } -int test_work_group_reduce_add(cl_device_id device, cl_context context, - cl_command_queue queue, int n_elems) -{ - int result = TEST_PASS; - - result |= run_test>>(device, context, queue, n_elems); - result |= run_test>>(device, context, queue, n_elems); - - if (gHasLong) - { - result |= - run_test>>(device, context, queue, n_elems); - result |= - run_test>>(device, context, queue, n_elems); - } - - return result; -} - -int test_work_group_reduce_max(cl_device_id device, cl_context context, - cl_command_queue queue, int n_elems) -{ - int result = TEST_PASS; - - result |= run_test>>(device, context, queue, n_elems); - result |= run_test>>(device, context, queue, n_elems); - - if (gHasLong) - { - result |= - run_test>>(device, context, queue, n_elems); - result |= - run_test>>(device, context, queue, n_elems); - } - - return result; -} - -int test_work_group_reduce_min(cl_device_id device, cl_context context, - cl_command_queue queue, int n_elems) -{ - int result = TEST_PASS; - - result |= run_test>>(device, context, queue, n_elems); - result |= run_test>>(device, context, queue, n_elems); - - if (gHasLong) - { - result |= - run_test>>(device, context, queue, n_elems); - result |= - run_test>>(device, context, queue, n_elems); - } - - return result; -} - -int test_work_group_scan_inclusive_add(cl_device_id device, cl_context context, - cl_command_queue queue, int n_elems) +REGISTER_TEST_VERSION(work_group_reduce_add, Version(2, 0)) { int result = TEST_PASS; result |= - run_test>>(device, context, queue, n_elems); + run_test>>(device, context, queue, num_elements); result |= - run_test>>(device, context, queue, n_elems); + run_test>>(device, context, queue, num_elements); + + if (gHasLong) + { + result |= run_test>>(device, context, queue, + num_elements); + result |= run_test>>(device, context, queue, + num_elements); + } + + return result; +} + +REGISTER_TEST_VERSION(work_group_reduce_max, Version(2, 0)) +{ + int result = TEST_PASS; + + result |= + run_test>>(device, context, queue, num_elements); + result |= + run_test>>(device, context, queue, num_elements); + + if (gHasLong) + { + result |= run_test>>(device, context, queue, + num_elements); + result |= run_test>>(device, context, queue, + num_elements); + } + + return result; +} + +REGISTER_TEST_VERSION(work_group_reduce_min, Version(2, 0)) +{ + int result = TEST_PASS; + + result |= + run_test>>(device, context, queue, num_elements); + result |= + run_test>>(device, context, queue, num_elements); + + if (gHasLong) + { + result |= run_test>>(device, context, queue, + num_elements); + result |= run_test>>(device, context, queue, + num_elements); + } + + return result; +} + +REGISTER_TEST_VERSION(work_group_scan_inclusive_add, Version(2, 0)) +{ + int result = TEST_PASS; + + result |= run_test>>(device, context, queue, + num_elements); + result |= run_test>>(device, context, queue, + num_elements); if (gHasLong) { result |= run_test>>(device, context, queue, - n_elems); + num_elements); result |= run_test>>(device, context, queue, - n_elems); + num_elements); } return result; } -int test_work_group_scan_inclusive_max(cl_device_id device, cl_context context, - cl_command_queue queue, int n_elems) +REGISTER_TEST_VERSION(work_group_scan_inclusive_max, Version(2, 0)) { int result = TEST_PASS; - result |= - run_test>>(device, context, queue, n_elems); - result |= - run_test>>(device, context, queue, n_elems); + result |= run_test>>(device, context, queue, + num_elements); + result |= run_test>>(device, context, queue, + num_elements); if (gHasLong) { result |= run_test>>(device, context, queue, - n_elems); + num_elements); result |= run_test>>(device, context, queue, - n_elems); + num_elements); } return result; } -int test_work_group_scan_inclusive_min(cl_device_id device, cl_context context, - cl_command_queue queue, int n_elems) +REGISTER_TEST_VERSION(work_group_scan_inclusive_min, Version(2, 0)) { int result = TEST_PASS; - result |= - run_test>>(device, context, queue, n_elems); - result |= - run_test>>(device, context, queue, n_elems); + result |= run_test>>(device, context, queue, + num_elements); + result |= run_test>>(device, context, queue, + num_elements); if (gHasLong) { result |= run_test>>(device, context, queue, - n_elems); + num_elements); result |= run_test>>(device, context, queue, - n_elems); + num_elements); } return result; } -int test_work_group_scan_exclusive_add(cl_device_id device, cl_context context, - cl_command_queue queue, int n_elems) +REGISTER_TEST_VERSION(work_group_scan_exclusive_add, Version(2, 0)) { int result = TEST_PASS; - result |= - run_test>>(device, context, queue, n_elems); - result |= - run_test>>(device, context, queue, n_elems); + result |= run_test>>(device, context, queue, + num_elements); + result |= run_test>>(device, context, queue, + num_elements); if (gHasLong) { result |= run_test>>(device, context, queue, - n_elems); + num_elements); result |= run_test>>(device, context, queue, - n_elems); + num_elements); } return result; } -int test_work_group_scan_exclusive_max(cl_device_id device, cl_context context, - cl_command_queue queue, int n_elems) +REGISTER_TEST_VERSION(work_group_scan_exclusive_max, Version(2, 0)) { int result = TEST_PASS; - result |= - run_test>>(device, context, queue, n_elems); - result |= - run_test>>(device, context, queue, n_elems); + result |= run_test>>(device, context, queue, + num_elements); + result |= run_test>>(device, context, queue, + num_elements); if (gHasLong) { result |= run_test>>(device, context, queue, - n_elems); + num_elements); result |= run_test>>(device, context, queue, - n_elems); + num_elements); } return result; } -int test_work_group_scan_exclusive_min(cl_device_id device, cl_context context, - cl_command_queue queue, int n_elems) +REGISTER_TEST_VERSION(work_group_scan_exclusive_min, Version(2, 0)) { int result = TEST_PASS; - result |= - run_test>>(device, context, queue, n_elems); - result |= - run_test>>(device, context, queue, n_elems); + result |= run_test>>(device, context, queue, + num_elements); + result |= run_test>>(device, context, queue, + num_elements); if (gHasLong) { result |= run_test>>(device, context, queue, - n_elems); + num_elements); result |= run_test>>(device, context, queue, - n_elems); + num_elements); } return result;