| // |
| // 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 <stdlib.h> |
| #include <stdio.h> |
| #include <string.h> |
| #include <sys/stat.h> |
| |
| #include "procs.h" |
| #include "harness/clImageHelper.h" |
| |
| static const char* rw_kernel_code = |
| "kernel void test_rw_images(read_write image2d_t src_image) {\n" |
| " int tid_x = get_global_id(0);\n" |
| " int tid_y = get_global_id(1);\n" |
| "\n" |
| " int2 coords = (int2)(tid_x, tid_y);\n" |
| "\n" |
| " uint4 src_val = read_imageui(src_image, coords);\n" |
| " src_val += 3;\n" |
| "\n" |
| " // required to ensure that following read from image at\n" |
| " // location coord returns the latest color value.\n" |
| " atomic_work_item_fence(CLK_IMAGE_MEM_FENCE,\n" |
| " memory_order_acq_rel,\n" |
| " memory_scope_work_item);\n" |
| "\n" |
| " write_imageui(src_image, coords, src_val);\n" |
| "}\n"; |
| |
| |
| int test_rw_image_access_qualifier(cl_device_id device_id, cl_context context, cl_command_queue commands, int num_elements) |
| { |
| // This test should be skipped if images are not supported. |
| if (checkForImageSupport(device_id)) |
| { |
| return TEST_SKIPPED_ITSELF; |
| } |
| |
| // Support for read-write image arguments is required for an |
| // or 2.X device if the device supports images. In OpenCL-3.0 |
| // read-write images are optional. This test is already being skipped |
| // for 1.X devices. |
| if (get_device_cl_version(device_id) >= Version(3, 0)) |
| { |
| cl_uint are_rw_images_supported{}; |
| test_error( |
| clGetDeviceInfo(device_id, CL_DEVICE_MAX_READ_WRITE_IMAGE_ARGS, |
| sizeof(are_rw_images_supported), |
| &are_rw_images_supported, nullptr), |
| "clGetDeviceInfo failed for CL_DEVICE_MAX_READ_IMAGE_ARGS\n"); |
| if (0 == are_rw_images_supported) |
| { |
| return TEST_SKIPPED_ITSELF; |
| } |
| } |
| |
| unsigned int i; |
| |
| unsigned int size_x; |
| unsigned int size_y; |
| unsigned int size; |
| |
| cl_int err; |
| |
| cl_program program; |
| cl_kernel kernel; |
| |
| cl_mem_flags flags; |
| cl_image_format format; |
| cl_mem src_image; |
| |
| unsigned int *input; |
| unsigned int *output; |
| |
| /* Create test input */ |
| size_x = 4; |
| size_y = 4; |
| size = size_x * size_y * 4; |
| |
| input = (unsigned int *)malloc(size*sizeof(unsigned int)); |
| output = (unsigned int *)malloc(size*sizeof(unsigned int)); |
| |
| if (!input && !output) { |
| log_error("Error: memory allocation failed\n"); |
| return -1; |
| } |
| |
| /* Fill input array with random values */ |
| for (i = 0; i < size; i++) { |
| input[i] = (unsigned int)(rand()/((double)RAND_MAX + 1)*255); |
| } |
| |
| /* Zero out output array */ |
| for (i = 0; i < size; i++) { |
| output[i] = 0.0f; |
| } |
| |
| /* Build the program executable */ |
| err = create_single_kernel_helper(context, &program, &kernel, 1, |
| &rw_kernel_code, "test_rw_images"); |
| if (err != CL_SUCCESS || !program) { |
| log_error("Error: clCreateProgramWithSource failed\n"); |
| return err; |
| } |
| |
| /* Create arrays for input and output data */ |
| format.image_channel_order = CL_RGBA; |
| format.image_channel_data_type = CL_UNSIGNED_INT32; |
| |
| /* Create input image */ |
| flags = CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR; |
| src_image = create_image_2d(context, flags, &format, |
| size_x, size_y, 0, |
| (void *)input, &err); |
| if (err != CL_SUCCESS || !src_image) { |
| log_error("Error: clCreateImage2D failed\n"); |
| return err; |
| } |
| |
| /* Set kernel arguments */ |
| err = clSetKernelArg(kernel, 0, sizeof(src_image), &src_image); |
| if (err != CL_SUCCESS) { |
| log_error("Error: clSetKernelArg failed\n"); |
| return err; |
| } |
| |
| /* Set kernel execution parameters */ |
| int dim_count = 2; |
| size_t global_dim[2]; |
| size_t local_dim[2]; |
| |
| global_dim[0] = size_x; |
| global_dim[1] = size_y; |
| |
| local_dim[0] = 1; |
| local_dim[1] = 1; |
| |
| /* Execute kernel */ |
| err = CL_SUCCESS; |
| unsigned int num_iter = 1; |
| for(i = 0; i < num_iter; i++) { |
| err |= clEnqueueNDRangeKernel(commands, kernel, dim_count, |
| NULL, global_dim, local_dim, |
| 0, NULL, NULL); |
| } |
| |
| /* Read back the results from the device to verify the output */ |
| const size_t origin[3] = {0, 0, 0}; |
| const size_t region[3] = {size_x, size_y, 1}; |
| err |= clEnqueueReadImage(commands, src_image, CL_TRUE, origin, region, 0, 0, |
| output, 0, NULL, NULL); |
| if (err != CL_SUCCESS) { |
| log_error("Error: clEnqueueReadBuffer failed\n"); |
| return err; |
| } |
| |
| /* Verify the correctness of kernel result */ |
| err = 0; |
| for (i = 0; i < size; i++) { |
| if (output[i] != (input[i] + 3)) { |
| log_error("Error: mismatch at index %d\n", i); |
| err++; |
| break; |
| } |
| } |
| |
| /* Release programs, kernel, contect, and memory objects */ |
| clReleaseMemObject(src_image); |
| clReleaseProgram(program); |
| clReleaseKernel(kernel); |
| |
| /* Deallocate arrays */ |
| free(input); |
| free(output); |
| |
| return err; |
| } |