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 "harness/compat.h"
17 #include "harness/imageHelpers.h"
18
19 #include <stdio.h>
20 #include <stdlib.h>
21 #include <math.h>
22 #include <string.h>
23 #include <sys/types.h>
24 #include <sys/stat.h>
25
26 #include "procs.h"
27
28 static const char *bgra8888_kernel_code =
29 "\n"
30 "__kernel void test_bgra8888(read_only image3d_t srcimg, __global float4 *dst, sampler_t sampler)\n"
31 "{\n"
32 " int tid_x = get_global_id(0);\n"
33 " int tid_y = get_global_id(1);\n"
34 " int tid_z = get_global_id(2);\n"
35 " int indx = (tid_z * get_image_height(srcimg) + tid_y) * get_image_width(srcimg) + tid_x;\n"
36 " float4 color;\n"
37 "\n"
38 " color = read_imagef(srcimg, sampler, (int4)(tid_x, tid_y, tid_z, 0));\n"
39 " dst[indx].x = color.z;\n"
40 " dst[indx].y = color.y;\n"
41 " dst[indx].z = color.x;\n"
42 " dst[indx].w = color.w;\n"
43 "\n"
44 "}\n";
45
46 static const char *rgba8888_kernel_code =
47 "\n"
48 "__kernel void test_rgba8888(read_only image3d_t srcimg, __global float4 *dst, sampler_t sampler)\n"
49 "{\n"
50 " int tid_x = get_global_id(0);\n"
51 " int tid_y = get_global_id(1);\n"
52 " int tid_z = get_global_id(2);\n"
53 " int indx = (tid_z * get_image_height(srcimg) + tid_y) * get_image_width(srcimg) + tid_x;\n"
54 " float4 color;\n"
55 "\n"
56 " color = read_imagef(srcimg, sampler, (int4)(tid_x, tid_y, tid_z, 0));\n"
57 " //indx *= 4;\n"
58 " dst[indx].x = color.x;\n"
59 " dst[indx].y = color.y;\n"
60 " dst[indx].z = color.z;\n"
61 " dst[indx].w = color.w;\n"
62 "\n"
63 "}\n";
64
65 static unsigned char *
generate_3d_image8(int w,int h,int d,MTdata data)66 generate_3d_image8(int w, int h, int d, MTdata data)
67 {
68 unsigned char *ptr = (unsigned char*)malloc(w * h * d * 4);
69 int i;
70
71 for (i=0; i<w*h*d*4; i++)
72 ptr[i] = (unsigned char)genrand_int32(data);
73
74 return ptr;
75 }
76
77 static int
verify_3d_image8(double * image,float * outptr,int w,int h,int d)78 verify_3d_image8(double *image, float *outptr, int w, int h, int d)
79 {
80 int i;
81
82 for (i=0; i<w*h*d*4; i++)
83 {
84 if (outptr[i] != (float)image[i])
85 {
86 float ulps = Ulp_Error( outptr[i], image[i]);
87
88 if(! (fabsf(ulps) < 1.5f) )
89 {
90 log_error( "ERROR: Data sample %d does not validate! Expected (%a), got (%a), ulp %f\n",
91 (int)i, image[i], outptr[ i ], ulps );
92 return -1;
93 }
94 }
95 }
96
97 return 0;
98 }
99
100 static double *
prepare_reference(unsigned char * input_ptr,int w,int h,int d)101 prepare_reference(unsigned char * input_ptr, int w, int h, int d)
102 {
103 double *ptr = (double*)malloc(w * h * d * 4 * sizeof(double));
104 int i;
105 for (i=0; i<w*h*d*4; i++)
106 ptr[i] = ((double)input_ptr[i]/255);
107
108 return ptr;
109 }
110
test_readimage3d(cl_device_id device,cl_context context,cl_command_queue queue,int num_elements)111 int test_readimage3d(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
112 {
113 cl_mem streams[2];
114 cl_program program;
115 cl_kernel kernel;
116 cl_sampler sampler;
117 struct testFormat
118 {
119 const char* kernelName;
120 const char* kernelSourceString;
121 const cl_image_format img_format;
122 };
123
124 static testFormat formatsToTest[] =
125 {
126 {
127 "test_bgra8888",
128 bgra8888_kernel_code,
129 {CL_BGRA, CL_UNORM_INT8},
130 },
131 {
132 "test_rgba8888",
133 rgba8888_kernel_code,
134 {CL_RGBA, CL_UNORM_INT8},
135 },
136 };
137
138 unsigned char *input_ptr;
139 float *output_ptr;
140 double *ref_ptr;
141 size_t threads[3];
142 int img_width = 64;
143 int img_height = 64;
144 int img_depth = 64;
145 int i, err;
146 size_t origin[3] = {0, 0, 0};
147 size_t region[3] = {img_width, img_height, img_depth};
148 size_t length = img_width * img_height * img_depth * 4 * sizeof(float);
149
150 PASSIVE_REQUIRE_3D_IMAGE_SUPPORT( device )
151
152 for (uint32_t i = 0; i < ARRAY_SIZE(formatsToTest); i++)
153 {
154 if (!is_image_format_required(formatsToTest[i].img_format, CL_MEM_READ_ONLY, CL_MEM_OBJECT_IMAGE3D, device))
155 continue;
156
157 MTdata d = init_genrand( gRandomSeed );
158 input_ptr = generate_3d_image8(img_width, img_height, img_depth, d);
159 ref_ptr = prepare_reference(input_ptr, img_width, img_height, img_depth);
160 output_ptr = (float*)malloc(length);
161
162 streams[0] = create_image_3d(context, CL_MEM_READ_ONLY, &formatsToTest[i].img_format, img_width, img_height, img_depth, 0, 0, NULL, &err);
163 test_error(err, "create_image_3d failed");
164
165 streams[1] = clCreateBuffer(context, CL_MEM_READ_WRITE, length, NULL, &err);
166 test_error(err, "clCreateBuffer failed");
167
168 sampler = clCreateSampler(context, CL_FALSE, CL_ADDRESS_CLAMP_TO_EDGE, CL_FILTER_NEAREST, &err);
169 test_error(err, "clCreateSampler failed");
170
171 err = clEnqueueWriteImage(queue, streams[0], CL_TRUE, origin, region, 0, 0, input_ptr, 0, NULL, NULL);
172 test_error(err, "clEnqueueWriteImage failed");
173
174 err = create_single_kernel_helper(context, &program, &kernel, 1, &formatsToTest[i].kernelSourceString, formatsToTest[i].kernelName);
175 test_error(err, "create_single_kernel_helper failed");
176
177 err = clSetKernelArg(kernel, 0, sizeof streams[0], &streams[0]);
178 err |= clSetKernelArg(kernel, 1, sizeof streams[1], &streams[1]);
179 err |= clSetKernelArg(kernel, 2, sizeof sampler, &sampler);
180 test_error(err, "clSetKernelArg failed");
181
182 threads[0] = (unsigned int)img_width;
183 threads[1] = (unsigned int)img_height;
184 threads[2] = (unsigned int)img_depth;
185
186 err = clEnqueueNDRangeKernel(queue, kernel, 3, NULL, threads, NULL, 0, NULL, NULL);
187 test_error(err, "clEnqueueNDRangeKernel failed");
188
189 err = clEnqueueReadBuffer(queue, streams[1], CL_TRUE, 0, length, output_ptr, 0, NULL, NULL);
190 test_error(err, "clEnqueueReadBuffer failed");
191
192 err = verify_3d_image8(ref_ptr, output_ptr, img_width, img_height, img_depth);
193 if ( err == 0 )
194 {
195 log_info("READ_IMAGE3D_%s_%s test passed\n",
196 GetChannelTypeName(formatsToTest[i].img_format.image_channel_data_type),
197 GetChannelOrderName(formatsToTest[i].img_format.image_channel_order));
198 }
199
200 clReleaseSampler(sampler);
201 clReleaseMemObject(streams[0]);
202 clReleaseMemObject(streams[1]);
203 clReleaseKernel(kernel);
204 clReleaseProgram(program);
205 free_mtdata(d);
206 d = NULL;
207 free(input_ptr);
208 free(ref_ptr);
209 free(output_ptr);
210 }
211
212 return err;
213 }
214