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 static const char *sample_binary_kernel_source[] = {
19 "__kernel void sample_test(__global float *src, __global int *dst)\n"
20 "{\n"
21 " int tid = get_global_id(0);\n"
22 "\n"
23 " dst[tid] = (int)src[tid] + 1;\n"
24 "\n"
25 "}\n" };
26
27
test_binary_get(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)28 int test_binary_get(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
29 {
30 int error;
31 clProgramWrapper program;
32 size_t binarySize;
33
34
35 error = create_single_kernel_helper(context, &program, NULL, 1, sample_binary_kernel_source, NULL);
36 test_error( error, "Unable to build test program" );
37
38 // Get the size of the resulting binary (only one device)
39 error = clGetProgramInfo( program, CL_PROGRAM_BINARY_SIZES, sizeof( binarySize ), &binarySize, NULL );
40 test_error( error, "Unable to get binary size" );
41
42 // Sanity check
43 if( binarySize == 0 )
44 {
45 log_error( "ERROR: Binary size of program is zero\n" );
46 return -1;
47 }
48
49 // Create a buffer and get the actual binary
50 unsigned char *binary;
51 binary = (unsigned char*)malloc(sizeof(unsigned char)*binarySize);
52 unsigned char *buffers[ 1 ] = { binary };
53
54 // Do another sanity check here first
55 size_t size;
56 error = clGetProgramInfo( program, CL_PROGRAM_BINARIES, 0, NULL, &size );
57 test_error( error, "Unable to get expected size of binaries array" );
58 if( size != sizeof( buffers ) )
59 {
60 log_error( "ERROR: Expected size of binaries array in clGetProgramInfo is incorrect (should be %d, got %d)\n", (int)sizeof( buffers ), (int)size );
61 free(binary);
62 return -1;
63 }
64
65 error = clGetProgramInfo( program, CL_PROGRAM_BINARIES, sizeof( buffers ), &buffers, NULL );
66 test_error( error, "Unable to get program binary" );
67
68 // No way to verify the binary is correct, so just be good with that
69 free(binary);
70 return 0;
71 }
72
73
test_binary_create(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)74 int test_binary_create(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
75 {
76 /* To test this in a self-contained fashion, we have to create a program with
77 source, then get the binary, then use that binary to reload the program, and then verify */
78
79 int error;
80 clProgramWrapper program, program_from_binary;
81 size_t binarySize;
82
83
84 error = create_single_kernel_helper(context, &program, NULL, 1, sample_binary_kernel_source, NULL);
85 test_error( error, "Unable to build test program" );
86
87 // Get the size of the resulting binary (only one device)
88 error = clGetProgramInfo( program, CL_PROGRAM_BINARY_SIZES, sizeof( binarySize ), &binarySize, NULL );
89 test_error( error, "Unable to get binary size" );
90
91 // Sanity check
92 if( binarySize == 0 )
93 {
94 log_error( "ERROR: Binary size of program is zero\n" );
95 return -1;
96 }
97
98 // Create a buffer and get the actual binary
99 unsigned char *binary = (unsigned char*)malloc(binarySize);
100 const unsigned char *buffers[ 1 ] = { binary };
101
102 error = clGetProgramInfo( program, CL_PROGRAM_BINARIES, sizeof( buffers ), &buffers, NULL );
103 test_error( error, "Unable to get program binary" );
104
105 cl_int loadErrors[ 1 ];
106 program_from_binary = clCreateProgramWithBinary( context, 1, &deviceID, &binarySize, buffers, loadErrors, &error );
107 test_error( error, "Unable to load valid program binary" );
108 test_error( loadErrors[ 0 ], "Unable to load valid device binary into program" );
109
110 error = clBuildProgram( program_from_binary, 1, &deviceID, NULL, NULL, NULL );
111 test_error( error, "Unable to build binary program" );
112
113 // Get the size of the binary built from the first binary
114 size_t binary2Size;
115 error = clGetProgramInfo( program_from_binary, CL_PROGRAM_BINARY_SIZES, sizeof( binary2Size ), &binary2Size, NULL );
116 test_error( error, "Unable to get size for the binary program" );
117
118 // Now get the binary one more time and verify it loaded the right binary
119 unsigned char *binary2 = (unsigned char*)malloc(binary2Size);
120 buffers[ 0 ] = binary2;
121 error = clGetProgramInfo( program_from_binary, CL_PROGRAM_BINARIES, sizeof( buffers ), &buffers, NULL );
122 test_error( error, "Unable to get program binary second time" );
123
124 // Try again, this time without passing the status ptr in, to make sure we still
125 // get a valid binary
126 clProgramWrapper programWithoutStatus = clCreateProgramWithBinary( context, 1, &deviceID, &binary2Size, buffers, NULL, &error );
127 test_error( error, "Unable to load valid program binary when binary_status pointer is NULL" );
128
129 error = clBuildProgram( programWithoutStatus, 1, &deviceID, NULL, NULL, NULL );
130 test_error( error, "Unable to build binary program created without binary_status" );
131
132 // Get the size of the binary created without passing binary_status
133 size_t binary3Size;
134 error = clGetProgramInfo( programWithoutStatus, CL_PROGRAM_BINARY_SIZES, sizeof( binary3Size ), &binary3Size, NULL );
135 test_error( error, "Unable to get size for the binary program created without binary_status" );
136
137 // Now get the binary one more time
138 unsigned char *binary3 = (unsigned char*)malloc(binary3Size);
139 buffers[ 0 ] = binary3;
140 error = clGetProgramInfo( programWithoutStatus, CL_PROGRAM_BINARIES, sizeof( buffers ), &buffers, NULL );
141 test_error( error, "Unable to get program binary from the program created without binary_status" );
142
143 // We no longer need these intermediate binaries
144 free(binary);
145 free(binary2);
146 free(binary3);
147
148 // Now execute them both to see that they both do the same thing.
149 clMemWrapper in, out, out_binary;
150 clKernelWrapper kernel, kernel_binary;
151 cl_int *out_data, *out_data_binary;
152 cl_float *in_data;
153 size_t size_to_run = 1000;
154
155 // Allocate some data
156 in_data = (cl_float*)malloc(sizeof(cl_float)*size_to_run);
157 out_data = (cl_int*)malloc(sizeof(cl_int)*size_to_run);
158 out_data_binary = (cl_int*)malloc(sizeof(cl_int)*size_to_run);
159 memset(out_data, 0, sizeof(cl_int)*size_to_run);
160 memset(out_data_binary, 0, sizeof(cl_int)*size_to_run);
161 for (size_t i=0; i<size_to_run; i++)
162 in_data[i] = (cl_float)i;
163
164 // Create the buffers
165 in = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR, sizeof(cl_float)*size_to_run, in_data, &error);
166 test_error( error, "clCreateBuffer failed");
167 out = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR, sizeof(cl_int)*size_to_run, out_data, &error);
168 test_error( error, "clCreateBuffer failed");
169 out_binary = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR, sizeof(cl_int)*size_to_run, out_data_binary, &error);
170 test_error( error, "clCreateBuffer failed");
171
172 // Create the kernels
173 kernel = clCreateKernel(program, "sample_test", &error);
174 test_error( error, "clCreateKernel failed");
175 kernel_binary = clCreateKernel(program_from_binary, "sample_test", &error);
176 test_error( error, "clCreateKernel from binary failed");
177
178 // Set the arguments
179 error = clSetKernelArg(kernel, 0, sizeof(in), &in);
180 test_error( error, "clSetKernelArg failed");
181 error = clSetKernelArg(kernel, 1, sizeof(out), &out);
182 test_error( error, "clSetKernelArg failed");
183 error = clSetKernelArg(kernel_binary, 0, sizeof(in), &in);
184 test_error( error, "clSetKernelArg failed");
185 error = clSetKernelArg(kernel_binary, 1, sizeof(out_binary), &out_binary);
186 test_error( error, "clSetKernelArg failed");
187
188 // Execute the kernels
189 error = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &size_to_run, NULL, 0, NULL, NULL);
190 test_error( error, "clEnqueueNDRangeKernel failed");
191 error = clEnqueueNDRangeKernel(queue, kernel_binary, 1, NULL, &size_to_run, NULL, 0, NULL, NULL);
192 test_error( error, "clEnqueueNDRangeKernel for binary kernel failed");
193
194 // Finish up
195 error = clFinish(queue);
196 test_error( error, "clFinish failed");
197
198 // Get the results back
199 error = clEnqueueReadBuffer(queue, out, CL_TRUE, 0, sizeof(cl_int)*size_to_run, out_data, 0, NULL, NULL);
200 test_error( error, "clEnqueueReadBuffer failed");
201 error = clEnqueueReadBuffer(queue, out_binary, CL_TRUE, 0, sizeof(cl_int)*size_to_run, out_data_binary, 0, NULL, NULL);
202 test_error( error, "clEnqueueReadBuffer failed");
203
204 // Compare the results
205 if( memcmp( out_data, out_data_binary, sizeof(cl_int)*size_to_run ) != 0 )
206 {
207 log_error( "ERROR: Results from executing binary and regular kernel differ.\n" );
208 return -1;
209 }
210
211 // All done!
212 free(in_data);
213 free(out_data);
214 free(out_data_binary);
215 return 0;
216 }
217
218
219