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
18
19 struct image_kernel_data
20 {
21 cl_int width;
22 cl_int height;
23 cl_int depth;
24 cl_int widthDim;
25 cl_int heightDim;
26 cl_int depthDim;
27 cl_int channelType;
28 cl_int channelOrder;
29 cl_int expectedChannelType;
30 cl_int expectedChannelOrder;
31 };
32
33 static const char *methodTestKernelPattern =
34 "typedef struct {\n"
35 " int width;\n"
36 " int height;\n"
37 " int depth;\n"
38 " int widthDim;\n"
39 " int heightDim;\n"
40 " int depthDim;\n"
41 " int channelType;\n"
42 " int channelOrder;\n"
43 " int expectedChannelType;\n"
44 " int expectedChannelOrder;\n"
45 " } image_kernel_data;\n"
46 " %s\n"
47 "__kernel void sample_kernel( %s image%dd%s_t input, __global "
48 "image_kernel_data *outData )\n"
49 "{\n"
50 " outData->width = get_image_width( input );\n"
51 " outData->height = get_image_height( input );\n"
52 "%s\n"
53 " int%d dim = get_image_dim( input );\n"
54 " outData->widthDim = dim.x;\n"
55 " outData->heightDim = dim.y;\n"
56 "%s\n"
57 " outData->channelType = get_image_channel_data_type( input );\n"
58 " outData->channelOrder = get_image_channel_order( input );\n"
59 "\n"
60 " outData->expectedChannelType = %s;\n"
61 " outData->expectedChannelOrder = %s;\n"
62 "}";
63
64 static const char *depthKernelLine = " outData->depth = get_image_depth( input );\n";
65 static const char *depthDimKernelLine = " outData->depthDim = dim.z;\n";
66
test_get_image_info_single(cl_context context,cl_command_queue queue,image_descriptor * imageInfo,MTdata d,cl_mem_flags flags)67 int test_get_image_info_single(cl_context context, cl_command_queue queue,
68 image_descriptor *imageInfo, MTdata d,
69 cl_mem_flags flags)
70 {
71 int error = 0;
72
73 clProgramWrapper program;
74 clKernelWrapper kernel;
75 clMemWrapper image, outDataBuffer;
76 char programSrc[ 10240 ];
77
78 image_kernel_data outKernelData;
79
80 // Generate some data to test against
81 BufferOwningPtr<char> imageValues;
82 generate_random_image_data( imageInfo, imageValues, d );
83
84 // Construct testing source
85 if( gDebugTrace )
86 log_info( " - Creating image %d by %d...\n", (int)imageInfo->width, (int)imageInfo->height );
87
88 if( imageInfo->depth != 0 )
89 image = create_image_3d(context, flags, imageInfo->format,
90 imageInfo->width, imageInfo->height,
91 imageInfo->depth, 0, 0, NULL, &error);
92 else
93 image =
94 create_image_2d(context, flags, imageInfo->format, imageInfo->width,
95 imageInfo->height, 0, NULL, &error);
96 if( image == NULL )
97 {
98 log_error( "ERROR: Unable to create image of size %d x %d x %d (%s)", (int)imageInfo->width, (int)imageInfo->height, (int)imageInfo->depth, IGetErrorString( error ) );
99 return -1;
100 }
101
102 char channelTypeConstantString[256] = {0};
103 char channelOrderConstantString[256] = {0};
104
105 const char* channelTypeName = GetChannelTypeName( imageInfo->format->image_channel_data_type );
106 const char* channelOrderName = GetChannelOrderName( imageInfo->format->image_channel_order );
107 const char *image_access_qualifier =
108 (flags == CL_MEM_READ_ONLY) ? "read_only" : "write_only";
109 const char *cl_khr_3d_image_writes_enabler = "";
110 if ((flags != CL_MEM_READ_ONLY) && (imageInfo->depth != 0))
111 cl_khr_3d_image_writes_enabler =
112 "#pragma OPENCL EXTENSION cl_khr_3d_image_writes : enable";
113
114 if(channelTypeName && strlen(channelTypeName))
115 sprintf(channelTypeConstantString, "CLK_%s", &channelTypeName[3]); // replace CL_* with CLK_*
116
117 if(channelOrderName && strlen(channelOrderName))
118 sprintf(channelOrderConstantString, "CLK_%s", &channelOrderName[3]); // replace CL_* with CLK_*
119
120 // Create a program to run against
121 sprintf(programSrc, methodTestKernelPattern, cl_khr_3d_image_writes_enabler,
122 image_access_qualifier, (imageInfo->depth != 0) ? 3 : 2,
123 (imageInfo->format->image_channel_order == CL_DEPTH) ? "_depth"
124 : "",
125 (imageInfo->depth != 0) ? depthKernelLine : "",
126 (imageInfo->depth != 0) ? 4 : 2,
127 (imageInfo->depth != 0) ? depthDimKernelLine : "",
128 channelTypeConstantString, channelOrderConstantString);
129
130 //log_info("-----------------------------------\n%s\n", programSrc);
131 error = clFinish(queue);
132 if (error)
133 print_error(error, "clFinish failed.\n");
134 const char *ptr = programSrc;
135 error = create_single_kernel_helper(context, &program, &kernel, 1, &ptr,
136 "sample_kernel");
137 test_error( error, "Unable to create kernel to test against" );
138
139 // Create an output buffer
140 outDataBuffer = clCreateBuffer(context, CL_MEM_READ_WRITE,
141 sizeof(outKernelData), NULL, &error);
142 test_error( error, "Unable to create output buffer" );
143
144 // Set up arguments and run
145 error = clSetKernelArg( kernel, 0, sizeof( image ), &image );
146 test_error( error, "Unable to set kernel argument" );
147 error = clSetKernelArg( kernel, 1, sizeof( outDataBuffer ), &outDataBuffer );
148 test_error( error, "Unable to set kernel argument" );
149
150 size_t threads[1] = { 1 }, localThreads[1] = { 1 };
151
152 error = clEnqueueNDRangeKernel( queue, kernel, 1, NULL, threads, localThreads, 0, NULL, NULL );
153 test_error( error, "Unable to run kernel" );
154
155 error = clEnqueueReadBuffer( queue, outDataBuffer, CL_TRUE, 0, sizeof( outKernelData ), &outKernelData, 0, NULL, NULL );
156 test_error( error, "Unable to read data buffer" );
157
158
159 // Verify the results now
160 if( outKernelData.width != (cl_int)imageInfo->width )
161 {
162 log_error( "ERROR: Returned width did not validate (expected %d, got %d)\n", (int)imageInfo->width, (int)outKernelData.width );
163 error = -1;
164 }
165 if( outKernelData.height != (cl_int)imageInfo->height )
166 {
167 log_error( "ERROR: Returned height did not validate (expected %d, got %d)\n", (int)imageInfo->height, (int)outKernelData.height );
168 error = -1;
169 }
170 if( ( imageInfo->depth != 0 ) && ( outKernelData.depth != (cl_int)imageInfo->depth ) )
171 {
172 log_error( "ERROR: Returned depth did not validate (expected %d, got %d)\n", (int)imageInfo->depth, (int)outKernelData.depth );
173 error = -1;
174 }
175 if( outKernelData.widthDim != (cl_int)imageInfo->width )
176 {
177 log_error( "ERROR: Returned width from get_image_dim did not validate (expected %d, got %d)\n", (int)imageInfo->width, (int)outKernelData.widthDim );
178 error = -1;
179 }
180 if( outKernelData.heightDim != (cl_int)imageInfo->height )
181 {
182 log_error( "ERROR: Returned height from get_image_dim did not validate (expected %d, got %d)\n", (int)imageInfo->height, (int)outKernelData.heightDim );
183 error = -1;
184 }
185 if( ( imageInfo->depth != 0 ) && ( outKernelData.depthDim != (cl_int)imageInfo->depth ) )
186 {
187 log_error( "ERROR: Returned depth from get_image_dim did not validate (expected %d, got %d)\n", (int)imageInfo->depth, (int)outKernelData.depthDim );
188 error = -1;
189 }
190 if( outKernelData.channelType != (cl_int)outKernelData.expectedChannelType )
191 {
192 log_error( "ERROR: Returned channel type did not validate (expected %s (%d), got %d)\n", GetChannelTypeName( imageInfo->format->image_channel_data_type ),
193 (int)outKernelData.expectedChannelType, (int)outKernelData.channelType );
194 error = -1;
195 }
196 if( outKernelData.channelOrder != (cl_int)outKernelData.expectedChannelOrder )
197 {
198 log_error( "ERROR: Returned channel order did not validate (expected %s (%d), got %d)\n", GetChannelOrderName( imageInfo->format->image_channel_order ),
199 (int)outKernelData.expectedChannelOrder, (int)outKernelData.channelOrder );
200 error = -1;
201 }
202
203 if( clFinish(queue) != CL_SUCCESS )
204 {
205 log_error( "ERROR: CL Finished failed in %s \n", __FUNCTION__);
206 error = -1;
207 }
208
209 return error;
210 }
211
test_get_image_info_2D(cl_device_id device,cl_context context,cl_command_queue queue,cl_image_format * format,cl_mem_flags flags)212 int test_get_image_info_2D(cl_device_id device, cl_context context,
213 cl_command_queue queue, cl_image_format *format,
214 cl_mem_flags flags)
215 {
216 size_t maxWidth, maxHeight;
217 cl_ulong maxAllocSize, memSize;
218 image_descriptor imageInfo = { 0 };
219 RandomSeed seed( gRandomSeed );
220 size_t pixelSize;
221
222 imageInfo.type = CL_MEM_OBJECT_IMAGE2D;
223 imageInfo.format = format;
224 imageInfo.depth = imageInfo.slicePitch = 0;
225 pixelSize = get_pixel_size( imageInfo.format );
226
227 int error = clGetDeviceInfo( device, CL_DEVICE_IMAGE2D_MAX_WIDTH, sizeof( maxWidth ), &maxWidth, NULL );
228 error |= clGetDeviceInfo( device, CL_DEVICE_IMAGE2D_MAX_HEIGHT, sizeof( maxHeight ), &maxHeight, NULL );
229 error |= clGetDeviceInfo( device, CL_DEVICE_MAX_MEM_ALLOC_SIZE, sizeof( maxAllocSize ), &maxAllocSize, NULL );
230 error |= clGetDeviceInfo( device, CL_DEVICE_GLOBAL_MEM_SIZE, sizeof( memSize ), &memSize, NULL );
231 test_error( error, "Unable to get max image 2D size from device" );
232
233 if (memSize > (cl_ulong)SIZE_MAX) {
234 memSize = (cl_ulong)SIZE_MAX;
235 maxAllocSize = (cl_ulong)SIZE_MAX;
236 }
237
238 if( gTestSmallImages )
239 {
240 for( imageInfo.width = 1; imageInfo.width < 13; imageInfo.width++ )
241 {
242 imageInfo.rowPitch = imageInfo.width * pixelSize;
243 for( imageInfo.height = 1; imageInfo.height < 9; imageInfo.height++ )
244 {
245 if( gDebugTrace )
246 log_info( " at size %d,%d\n", (int)imageInfo.width, (int)imageInfo.height );
247
248 int ret = test_get_image_info_single(context, queue, &imageInfo,
249 seed, flags);
250 if( ret )
251 return -1;
252 }
253 }
254 }
255 else if( gTestMaxImages )
256 {
257 // Try a specific set of maximum sizes
258 size_t numbeOfSizes;
259 size_t sizes[100][3];
260
261 get_max_sizes(&numbeOfSizes, 100, sizes, maxWidth, maxHeight, 1, 1, maxAllocSize, memSize, CL_MEM_OBJECT_IMAGE2D, imageInfo.format);
262
263 for( size_t idx = 0; idx < numbeOfSizes; idx++ )
264 {
265 imageInfo.width = sizes[ idx ][ 0 ];
266 imageInfo.height = sizes[ idx ][ 1 ];
267 imageInfo.rowPitch = imageInfo.width * pixelSize;
268
269 log_info( "Testing %d x %d\n", (int)sizes[ idx ][ 0 ], (int)sizes[ idx ][ 1 ]);
270 if( gDebugTrace )
271 log_info( " at max size %d,%d\n", (int)sizes[ idx ][ 0 ], (int)sizes[ idx ][ 1 ] );
272 if (test_get_image_info_single(context, queue, &imageInfo, seed,
273 flags))
274 return -1;
275 }
276 }
277 else
278 {
279 for( int i = 0; i < NUM_IMAGE_ITERATIONS; i++ )
280 {
281 cl_ulong size;
282 // Loop until we get a size that a) will fit in the max alloc size and b) that an allocation of that
283 // image, the result array, plus offset arrays, will fit in the global ram space
284 do
285 {
286 imageInfo.width = (size_t)random_log_in_range( 16, (int)maxWidth / 32, seed );
287 imageInfo.height = (size_t)random_log_in_range( 16, (int)maxHeight / 32, seed );
288
289 imageInfo.rowPitch = imageInfo.width * pixelSize;
290 size_t extraWidth = (int)random_log_in_range( 0, 64, seed );
291 imageInfo.rowPitch += extraWidth;
292
293 do {
294 extraWidth++;
295 imageInfo.rowPitch += extraWidth;
296 } while ((imageInfo.rowPitch % pixelSize) != 0);
297
298 size = (cl_ulong)imageInfo.rowPitch * (cl_ulong)imageInfo.height * 4;
299 } while( size > maxAllocSize || ( size * 3 ) > memSize );
300
301 if( gDebugTrace )
302 log_info( " at size %d,%d (row pitch %d) out of %d,%d\n", (int)imageInfo.width, (int)imageInfo.height, (int)imageInfo.rowPitch, (int)maxWidth, (int)maxHeight );
303 int ret = test_get_image_info_single(context, queue, &imageInfo,
304 seed, flags);
305 if( ret )
306 return -1;
307 }
308 }
309
310 return 0;
311 }
312