| // |
| // Copyright (c) 2017 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/typeWrappers.h" |
| #include "harness/mt19937.h" |
| #include "base.h" |
| |
| #include <string> |
| #include <vector> |
| #include <algorithm> |
| #include <sstream> |
| |
| class CStressTest : public CTest { |
| public: |
| CStressTest(const std::vector<std::string>& kernel) : CTest(), _kernels(kernel) { |
| |
| } |
| |
| CStressTest(const std::string& kernel) : CTest(), _kernels(1, kernel) { |
| |
| } |
| |
| int ExecuteSubcase(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements, const std::string& src) { |
| cl_int error; |
| |
| clProgramWrapper program; |
| clKernelWrapper kernel; |
| |
| const char *srcPtr = src.c_str(); |
| |
| if (create_single_kernel_helper_with_build_options(context, &program, &kernel, 1, &srcPtr, "testKernel", "-cl-std=CL2.0")) { |
| log_error("create_single_kernel_helper failed"); |
| return -1; |
| } |
| |
| size_t bufferSize = num_elements * sizeof(cl_uint); |
| clMemWrapper buffer = clCreateBuffer(context, CL_MEM_WRITE_ONLY, bufferSize, NULL, &error); |
| test_error(error, "clCreateBuffer failed"); |
| |
| error = clSetKernelArg(kernel, 0, sizeof(buffer), &buffer); |
| test_error(error, "clSetKernelArg failed"); |
| |
| size_t globalWorkGroupSize = num_elements; |
| size_t localWorkGroupSize = 0; |
| error = get_max_common_work_group_size(context, kernel, globalWorkGroupSize, &localWorkGroupSize); |
| test_error(error, "Unable to get common work group size"); |
| |
| error = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &globalWorkGroupSize, &localWorkGroupSize, 0, NULL, NULL); |
| test_error(error, "clEnqueueNDRangeKernel failed"); |
| |
| // verify results |
| std::vector<cl_uint> results(num_elements); |
| |
| error = clEnqueueReadBuffer(queue, buffer, CL_TRUE, 0, bufferSize, &results[0], 0, NULL, NULL); |
| test_error(error, "clEnqueueReadBuffer failed"); |
| |
| size_t passCount = std::count(results.begin(), results.end(), 1); |
| if (passCount != results.size()) { |
| std::vector<cl_uint>::iterator iter = std::find(results.begin(), results.end(), 0); |
| log_error("Verification on device failed at index %ld\n", std::distance(results.begin(), iter)); |
| log_error("%ld out of %ld failed\n", (results.size()-passCount), results.size()); |
| return -1; |
| } |
| |
| return CL_SUCCESS; |
| } |
| |
| int Execute(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) { |
| cl_int result = CL_SUCCESS; |
| |
| for (std::vector<std::string>::const_iterator it = _kernels.begin(); it != _kernels.end(); ++it) { |
| log_info("Executing subcase #%ld out of %ld\n", (it - _kernels.begin() + 1), _kernels.size()); |
| |
| result |= ExecuteSubcase(deviceID, context, queue, num_elements, *it); |
| } |
| |
| return result; |
| } |
| |
| private: |
| const std::vector<std::string> _kernels; |
| }; |
| |
| int test_max_number_of_params(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) { |
| cl_int error; |
| |
| size_t deviceMaxParameterSize; |
| error = clGetDeviceInfo(deviceID, CL_DEVICE_MAX_PARAMETER_SIZE, sizeof(deviceMaxParameterSize), &deviceMaxParameterSize, NULL); |
| test_error(error, "clGetDeviceInfo failed"); |
| |
| size_t deviceAddressBits; |
| error = clGetDeviceInfo(deviceID, CL_DEVICE_ADDRESS_BITS, sizeof(deviceAddressBits), &deviceAddressBits, NULL); |
| test_error(error, "clGetDeviceInfo failed"); |
| |
| size_t maxParams = deviceMaxParameterSize / (deviceAddressBits / 8); |
| |
| const std::string KERNEL_FUNCTION_TEMPLATE[] = { |
| common::CONFORMANCE_VERIFY_FENCE + |
| NL |
| NL "bool helperFunction(int *ptr0 ", |
| // the rest of arguments goes here |
| ") {" |
| NL " // check first pointer only" |
| NL " if (!isFenceValid(get_fence(ptr0)))" |
| NL " return false;" |
| NL |
| NL " return true;" |
| NL "}" |
| NL |
| NL "__kernel void testKernel(__global uint *results) {" |
| NL " uint tid = get_global_id(0);" |
| NL |
| NL " __global int * gptr;" |
| NL " __local int * lptr;" |
| NL " __private int * pptr;" |
| NL |
| NL " size_t failures = 0;" |
| NL |
| NL, |
| // the function body goes here |
| NL |
| NL " results[tid] = (failures == 0);" |
| NL "}" |
| NL |
| }; |
| |
| std::ostringstream type_params; |
| std::ostringstream function_calls; |
| |
| for (size_t i = 0; i < maxParams; i++) { |
| type_params << ", int *ptr" << i+1; |
| } |
| |
| // use pseudo random generator to shuffle params |
| MTdata d = init_genrand(gRandomSeed); |
| if (!d) |
| return -1; |
| |
| std::string pointers[] = { "gptr", "lptr", "pptr" }; |
| |
| size_t totalCalls = maxParams / 2; |
| for (size_t i = 0; i < totalCalls; i++) { |
| function_calls << "\tif (!helperFunction(gptr"; |
| |
| for (size_t j = 0; j < maxParams; j++) { |
| function_calls << ", " << pointers[genrand_int32(d)%3]; |
| } |
| |
| function_calls << ")) failures++;" << NL; |
| } |
| |
| free_mtdata(d); |
| d = NULL; |
| |
| const std::string KERNEL_FUNCTION = KERNEL_FUNCTION_TEMPLATE[0] + type_params.str() + KERNEL_FUNCTION_TEMPLATE[1] + function_calls.str() + KERNEL_FUNCTION_TEMPLATE[2]; |
| |
| CStressTest test(KERNEL_FUNCTION); |
| |
| return test.Execute(deviceID, context, queue, num_elements); |
| } |