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