1 //
2 // Copyright (c) 2017 The Khronos Group Inc.
3 //
4 // Licensed under the Apache License, Version 2.0 (the "License");
5 // you may not use this file except in compliance with the License.
6 // You may obtain a copy of the License at
7 //
8 // http://www.apache.org/licenses/LICENSE-2.0
9 //
10 // Unless required by applicable law or agreed to in writing, software
11 // distributed under the License is distributed on an "AS IS" BASIS,
12 // WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
13 // See the License for the specific language governing permissions and
14 // limitations under the License.
15 //
16 #include <stdlib.h>
17 #include <stdio.h>
18 #include <string.h>
19 #include <sys/stat.h>
20
21 #include "procs.h"
22 #include "harness/clImageHelper.h"
23
24 static const char* rw_kernel_code =
25 "kernel void test_rw_images(read_write image2d_t src_image) {\n"
26 " int tid_x = get_global_id(0);\n"
27 " int tid_y = get_global_id(1);\n"
28 "\n"
29 " int2 coords = (int2)(tid_x, tid_y);\n"
30 "\n"
31 " uint4 src_val = read_imageui(src_image, coords);\n"
32 " src_val += 3;\n"
33 "\n"
34 " // required to ensure that following read from image at\n"
35 " // location coord returns the latest color value.\n"
36 " atomic_work_item_fence(CLK_IMAGE_MEM_FENCE,\n"
37 " memory_order_acq_rel,\n"
38 " memory_scope_work_item);\n"
39 "\n"
40 " write_imageui(src_image, coords, src_val);\n"
41 "}\n";
42
43
test_rw_image_access_qualifier(cl_device_id device_id,cl_context context,cl_command_queue commands,int num_elements)44 int test_rw_image_access_qualifier(cl_device_id device_id, cl_context context, cl_command_queue commands, int num_elements)
45 {
46 // This test should be skipped if images are not supported.
47 if (checkForImageSupport(device_id))
48 {
49 return TEST_SKIPPED_ITSELF;
50 }
51
52 // Support for read-write image arguments is required for an
53 // or 2.X device if the device supports images. In OpenCL-3.0
54 // read-write images are optional. This test is already being skipped
55 // for 1.X devices.
56 if (get_device_cl_version(device_id) >= Version(3, 0))
57 {
58 cl_uint are_rw_images_supported{};
59 test_error(
60 clGetDeviceInfo(device_id, CL_DEVICE_MAX_READ_WRITE_IMAGE_ARGS,
61 sizeof(are_rw_images_supported),
62 &are_rw_images_supported, nullptr),
63 "clGetDeviceInfo failed for CL_DEVICE_MAX_READ_IMAGE_ARGS\n");
64 if (0 == are_rw_images_supported)
65 {
66 return TEST_SKIPPED_ITSELF;
67 }
68 }
69
70 unsigned int i;
71
72 unsigned int size_x;
73 unsigned int size_y;
74 unsigned int size;
75
76 cl_int err;
77
78 cl_program program;
79 cl_kernel kernel;
80
81 cl_mem_flags flags;
82 cl_image_format format;
83 cl_mem src_image;
84
85 unsigned int *input;
86 unsigned int *output;
87
88 /* Create test input */
89 size_x = 4;
90 size_y = 4;
91 size = size_x * size_y * 4;
92
93 input = (unsigned int *)malloc(size*sizeof(unsigned int));
94 output = (unsigned int *)malloc(size*sizeof(unsigned int));
95
96 if (!input && !output) {
97 log_error("Error: memory allocation failed\n");
98 return -1;
99 }
100
101 /* Fill input array with random values */
102 for (i = 0; i < size; i++) {
103 input[i] = (unsigned int)(rand()/((double)RAND_MAX + 1)*255);
104 }
105
106 /* Zero out output array */
107 for (i = 0; i < size; i++) {
108 output[i] = 0.0f;
109 }
110
111 /* Build the program executable */
112 err = create_single_kernel_helper(context, &program, &kernel, 1,
113 &rw_kernel_code, "test_rw_images");
114 if (err != CL_SUCCESS || !program) {
115 log_error("Error: clCreateProgramWithSource failed\n");
116 return err;
117 }
118
119 /* Create arrays for input and output data */
120 format.image_channel_order = CL_RGBA;
121 format.image_channel_data_type = CL_UNSIGNED_INT32;
122
123 /* Create input image */
124 flags = CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR;
125 src_image = create_image_2d(context, flags, &format,
126 size_x, size_y, 0,
127 (void *)input, &err);
128 if (err != CL_SUCCESS || !src_image) {
129 log_error("Error: clCreateImage2D failed\n");
130 return err;
131 }
132
133 /* Set kernel arguments */
134 err = clSetKernelArg(kernel, 0, sizeof(src_image), &src_image);
135 if (err != CL_SUCCESS) {
136 log_error("Error: clSetKernelArg failed\n");
137 return err;
138 }
139
140 /* Set kernel execution parameters */
141 int dim_count = 2;
142 size_t global_dim[2];
143 size_t local_dim[2];
144
145 global_dim[0] = size_x;
146 global_dim[1] = size_y;
147
148 local_dim[0] = 1;
149 local_dim[1] = 1;
150
151 /* Execute kernel */
152 err = CL_SUCCESS;
153 unsigned int num_iter = 1;
154 for(i = 0; i < num_iter; i++) {
155 err |= clEnqueueNDRangeKernel(commands, kernel, dim_count,
156 NULL, global_dim, local_dim,
157 0, NULL, NULL);
158 }
159
160 /* Read back the results from the device to verify the output */
161 const size_t origin[3] = {0, 0, 0};
162 const size_t region[3] = {size_x, size_y, 1};
163 err |= clEnqueueReadImage(commands, src_image, CL_TRUE, origin, region, 0, 0,
164 output, 0, NULL, NULL);
165 if (err != CL_SUCCESS) {
166 log_error("Error: clEnqueueReadBuffer failed\n");
167 return err;
168 }
169
170 /* Verify the correctness of kernel result */
171 err = 0;
172 for (i = 0; i < size; i++) {
173 if (output[i] != (input[i] + 3)) {
174 log_error("Error: mismatch at index %d\n", i);
175 err++;
176 break;
177 }
178 }
179
180 /* Release programs, kernel, contect, and memory objects */
181 clReleaseMemObject(src_image);
182 clReleaseProgram(program);
183 clReleaseKernel(kernel);
184
185 /* Deallocate arrays */
186 free(input);
187 free(output);
188
189 return err;
190 }
191