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 "../testBase.h"
17 #include <float.h>
18 
19 #if defined( __APPLE__ )
20     #include <signal.h>
21     #include <sys/signal.h>
22     #include <setjmp.h>
23 #endif
24 
25 
26 const char *read1DBufferKernelSourcePattern =
27 "__kernel void sample_kernel( read_only image1d_buffer_t inputA, read_only image1d_t inputB, sampler_t sampler, __global int *results )\n"
28 "{\n"
29 "   int tidX = get_global_id(0);\n"
30 "   int offset = tidX;\n"
31 "   %s clr = read_image%s( inputA, tidX );\n"
32 "   int4 test = (clr != read_image%s( inputB, sampler, tidX ));\n"
33 "   if ( test.x || test.y || test.z || test.w )\n"
34 "      results[offset] = -1;\n"
35 "   else\n"
36 "      results[offset] = 0;\n"
37 "}";
38 
39 
test_read_image_1D_buffer(cl_context context,cl_command_queue queue,cl_kernel kernel,image_descriptor * imageInfo,image_sampler_data * imageSampler,ExplicitType outputType,MTdata d)40 int test_read_image_1D_buffer( cl_context context, cl_command_queue queue, cl_kernel kernel,
41                         image_descriptor *imageInfo, image_sampler_data *imageSampler,
42                         ExplicitType outputType, MTdata d )
43 {
44     int error;
45     size_t threads[2];
46     cl_sampler actualSampler;
47 
48     BufferOwningPtr<char> imageValues;
49     generate_random_image_data( imageInfo, imageValues, d );
50 
51     if ( gDebugTrace )
52         log_info( " - Creating 1D image from buffer %d ...\n", (int)imageInfo->width );
53 
54     // Construct testing sources
55     cl_mem image[2];
56     cl_image_desc image_desc;
57 
58     cl_mem imageBuffer = clCreateBuffer( context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, imageInfo->rowPitch, imageValues, &error);
59     if ( error != CL_SUCCESS )
60     {
61         log_error( "ERROR: Unable to create buffer of size %d bytes (%s)\n", (int)imageInfo->rowPitch, IGetErrorString( error ) );
62         return error;
63     }
64 
65     memset(&image_desc, 0x0, sizeof(cl_image_desc));
66     image_desc.image_type = CL_MEM_OBJECT_IMAGE1D_BUFFER;
67     image_desc.image_width = imageInfo->width;
68     image_desc.mem_object = imageBuffer;
69     image[0] = clCreateImage( context, CL_MEM_READ_ONLY, imageInfo->format,
70         &image_desc, NULL, &error );
71     if ( error != CL_SUCCESS )
72     {
73         log_error( "ERROR: Unable to create IMAGE1D_BUFFER of size %d pitch %d (%s)\n", (int)imageInfo->width, (int)imageInfo->rowPitch, IGetErrorString( error ) );
74         return error;
75     }
76 
77     cl_mem ret = NULL;
78     error = clGetMemObjectInfo(image[0], CL_MEM_ASSOCIATED_MEMOBJECT, sizeof(ret), &ret, NULL);
79     if ( error != CL_SUCCESS )
80     {
81       log_error( "ERROR: Unable to query CL_MEM_ASSOCIATED_MEMOBJECT\n", IGetErrorString( error ) );
82       return error;
83     }
84 
85     if (ret != imageBuffer) {
86       log_error("ERROR: clGetImageInfo for CL_IMAGE_BUFFER returned wrong value\n");
87       return -1;
88     }
89 
90     memset(&image_desc, 0x0, sizeof(cl_image_desc));
91     image_desc.image_type = CL_MEM_OBJECT_IMAGE1D;
92     image_desc.image_width = imageInfo->width;
93     image[1] = clCreateImage( context, CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR, imageInfo->format, &image_desc, imageValues, &error );
94     if ( error != CL_SUCCESS )
95     {
96         log_error( "ERROR: Unable to create IMAGE1D of size %d pitch %d (%s)\n", (int)imageInfo->width, (int)imageInfo->rowPitch, IGetErrorString( error ) );
97         return error;
98     }
99 
100     if ( gDebugTrace )
101         log_info( " - Creating kernel arguments...\n" );
102 
103     // Create sampler to use
104     actualSampler = clCreateSampler( context, CL_FALSE, CL_ADDRESS_NONE, CL_FILTER_NEAREST, &error );
105     test_error( error, "Unable to create image sampler" );
106 
107     // Create results buffer
108     cl_mem results = clCreateBuffer( context, 0, imageInfo->width * sizeof(cl_int), NULL, &error);
109     test_error( error, "Unable to create results buffer" );
110 
111     size_t resultValuesSize = imageInfo->width * sizeof(cl_int);
112     BufferOwningPtr<int> resultValues(malloc( resultValuesSize ));
113     memset( resultValues, 0xff, resultValuesSize );
114     clEnqueueWriteBuffer( queue, results, CL_TRUE, 0, resultValuesSize, resultValues, 0, NULL, NULL );
115 
116     // Set arguments
117     int idx = 0;
118     error = clSetKernelArg( kernel, idx++, sizeof( cl_mem ), &image[0] );
119     test_error( error, "Unable to set kernel arguments" );
120     error = clSetKernelArg( kernel, idx++, sizeof( cl_mem ), &image[1] );
121     test_error( error, "Unable to set kernel arguments" );
122     error = clSetKernelArg( kernel, idx++, sizeof( cl_sampler ), &actualSampler );
123     test_error( error, "Unable to set kernel arguments" );
124     error = clSetKernelArg( kernel, idx++, sizeof( cl_mem ), &results );
125     test_error( error, "Unable to set kernel arguments" );
126 
127     // Run the kernel
128     threads[0] = (size_t)imageInfo->width;
129     error = clEnqueueNDRangeKernel( queue, kernel, 1, NULL, threads, NULL, 0, NULL, NULL );
130     test_error( error, "Unable to run kernel" );
131 
132     if ( gDebugTrace )
133         log_info( "    reading results, %ld kbytes\n", (unsigned long)( imageInfo->width * sizeof(cl_int) / 1024 ) );
134 
135     error = clEnqueueReadBuffer( queue, results, CL_TRUE, 0, resultValuesSize, resultValues, 0, NULL, NULL );
136     test_error( error, "Unable to read results from kernel" );
137     if ( gDebugTrace )
138         log_info( "    results read\n" );
139 
140     // Check for non-zero comps
141     bool allZeroes = true;
142     for ( size_t ic = 0; ic < imageInfo->width; ++ic )
143     {
144         if ( resultValues[ic] ) {
145             allZeroes = false;
146             break;
147         }
148     }
149     if ( !allZeroes )
150     {
151         log_error( " Sampler-less reads differ from reads with sampler.\n" );
152         return -1;
153     }
154 
155     clReleaseSampler(actualSampler);
156     clReleaseMemObject(results);
157     clReleaseMemObject(image[0]);
158     clReleaseMemObject(image[1]);
159     clReleaseMemObject(imageBuffer);
160     return 0;
161 }
162 
test_read_image_set_1D_buffer(cl_device_id device,cl_context context,cl_command_queue queue,const cl_image_format * format,image_sampler_data * imageSampler,ExplicitType outputType)163 int test_read_image_set_1D_buffer(cl_device_id device, cl_context context,
164                                   cl_command_queue queue,
165                                   const cl_image_format *format,
166                                   image_sampler_data *imageSampler,
167                                   ExplicitType outputType)
168 {
169     char programSrc[10240];
170     const char *ptr;
171     const char *readFormat;
172     const char *dataType;
173     clProgramWrapper program;
174     clKernelWrapper kernel;
175     RandomSeed seed( gRandomSeed );
176     int error;
177 
178     // Get our operating params
179     size_t maxWidth, maxWidth1D;
180     cl_ulong maxAllocSize, memSize;
181     image_descriptor imageInfo = { 0 };
182     size_t pixelSize;
183 
184     if (format->image_channel_order == CL_RGB || format->image_channel_order == CL_RGBx)
185     {
186         switch (format->image_channel_data_type)
187         {
188             case CL_UNORM_INT8:
189             case CL_UNORM_INT16:
190             case CL_SNORM_INT8:
191             case CL_SNORM_INT16:
192             case CL_HALF_FLOAT:
193             case CL_FLOAT:
194             case CL_SIGNED_INT8:
195             case CL_SIGNED_INT16:
196             case CL_SIGNED_INT32:
197             case CL_UNSIGNED_INT8:
198             case CL_UNSIGNED_INT16:
199             case CL_UNSIGNED_INT32:
200             case CL_UNORM_INT_101010:
201                 log_info( "Skipping image format: %s %s\n", GetChannelOrderName( format->image_channel_order ),
202                          GetChannelTypeName( format->image_channel_data_type ));
203                 return 0;
204             default:
205                 break;
206         }
207     }
208 
209     imageInfo.format = format;
210     imageInfo.height = imageInfo.depth = imageInfo.arraySize = imageInfo.slicePitch = 0;
211     imageInfo.type = CL_MEM_OBJECT_IMAGE1D;
212     pixelSize = get_pixel_size( imageInfo.format );
213 
214     error = clGetDeviceInfo( device, CL_DEVICE_IMAGE_MAX_BUFFER_SIZE, sizeof( maxWidth ), &maxWidth, NULL );
215     error |= clGetDeviceInfo( device, CL_DEVICE_MAX_MEM_ALLOC_SIZE, sizeof( maxAllocSize ), &maxAllocSize, NULL );
216     error |= clGetDeviceInfo( device, CL_DEVICE_GLOBAL_MEM_SIZE, sizeof( memSize ), &memSize, NULL );
217     error |= clGetDeviceInfo( device, CL_DEVICE_IMAGE2D_MAX_WIDTH, sizeof( maxWidth ), &maxWidth1D, NULL );
218     test_error( error, "Unable to get max image 1D buffer size from device" );
219 
220     if (memSize > (cl_ulong)SIZE_MAX) {
221       memSize = (cl_ulong)SIZE_MAX;
222     }
223 
224     // note: image_buffer test uses image1D for results validation.
225     // So the test can't use the biggest possible size for image_buffer if it's bigger than the max image1D size
226     maxWidth = (maxWidth > maxWidth1D) ? maxWidth1D : maxWidth;
227     // Determine types
228     if ( outputType == kInt )
229     {
230         readFormat = "i";
231         dataType = "int4";
232     }
233     else if ( outputType == kUInt )
234     {
235         readFormat = "ui";
236         dataType = "uint4";
237     }
238     else // kFloat
239     {
240         readFormat = "f";
241         dataType = "float4";
242     }
243 
244     sprintf( programSrc, read1DBufferKernelSourcePattern, dataType,
245              readFormat,
246              readFormat );
247 
248     ptr = programSrc;
249     error = create_single_kernel_helper(context, &program, &kernel, 1, &ptr,
250                                         "sample_kernel");
251     test_error( error, "Unable to create testing kernel" );
252 
253     if ( gTestSmallImages )
254     {
255         for ( imageInfo.width = 1; imageInfo.width < 13; imageInfo.width++ )
256         {
257             imageInfo.rowPitch = imageInfo.width * pixelSize;
258             {
259                 if ( gDebugTrace )
260                     log_info( "   at size %d\n", (int)imageInfo.width );
261 
262                 int retCode = test_read_image_1D_buffer( context, queue, kernel, &imageInfo, imageSampler, outputType, seed );
263                 if ( retCode )
264                     return retCode;
265             }
266         }
267     }
268     else if ( gTestMaxImages )
269     {
270         // Try a specific set of maximum sizes
271         size_t numbeOfSizes;
272         size_t sizes[100][3];
273 
274         get_max_sizes(&numbeOfSizes, 100, sizes, maxWidth, 1, 1, 1, maxAllocSize, memSize, CL_MEM_OBJECT_IMAGE1D, imageInfo.format);
275 
276         for ( size_t idx = 0; idx < numbeOfSizes; idx++ )
277         {
278             imageInfo.width = sizes[ idx ][ 0 ];
279             imageInfo.rowPitch = imageInfo.width * pixelSize;
280             log_info("Testing %d\n", (int)sizes[ idx ][ 0 ]);
281             if ( gDebugTrace )
282                 log_info( "   at max size %d\n", (int)sizes[ idx ][ 0 ] );
283             int retCode = test_read_image_1D_buffer( context, queue, kernel, &imageInfo, imageSampler, outputType, seed );
284             if ( retCode )
285                 return retCode;
286         }
287     }
288     else
289     {
290         for ( int i = 0; i < NUM_IMAGE_ITERATIONS; i++ )
291         {
292             cl_ulong size;
293             // Loop until we get a size that a) will fit in the max alloc size and b) that an allocation of that
294             // image, the result array, plus offset arrays, will fit in the global ram space
295             do
296             {
297                 imageInfo.width = (size_t)random_log_in_range( 16, (int)maxWidth / 32, seed );
298                 imageInfo.rowPitch = imageInfo.width * pixelSize;
299                 size = (size_t)imageInfo.rowPitch * 4;
300             } while (  size > maxAllocSize || ( size * 3 ) > memSize );
301 
302             if ( gDebugTrace )
303                 log_info( "   at size %d (row pitch %d) out of %d\n", (int)imageInfo.width, (int)imageInfo.rowPitch, (int)maxWidth );
304             int retCode = test_read_image_1D_buffer( context, queue, kernel, &imageInfo, imageSampler, outputType, seed );
305             if ( retCode )
306                 return retCode;
307         }
308     }
309 
310     return 0;
311 }
312 
313 
314