| // |
| // 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 "testBase.h" |
| #include "harness/typeWrappers.h" |
| #include "harness/conversions.h" |
| |
| const char* zero_sized_enqueue_test_kernel[] = { |
| "__kernel void foo_kernel(__global int *dst)\n" |
| "{\n" |
| " int tid = get_global_id(0);\n" |
| "\n" |
| " dst[tid] = 1;\n" |
| "\n" |
| "}\n" |
| }; |
| |
| const int bufSize = 128; |
| |
| cl_int test_zero_sized_enqueue_and_test_output_buffer(cl_command_queue queue, clKernelWrapper& kernel, clMemWrapper& buf, size_t dim, size_t ndrange[]) |
| { |
| cl_int error = clEnqueueNDRangeKernel(queue, kernel, dim, NULL, ndrange, NULL, 0, NULL, NULL); |
| if (error != CL_SUCCESS) |
| { |
| return error; |
| } |
| |
| clFinish(queue); |
| |
| // check output buffer has not changed. |
| int* output = reinterpret_cast<int*>(clEnqueueMapBuffer(queue, buf, CL_TRUE, CL_MAP_READ, 0, sizeof(int) * bufSize, 0, NULL, NULL, &error)); |
| if (error != CL_SUCCESS) |
| { |
| return error; |
| } |
| |
| for (int i = 0; i < bufSize; ++i) |
| { |
| if (output[i] != 0) |
| { |
| log_error( "ERROR: output buffer value has changed.\n" ); |
| return CL_INVALID_OPERATION; |
| } |
| } |
| |
| return clEnqueueUnmapMemObject(queue, buf, output, 0, NULL, NULL); |
| } |
| |
| int test_zero_sized_enqueue_helper(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) |
| { |
| int error; |
| clProgramWrapper program; |
| clKernelWrapper kernel; |
| clMemWrapper output_stream; |
| size_t ndrange1 = 0; |
| size_t ndrange20[2] = {0, 0}; |
| size_t ndrange21[2] = {1, 0}; |
| size_t ndrange22[2] = {0, 1}; |
| |
| size_t ndrange30[3] = {0, 0, 0}; |
| size_t ndrange31[3] = {1, 0, 0}; |
| size_t ndrange32[3] = {0, 1, 0}; |
| size_t ndrange33[3] = {0, 0, 1}; |
| size_t ndrange34[3] = {0, 1, 1}; |
| size_t ndrange35[3] = {1, 0, 1}; |
| size_t ndrange36[3] = {1, 1, 0}; |
| |
| output_stream = |
| clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR, |
| bufSize * sizeof(int), NULL, &error); |
| |
| // Initialise output buffer. |
| int output_buffer_data = 0; |
| error = clEnqueueFillBuffer(queue, output_stream, &output_buffer_data, |
| sizeof(int), 0, sizeof(int) * bufSize, 0, NULL, |
| NULL); |
| |
| /* Create a kernel to test with */ |
| if( create_single_kernel_helper( context, &program, &kernel, 1, zero_sized_enqueue_test_kernel, "foo_kernel" ) != 0 ) |
| { |
| return -1; |
| } |
| |
| error = clSetKernelArg(kernel, 0, sizeof(cl_mem), &output_stream); |
| test_error( error, "clSetKernelArg failed." ); |
| |
| // Simple API return code tests for 1D, 2D and 3D zero sized ND range. |
| error = test_zero_sized_enqueue_and_test_output_buffer( |
| queue, kernel, output_stream, 1, &ndrange1); |
| test_error( error, "1D zero sized kernel enqueue failed." ); |
| |
| error = test_zero_sized_enqueue_and_test_output_buffer( |
| queue, kernel, output_stream, 2, ndrange20); |
| test_error( error, "2D zero sized kernel enqueue failed." ); |
| |
| error = test_zero_sized_enqueue_and_test_output_buffer( |
| queue, kernel, output_stream, 2, ndrange21); |
| test_error( error, "2D zero sized kernel enqueue failed." ); |
| |
| error = test_zero_sized_enqueue_and_test_output_buffer( |
| queue, kernel, output_stream, 2, ndrange22); |
| test_error( error, "2D zero sized kernel enqueue failed." ); |
| |
| |
| error = test_zero_sized_enqueue_and_test_output_buffer( |
| queue, kernel, output_stream, 3, ndrange30); |
| test_error( error, "3D zero sized kernel enqueue failed." ); |
| |
| error = test_zero_sized_enqueue_and_test_output_buffer( |
| queue, kernel, output_stream, 3, ndrange31); |
| test_error( error, "3D zero sized kernel enqueue failed." ); |
| |
| error = test_zero_sized_enqueue_and_test_output_buffer( |
| queue, kernel, output_stream, 3, ndrange32); |
| test_error( error, "3D zero sized kernel enqueue failed." ); |
| |
| error = test_zero_sized_enqueue_and_test_output_buffer( |
| queue, kernel, output_stream, 3, ndrange33); |
| test_error( error, "3D zero sized kernel enqueue failed." ); |
| |
| error = test_zero_sized_enqueue_and_test_output_buffer( |
| queue, kernel, output_stream, 3, ndrange34); |
| test_error( error, "3D zero sized kernel enqueue failed." ); |
| |
| error = test_zero_sized_enqueue_and_test_output_buffer( |
| queue, kernel, output_stream, 3, ndrange35); |
| test_error( error, "3D zero sized kernel enqueue failed." ); |
| |
| error = test_zero_sized_enqueue_and_test_output_buffer( |
| queue, kernel, output_stream, 3, ndrange36); |
| test_error( error, "3D zero sized kernel enqueue failed." ); |
| |
| // Verify zero-sized ND range kernel still satisfy event wait list and correct event object |
| // is returned |
| clEventWrapper ev = NULL; |
| clEventWrapper user_ev = clCreateUserEvent(context, &error); |
| test_error( error, "user event creation failed." ); |
| error = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, ndrange30, NULL, 1, &user_ev, &ev); |
| test_error( error, "3D zero sized kernel enqueue failed." ); |
| if (ev == NULL) |
| { |
| log_error( "ERROR: failed to create an event object\n" ); |
| return -1; |
| } |
| |
| cl_int sta; |
| error = clGetEventInfo(ev, CL_EVENT_COMMAND_EXECUTION_STATUS, sizeof(cl_int), &sta, NULL); |
| test_error( error, "Failed to get event status."); |
| |
| if (sta != CL_QUEUED && sta != CL_SUBMITTED) |
| { |
| log_error( "ERROR: incorrect zero sized kernel enqueue event status.\n" ); |
| return -1; |
| } |
| |
| // now unblock zero-sized enqueue |
| error = clSetUserEventStatus(user_ev, CL_COMPLETE); |
| test_error( error, "Failed to set user event status."); |
| |
| clFinish(queue); |
| |
| // now check zero sized enqueue event status |
| error = clGetEventInfo(ev, CL_EVENT_COMMAND_EXECUTION_STATUS, sizeof(cl_int), &sta, NULL); |
| test_error( error, "Failed to get event status."); |
| |
| if (sta != CL_COMPLETE) |
| { |
| log_error( "ERROR: incorrect zero sized kernel enqueue event status.\n" ); |
| return -1; |
| } |
| |
| return 0; |
| } |
| |
| |
| int test_zero_sized_enqueue(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) |
| { |
| int res = test_zero_sized_enqueue_helper(deviceID, context, queue, num_elements); |
| if (res != 0) |
| { |
| return res; |
| } |
| |
| // now test out of order queue |
| cl_command_queue_properties props; |
| cl_int error = clGetDeviceInfo(deviceID, CL_DEVICE_QUEUE_PROPERTIES, sizeof(cl_command_queue_properties), &props, NULL); |
| test_error( error, "clGetDeviceInfo failed."); |
| |
| if (props & CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE) |
| { |
| // test out of order queue |
| cl_queue_properties queue_prop_def[] = |
| { |
| CL_QUEUE_PROPERTIES, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, |
| 0 |
| }; |
| |
| clCommandQueueWrapper ooqueue = clCreateCommandQueueWithProperties(context, deviceID, queue_prop_def, &error); |
| test_error( error, "clCreateCommandQueueWithProperties failed."); |
| |
| res = test_zero_sized_enqueue_helper(deviceID, context, ooqueue, num_elements); |
| } |
| |
| return res; |
| } |