From e2c7e901b8cdf2ec6ebed685224ad1e56b89a065 Mon Sep 17 00:00:00 2001 From: Pedro Olsen Ferreira Date: Tue, 25 Jun 2024 17:41:23 +0100 Subject: [PATCH] Add special test for generic address space atomics (#1959) The atomic operations are tested with generic pointers but in a way where the compiler can infer the original pointer address space. This commit adds tests specifically for the case where the compiler cannot make inferences. Test that the correct address is used with atomics when the address is group variant and invariant. --- .../generic_address_space/CMakeLists.txt | 1 + .../generic_address_space/atomic_tests.cpp | 224 ++++++++++++++++++ .../generic_address_space/main.cpp | 50 ++-- 3 files changed, 254 insertions(+), 21 deletions(-) create mode 100644 test_conformance/generic_address_space/atomic_tests.cpp diff --git a/test_conformance/generic_address_space/CMakeLists.txt b/test_conformance/generic_address_space/CMakeLists.txt index e74bcf4a5..951c5ab4e 100644 --- a/test_conformance/generic_address_space/CMakeLists.txt +++ b/test_conformance/generic_address_space/CMakeLists.txt @@ -2,6 +2,7 @@ set(MODULE_NAME GENERIC_ADDRESS_SPACE) set(${MODULE_NAME}_SOURCES advanced_tests.cpp + atomic_tests.cpp basic_tests.cpp main.cpp stress_tests.cpp diff --git a/test_conformance/generic_address_space/atomic_tests.cpp b/test_conformance/generic_address_space/atomic_tests.cpp new file mode 100644 index 000000000..a24c6ae2f --- /dev/null +++ b/test_conformance/generic_address_space/atomic_tests.cpp @@ -0,0 +1,224 @@ +// +// Copyright (c) 2024 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 + +#ifdef __APPLE__ +#include +#else +#include +#endif + +#include "errorHelpers.h" +#include "typeWrappers.h" + +namespace { +// In this source, each workgroup will generate one value. +// Every other workgroup will use either a global or local +// pointer on an atomic operation. +const char* KernelSourceInvariant = R"OpenCLC( +kernel void testKernel(global atomic_int* globalPtr, local atomic_int* localPtr) { + int gid = get_global_id(0); + int tid = get_local_id(0); + int wgid = get_group_id(0); + int wgsize = get_local_size(0); + + if (tid == 0) atomic_store(localPtr, 0); + + barrier(CLK_LOCAL_MEM_FENCE); + + // Initialise the generic pointer to + // the global. + generic atomic_int* ptr = globalPtr + wgid; + + // In a workgroup-invariant way, select a localPtr instead. + if ((wgid % 2) == 0) + ptr = localPtr; + + int inc = atomic_fetch_add(ptr, 1); + + // In the cases where the local memory ptr was used, + // save off the final value. + if ((wgid % 2) == 0 && inc == (wgsize-1)) + atomic_store(&globalPtr[wgid], inc); +} +)OpenCLC"; + +// In this source, each workgroup will generate two values. +// Every other work item in the workgroup will select either +// a local or global memory pointer and perform an atomic +// operation on that. +const char* KernelSourceVariant = R"OpenCLC( +kernel void testKernel(global atomic_int* globalPtr, local atomic_int* localPtr) { + int gid = get_global_id(0); + int tid = get_local_id(0); + int wgid = get_group_id(0); + int wgsize = get_local_size(0); + + if (tid == 0) atomic_store(localPtr, 0); + + barrier(CLK_LOCAL_MEM_FENCE); + + // Initialise the generic pointer to + // the global. Two values are written per WG. + generic atomic_int* ptr = globalPtr + (wgid * 2); + + // In a workgroup-invariant way, select a localPtr instead. + if ((tid % 2) == 0) + ptr = localPtr; + + atomic_fetch_add(ptr, 1); + + barrier(CLK_LOCAL_MEM_FENCE); + + // In the cases where the local memory ptr was used, + // save off the final value. + if (tid == 0) + atomic_store(&globalPtr[(wgid * 2) + 1], atomic_load(localPtr)); +} +)OpenCLC"; +} + +int test_generic_atomics_invariant(cl_device_id deviceID, cl_context context, + cl_command_queue queue, int) +{ + const auto version = get_device_cl_version(deviceID); + + if (version < Version(2, 0)) return TEST_SKIPPED_ITSELF; + + cl_int err = CL_SUCCESS; + clProgramWrapper program; + clKernelWrapper kernel; + + err = create_single_kernel_helper(context, &program, &kernel, 1, + &KernelSourceInvariant, "testKernel"); + test_error(err, "Failed to create test kernel"); + + size_t wgSize, retSize; + // Attempt to find the simd unit size for the device. + err = clGetKernelWorkGroupInfo(kernel, deviceID, + CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE, + sizeof(wgSize), &wgSize, &retSize); + test_error(err, "clGetKernelWorkGroupInfo failed"); + + // How many workgroups to run for the test. + const int numWGs = 2; + const size_t bufferSize = numWGs * sizeof(cl_uint); + clMemWrapper buffer = + clCreateBuffer(context, CL_MEM_WRITE_ONLY, bufferSize, nullptr, &err); + test_error(err, "clCreateBuffer failed"); + const cl_int zero = 0; + err = clEnqueueFillBuffer(queue, buffer, &zero, sizeof(zero), 0, bufferSize, + 0, nullptr, nullptr); + test_error(err, "clEnqueueFillBuffer failed"); + + err = clSetKernelArg(kernel, 0, sizeof(cl_mem), &buffer); + test_error(err, "clSetKernelArg failed"); + err = clSetKernelArg(kernel, 1, bufferSize, nullptr); + test_error(err, "clSetKernelArg failed"); + + const size_t globalSize = wgSize * numWGs; + err = clEnqueueNDRangeKernel(queue, kernel, 1, nullptr, &globalSize, + &wgSize, 0, nullptr, nullptr); + test_error(err, "clEnqueueNDRangeKernel failed"); + + std::vector results(numWGs); + err = clEnqueueReadBuffer(queue, buffer, CL_TRUE, 0, bufferSize, + results.data(), 0, nullptr, nullptr); + test_error(err, "clEnqueueReadBuffer failed"); + + clFinish(queue); + + for (size_t i = 0; i < numWGs; ++i) + { + const cl_int expected = ((i % 2) == 0) ? wgSize - 1 : wgSize; + if (results[i] != expected) + { + log_error("Verification on device failed at index %zu\n", i); + return TEST_FAIL; + } + } + + return CL_SUCCESS; +} + +int test_generic_atomics_variant(cl_device_id deviceID, cl_context context, + cl_command_queue queue, int) +{ + const auto version = get_device_cl_version(deviceID); + + if (version < Version(2, 0)) return TEST_SKIPPED_ITSELF; + + cl_int err = CL_SUCCESS; + clProgramWrapper program; + clKernelWrapper kernel; + + err = create_single_kernel_helper(context, &program, &kernel, 1, + &KernelSourceVariant, "testKernel"); + test_error(err, "Failed to create test kernel"); + + size_t wgSize, retSize; + // Attempt to find the simd unit size for the device. + err = clGetKernelWorkGroupInfo(kernel, deviceID, + CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE, + sizeof(wgSize), &wgSize, &retSize); + test_error(err, "clGetKernelWorkGroupInfo failed"); + + // How many workgroups to run for the test. + const int numWGs = 2; + const size_t bufferSize = numWGs * sizeof(cl_uint) * 2; + clMemWrapper buffer = + clCreateBuffer(context, CL_MEM_WRITE_ONLY, bufferSize, nullptr, &err); + test_error(err, "clCreateBuffer failed"); + const cl_int zero = 0; + err = clEnqueueFillBuffer(queue, buffer, &zero, sizeof(zero), 0, bufferSize, + 0, nullptr, nullptr); + test_error(err, "clEnqueueFillBuffer failed"); + + err = clSetKernelArg(kernel, 0, sizeof(cl_mem), &buffer); + test_error(err, "clSetKernelArg failed"); + err = clSetKernelArg(kernel, 1, bufferSize, nullptr); + test_error(err, "clSetKernelArg failed"); + + const size_t globalSize = wgSize * numWGs; + err = clEnqueueNDRangeKernel(queue, kernel, 1, nullptr, &globalSize, + &wgSize, 0, nullptr, nullptr); + test_error(err, "clEnqueueNDRangeKernel failed"); + + std::vector results(numWGs * 2); + err = clEnqueueReadBuffer(queue, buffer, CL_TRUE, 0, bufferSize, + results.data(), 0, nullptr, nullptr); + test_error(err, "clEnqueueReadBuffer failed"); + + clFinish(queue); + + const cl_int expected = wgSize / 2; + for (size_t i = 0; i < (numWGs * 2); i += 2) + { + if (results[i] != expected) + { + log_error("Verification on device failed at index %zu\n", i); + return TEST_FAIL; + } + if (results[i + 1] != expected) + { + const size_t index = i + 1; + log_error("Verification on device failed at index %zu\n", index); + return TEST_FAIL; + } + } + + return CL_SUCCESS; +} diff --git a/test_conformance/generic_address_space/main.cpp b/test_conformance/generic_address_space/main.cpp index 011475839..a7897367b 100644 --- a/test_conformance/generic_address_space/main.cpp +++ b/test_conformance/generic_address_space/main.cpp @@ -40,31 +40,39 @@ extern int test_generic_advanced_casting(cl_device_id deviceID, cl_context conte extern int test_generic_ptr_to_host_mem(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); extern int test_generic_ptr_to_host_mem_svm(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); extern int test_max_number_of_params(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); +// atomic tests +int test_generic_atomics_invariant(cl_device_id deviceID, cl_context context, + cl_command_queue queue, int num_elements); +int test_generic_atomics_variant(cl_device_id deviceID, cl_context context, + cl_command_queue queue, int num_elements); test_definition test_list[] = { // basic tests - ADD_TEST( function_get_fence ), - ADD_TEST( function_to_address_space ), - ADD_TEST( variable_get_fence ), - ADD_TEST( variable_to_address_space ), - ADD_TEST( casting ), - ADD_TEST( conditional_casting ), - ADD_TEST( chain_casting ), - ADD_TEST( ternary_operator_casting ), - ADD_TEST( language_struct ), - ADD_TEST( language_union ), - ADD_TEST( multiple_calls_same_function ), - ADD_TEST( compare_pointers ), + ADD_TEST(function_get_fence), + ADD_TEST(function_to_address_space), + ADD_TEST(variable_get_fence), + ADD_TEST(variable_to_address_space), + ADD_TEST(casting), + ADD_TEST(conditional_casting), + ADD_TEST(chain_casting), + ADD_TEST(ternary_operator_casting), + ADD_TEST(language_struct), + ADD_TEST(language_union), + ADD_TEST(multiple_calls_same_function), + ADD_TEST(compare_pointers), // advanced tests - ADD_TEST( library_function ), - ADD_TEST( generic_variable_volatile ), - ADD_TEST( generic_variable_const ), - ADD_TEST( generic_variable_gentype ), - ADD_TEST( builtin_functions ), - ADD_TEST( generic_advanced_casting ), - ADD_TEST( generic_ptr_to_host_mem ), - ADD_TEST( generic_ptr_to_host_mem_svm ), - ADD_TEST( max_number_of_params ), + ADD_TEST(library_function), + ADD_TEST(generic_variable_volatile), + ADD_TEST(generic_variable_const), + ADD_TEST(generic_variable_gentype), + ADD_TEST(builtin_functions), + ADD_TEST(generic_advanced_casting), + ADD_TEST(generic_ptr_to_host_mem), + ADD_TEST(generic_ptr_to_host_mem_svm), + ADD_TEST(max_number_of_params), + // atomic tests + ADD_TEST(generic_atomics_invariant), + ADD_TEST(generic_atomics_variant), }; const int test_num = ARRAY_SIZE( test_list );