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 #include "gl_headers.h"
19
20 static const char *imageReadKernelPattern =
21 "__kernel void sample_test( read_only image3d_t source, sampler_t sampler, __global %s4 *results )\n"
22 "{\n"
23 " int tidX = get_global_id(0);\n"
24 " int tidY = get_global_id(1);\n"
25 " int tidZ = get_global_id(2);\n"
26 " int width = get_image_width( source );\n"
27 " int height = get_image_height( source );\n"
28 " int offset = tidZ * width * height + tidY * width + tidX;\n"
29 "\n"
30 " results[ offset ] = read_image%s( source, sampler, (int4)( tidX, tidY, tidZ, 0 ) );\n"
31 "}\n";
32
test_image_read(cl_context context,cl_command_queue queue,GLenum glTarget,GLuint glTexture,size_t imageWidth,size_t imageHeight,size_t imageDepth,cl_image_format * outFormat,ExplicitType * outType,void ** outResultBuffer)33 static int test_image_read( cl_context context, cl_command_queue queue, GLenum glTarget, GLuint glTexture,
34 size_t imageWidth, size_t imageHeight, size_t imageDepth,
35 cl_image_format *outFormat, ExplicitType *outType, void **outResultBuffer )
36 {
37 clProgramWrapper program;
38 clKernelWrapper kernel;
39 clMemWrapper streams[ 2 ];
40
41 int error;
42 size_t threads[ 3 ], localThreads[ 3 ];
43 char kernelSource[1024];
44 char *programPtr;
45
46
47 // Create a CL image from the supplied GL texture
48 streams[ 0 ] = (*clCreateFromGLTexture_ptr)( context, CL_MEM_READ_ONLY, glTarget, 0, glTexture, &error );
49 if( error != CL_SUCCESS )
50 {
51 print_error( error, "Unable to create CL image from GL texture" );
52 #ifndef GL_ES_VERSION_2_0
53 GLint fmt;
54 glGetTexLevelParameteriv( glTarget, 0, GL_TEXTURE_INTERNAL_FORMAT, &fmt );
55 log_error( " Supplied GL texture was format %s\n", GetGLFormatName( fmt ) );
56 #endif
57 return error;
58 }
59
60 // Determine data type and format that CL came up with
61 error = clGetImageInfo( streams[ 0 ], CL_IMAGE_FORMAT, sizeof( cl_image_format ), outFormat, NULL );
62 test_error( error, "Unable to get CL image format" );
63
64 /* Create the source */
65 *outType = get_read_kernel_type( outFormat );
66 size_t channelSize = get_explicit_type_size( *outType );
67
68 sprintf( kernelSource, imageReadKernelPattern, get_explicit_type_name( *outType ), get_kernel_suffix( outFormat ) );
69
70 /* Create kernel */
71 programPtr = kernelSource;
72 if( create_single_kernel_helper( context, &program, &kernel, 1, (const char **)&programPtr, "sample_test" ) )
73 {
74 return -1;
75 }
76
77
78 // Create a vanilla output buffer
79 streams[ 1 ] = clCreateBuffer( context, CL_MEM_READ_WRITE, channelSize * 4 * imageWidth * imageHeight * imageDepth, NULL, &error );
80 test_error( error, "Unable to create output buffer" );
81
82
83 /* Assign streams and execute */
84 clSamplerWrapper sampler = clCreateSampler( context, CL_FALSE, CL_ADDRESS_NONE, CL_FILTER_NEAREST, &error );
85 test_error( error, "Unable to create sampler" );
86
87 error = clSetKernelArg( kernel, 0, sizeof( streams[ 0 ] ), &streams[ 0 ] );
88 test_error( error, "Unable to set kernel arguments" );
89 error = clSetKernelArg( kernel, 1, sizeof( sampler ), &sampler );
90 test_error( error, "Unable to set kernel arguments" );
91 error = clSetKernelArg( kernel, 2, sizeof( streams[ 1 ] ), &streams[ 1 ] );
92 test_error( error, "Unable to set kernel arguments" );
93
94 glFlush();
95
96 error = (*clEnqueueAcquireGLObjects_ptr)( queue, 1, &streams[ 0 ], 0, NULL, NULL);
97 test_error( error, "Unable to acquire GL obejcts");
98
99 /* Run the kernel */
100 threads[ 0 ] = imageWidth;
101 threads[ 1 ] = imageHeight;
102 threads[ 2 ] = imageDepth;
103
104 error = get_max_common_3D_work_group_size( context, kernel, threads, localThreads );
105 test_error( error, "Unable to get work group size to use" );
106
107 error = clEnqueueNDRangeKernel( queue, kernel, 3, NULL, threads, localThreads, 0, NULL, NULL );
108 test_error( error, "Unable to execute test kernel" );
109
110
111 error = (*clEnqueueReleaseGLObjects_ptr)( queue, 1, &streams[ 0 ], 0, NULL, NULL );
112 test_error(error, "clEnqueueReleaseGLObjects failed");
113
114 // Read results from the CL buffer
115 *outResultBuffer = (void *)( new char[ channelSize * 4 * imageWidth * imageHeight * imageDepth ] );
116 error = clEnqueueReadBuffer( queue, streams[ 1 ], CL_TRUE, 0, channelSize * 4 * imageWidth * imageHeight * imageDepth,
117 *outResultBuffer, 0, NULL, NULL );
118 test_error( error, "Unable to read output CL buffer!" );
119
120 return 0;
121 }
122
test_image_format_read(cl_context context,cl_command_queue queue,size_t width,size_t height,size_t depth,GLenum target,GLenum format,GLenum internalFormat,GLenum glType,ExplicitType type,MTdata d)123 int test_image_format_read( cl_context context, cl_command_queue queue,
124 size_t width, size_t height, size_t depth,
125 GLenum target, GLenum format, GLenum internalFormat,
126 GLenum glType, ExplicitType type, MTdata d )
127 {
128 int error;
129
130
131 // Create the GL texture
132 glTextureWrapper glTexture;
133 void* tmp = CreateGLTexture3D( width, height, depth, target, format, internalFormat, glType, type, &glTexture, &error, d );
134 BufferOwningPtr<char> inputBuffer(tmp);
135 if( error != 0 )
136 {
137 return error;
138 }
139
140 /* skip formats not supported by OpenGL */
141 if(!tmp)
142 {
143 return 0;
144 }
145
146 // Run and get the results
147 cl_image_format clFormat;
148 ExplicitType actualType;
149 char *outBuffer;
150 error = test_image_read( context, queue, target, glTexture, width, height, depth, &clFormat, &actualType, (void **)&outBuffer );
151 if( error != 0 )
152 return error;
153 BufferOwningPtr<char> actualResults(outBuffer);
154
155 log_info( "- Read [%4d x %4d x %4d] : GL Texture : %s : %s : %s => CL Image : %s : %s \n",
156 (int)width, (int)height, (int)depth,
157 GetGLFormatName( format ), GetGLFormatName( internalFormat ), GetGLTypeName( glType),
158 GetChannelOrderName( clFormat.image_channel_order ), GetChannelTypeName( clFormat.image_channel_data_type ));
159
160 // We have to convert our input buffer to the returned type, so we can validate.
161 // This is necessary because OpenCL might not actually pick an internal format that actually matches our
162 // input format (for example, if it picks a normalized format, the results will come out as floats instead of
163 // going in as ints).
164
165 BufferOwningPtr<char> convertedInputs(convert_to_expected( inputBuffer, width * height * depth, type, actualType ));
166 if( convertedInputs == NULL )
167 return -1;
168
169 // Now we validate
170 if( actualType == kFloat )
171 return validate_float_results( convertedInputs, actualResults, width, height, depth );
172 else
173 return validate_integer_results( convertedInputs, actualResults, width, height, depth, get_explicit_type_size( actualType ) );
174 }
175
176
test_images_read_3D(cl_device_id device,cl_context context,cl_command_queue queue,int numElements)177 int test_images_read_3D( cl_device_id device, cl_context context, cl_command_queue queue, int numElements )
178 {
179 GLenum targets[] = { GL_TEXTURE_3D };
180
181 struct {
182 GLenum internal;
183 GLenum format;
184 GLenum datatype;
185 ExplicitType type;
186
187 } formats[] = {
188 #ifdef GL_ES_VERSION_2_0
189 { GL_RGBA, GL_RGBA, GL_UNSIGNED_BYTE, kUChar },
190 { GL_RGBA, GL_RGBA, GL_UNSIGNED_SHORT, kUShort },
191 // XXX add others
192 #else // GL_ES_VERSION_2_0
193 { GL_RGBA, GL_BGRA, GL_UNSIGNED_INT_8_8_8_8_REV, kUChar },
194 { GL_RGBA, GL_RGBA, GL_UNSIGNED_INT_8_8_8_8_REV, kUChar },
195 { GL_RGBA8, GL_RGBA, GL_UNSIGNED_BYTE, kUChar },
196 { GL_RGBA16, GL_RGBA, GL_UNSIGNED_SHORT, kUShort },
197 { GL_RGBA8I_EXT, GL_RGBA_INTEGER_EXT, GL_BYTE, kChar },
198 { GL_RGBA16I_EXT, GL_RGBA_INTEGER_EXT, GL_SHORT, kShort },
199 { GL_RGBA32I_EXT, GL_RGBA_INTEGER_EXT, GL_INT, kInt },
200 { GL_RGBA8UI_EXT, GL_RGBA_INTEGER_EXT, GL_UNSIGNED_BYTE, kUChar },
201 { GL_RGBA16UI_EXT, GL_RGBA_INTEGER_EXT, GL_UNSIGNED_SHORT, kUShort },
202 { GL_RGBA32UI_EXT, GL_RGBA_INTEGER_EXT, GL_UNSIGNED_INT, kUInt },
203 { GL_RGBA32F_ARB, GL_RGBA, GL_FLOAT, kFloat }
204 #endif
205 };
206
207 size_t sizes[] = { 2, 4, 8, 16, 32, 64, 128 };
208 size_t fmtIdx, tgtIdx;
209 int error = 0;
210 RandomSeed seed(gRandomSeed);
211
212 size_t iter = sizeof(sizes)/sizeof(sizes[0]);
213
214 // Check if images are supported
215 if (checkForImageSupport(device)) {
216 log_info("Device does not support images. Skipping test.\n");
217 return 0;
218 }
219
220 // Loop through a set of GL formats, testing a set of sizes against each one
221 for( fmtIdx = 0; fmtIdx < sizeof( formats ) / sizeof( formats[ 0 ] ); fmtIdx++ )
222 {
223 for( tgtIdx = 0; tgtIdx < sizeof( targets ) / sizeof( targets[ 0 ] ); tgtIdx++ )
224 {
225 size_t i;
226
227 log_info( "Testing image read for GL format %s : %s : %s : %s\n",
228 GetGLTargetName( targets[ tgtIdx ] ),
229 GetGLFormatName( formats[ fmtIdx ].internal ),
230 GetGLBaseFormatName( formats[ fmtIdx ].format ),
231 GetGLTypeName( formats[ fmtIdx ].datatype ) );
232
233 for( i = 0; i < iter; i++ )
234 {
235 if( test_image_format_read( context, queue, sizes[i], sizes[i], sizes[i],
236 targets[ tgtIdx ],
237 formats[ fmtIdx ].format,
238 formats[ fmtIdx ].internal,
239 formats[ fmtIdx ].datatype,
240 formats[ fmtIdx ].type, seed ) )
241 {
242 log_error( "ERROR: Image read test failed for %s : %s : %s : %s\n\n",
243 GetGLTargetName( targets[ tgtIdx ] ),
244 GetGLFormatName( formats[ fmtIdx ].internal ),
245 GetGLBaseFormatName( formats[ fmtIdx ].format ),
246 GetGLTypeName( formats[ fmtIdx ].datatype ) );
247
248 error++;
249 break; // Skip other sizes for this combination
250 }
251 }
252 if( i == sizeof (sizes) / sizeof( sizes[0] ) )
253 {
254 log_info( "passed: Image read test for GL format %s : %s : %s : %s\n\n",
255 GetGLTargetName( targets[ tgtIdx ] ),
256 GetGLFormatName( formats[ fmtIdx ].internal ),
257 GetGLBaseFormatName( formats[ fmtIdx ].format ),
258 GetGLTypeName( formats[ fmtIdx ].datatype ) );
259
260 }
261 }
262 }
263
264 return error;
265 }
266
267