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 MTdata mtData = init_genrand(gRandomSeed);
102 /* Fill input array with random values */
103 for (i = 0; i < size; i++) {
104 input[i] = genrand_int32(mtData);
105 }
106 free_mtdata(mtData);
107 mtData = NULL;
108
109 /* Zero out output array */
110 for (i = 0; i < size; i++) {
111 output[i] = 0.0f;
112 }
113
114 /* Build the program executable */
115 err = create_single_kernel_helper(context, &program, &kernel, 1,
116 &rw_kernel_code, "test_rw_images");
117 if (err != CL_SUCCESS || !program) {
118 log_error("Error: clCreateProgramWithSource failed\n");
119 return err;
120 }
121
122 /* Create arrays for input and output data */
123 format.image_channel_order = CL_RGBA;
124 format.image_channel_data_type = CL_UNSIGNED_INT32;
125
126 /* Create input image */
127 flags = CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR;
128 src_image = create_image_2d(context, flags, &format,
129 size_x, size_y, 0,
130 (void *)input, &err);
131 if (err != CL_SUCCESS || !src_image) {
132 log_error("Error: clCreateImage2D failed\n");
133 return err;
134 }
135
136 /* Set kernel arguments */
137 err = clSetKernelArg(kernel, 0, sizeof(src_image), &src_image);
138 if (err != CL_SUCCESS) {
139 log_error("Error: clSetKernelArg failed\n");
140 return err;
141 }
142
143 /* Set kernel execution parameters */
144 int dim_count = 2;
145 size_t global_dim[2];
146 size_t local_dim[2];
147
148 global_dim[0] = size_x;
149 global_dim[1] = size_y;
150
151 local_dim[0] = 1;
152 local_dim[1] = 1;
153
154 /* Execute kernel */
155 err = CL_SUCCESS;
156 unsigned int num_iter = 1;
157 for(i = 0; i < num_iter; i++) {
158 err |= clEnqueueNDRangeKernel(commands, kernel, dim_count,
159 NULL, global_dim, local_dim,
160 0, NULL, NULL);
161 }
162
163 /* Read back the results from the device to verify the output */
164 const size_t origin[3] = {0, 0, 0};
165 const size_t region[3] = {size_x, size_y, 1};
166 err |= clEnqueueReadImage(commands, src_image, CL_TRUE, origin, region, 0, 0,
167 output, 0, NULL, NULL);
168 if (err != CL_SUCCESS) {
169 log_error("Error: clEnqueueReadBuffer failed\n");
170 return err;
171 }
172
173 /* Verify the correctness of kernel result */
174 err = 0;
175 for (i = 0; i < size; i++) {
176 if (output[i] != (input[i] + 3)) {
177 log_error("Error: mismatch at index %d\n", i);
178 err++;
179 break;
180 }
181 }
182
183 /* Release programs, kernel, contect, and memory objects */
184 clReleaseMemObject(src_image);
185 clReleaseProgram(program);
186 clReleaseKernel(kernel);
187
188 /* Deallocate arrays */
189 free(input);
190 free(output);
191
192 return err;
193 }
194