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 "harness/typeWrappers.h"
18 #include "harness/testHarness.h"
19 #include "harness/conversions.h"
20
21 const char *test_kernels[] = { "__kernel void kernelA(__global uint *dst)\n"
22 "{\n"
23 "\n"
24 " dst[get_global_id(0)]*=3;\n"
25 "\n"
26 "}\n"
27 "__kernel void kernelB(__global uint *dst)\n"
28 "{\n"
29 "\n"
30 " dst[get_global_id(0)]++;\n"
31 "\n"
32 "}\n" };
33
34 #define TEST_SIZE 512
35 #define MAX_DEVICES 32
36 #define MAX_QUEUES 1000
37
test_device_set(size_t deviceCount,size_t queueCount,cl_device_id * devices,int num_elements)38 int test_device_set(size_t deviceCount, size_t queueCount, cl_device_id *devices, int num_elements)
39 {
40 int error;
41 clContextWrapper context;
42 clProgramWrapper program;
43 clKernelWrapper kernels[2];
44 clMemWrapper stream;
45 clCommandQueueWrapper queues[MAX_QUEUES];
46 size_t threads[1], localThreads[1];
47 cl_uint data[TEST_SIZE];
48 cl_uint outputData[TEST_SIZE];
49 cl_uint expectedResults[TEST_SIZE];
50 cl_uint expectedResultsOneDevice[MAX_DEVICES][TEST_SIZE];
51 size_t i;
52
53 memset(queues, 0, sizeof(queues));
54
55 RandomSeed seed( gRandomSeed );
56
57 if (deviceCount > MAX_DEVICES) {
58 log_error("Number of devices in set (%ld) is greater than the number for which the test was written (%d).", deviceCount, MAX_DEVICES);
59 return -1;
60 }
61
62 if (queueCount > MAX_QUEUES) {
63 log_error("Number of queues (%ld) is greater than the number for which the test was written (%d).", queueCount, MAX_QUEUES);
64 return -1;
65 }
66
67 log_info("Testing with %ld queues on %ld devices, %ld kernel executions.\n", queueCount, deviceCount, queueCount*num_elements/TEST_SIZE);
68
69 for (i=0; i<deviceCount; i++) {
70 char deviceName[4096] = "";
71 error = clGetDeviceInfo(devices[i], CL_DEVICE_NAME, sizeof(deviceName), deviceName, NULL);
72 test_error(error, "clGetDeviceInfo CL_DEVICE_NAME failed");
73 log_info("Device %ld is \"%s\".\n", i, deviceName);
74 }
75
76 /* Create a context */
77 context = clCreateContext( NULL, (cl_uint)deviceCount, devices, notify_callback, NULL, &error );
78 test_error( error, "Unable to create testing context" );
79
80 /* Create our kernels (they all have the same arguments so we don't need multiple ones for each device) */
81 if( create_single_kernel_helper( context, &program, &kernels[0], 1, test_kernels, "kernelA" ) != 0 )
82 {
83 return -1;
84 }
85
86 kernels[1] = clCreateKernel(program, "kernelB", &error);
87 test_error(error, "clCreateKernel failed");
88
89
90 /* Now create I/O streams */
91 for( i = 0; i < TEST_SIZE; i++ )
92 data[i] = genrand_int32(seed);
93
94 stream = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR,
95 sizeof(cl_uint) * TEST_SIZE, data, &error);
96 test_error(error, "Unable to create test array");
97
98 // Update the expected results
99 for( i = 0; i < TEST_SIZE; i++ ) {
100 expectedResults[i] = data[i];
101 for (size_t j=0; j<deviceCount; j++)
102 expectedResultsOneDevice[j][i] = data[i];
103 }
104
105
106 // Set the arguments
107 error = clSetKernelArg( kernels[0], 0, sizeof( stream ), &stream);
108 test_error( error, "Unable to set kernel arguments" );
109 error = clSetKernelArg( kernels[1], 0, sizeof( stream ), &stream);
110 test_error( error, "Unable to set kernel arguments" );
111
112 /* Run the test */
113 threads[0] = (size_t)TEST_SIZE;
114
115 error = get_max_common_work_group_size( context, kernels[0], threads[0], &localThreads[ 0 ] );
116 test_error( error, "Unable to calc work group size" );
117
118 /* Create work queues */
119 for( i = 0; i < queueCount; i++ )
120 {
121 queues[i] = clCreateCommandQueue( context, devices[ i % deviceCount ], 0, &error );
122 if (error != CL_SUCCESS || queues[i] == NULL) {
123 log_info("Could not create queue[%d].\n", (int)i);
124 queueCount = i;
125 break;
126 }
127 }
128 log_info("Testing with %d queues.\n", (int)queueCount);
129
130 /* Enqueue executions */
131 for( int z = 0; z<num_elements/TEST_SIZE; z++) {
132 for( i = 0; i < queueCount; i++ )
133 {
134 // Randomly choose a kernel to execute.
135 int kernel_selection = (int)get_random_float(0, 2, seed);
136 error = clEnqueueNDRangeKernel( queues[ i ], kernels[ kernel_selection ], 1, NULL, threads, localThreads, 0, NULL, NULL );
137 test_error( error, "Kernel execution failed" );
138
139 // Update the expected results
140 for( int j = 0; j < TEST_SIZE; j++ ) {
141 expectedResults[j] = (kernel_selection) ? expectedResults[j]+1 : expectedResults[j]*3;
142 expectedResultsOneDevice[i % deviceCount][j] = (kernel_selection) ? expectedResultsOneDevice[i % deviceCount][j]+1 : expectedResultsOneDevice[i % deviceCount][j]*3;
143 }
144
145 // Force the queue to finish so the next one will be in sync
146 error = clFinish(queues[i]);
147 test_error( error, "clFinish failed");
148 }
149 }
150
151 /* Read results */
152 int errors = 0;
153 for (int q = 0; q<(int)queueCount; q++) {
154 error = clEnqueueReadBuffer( queues[ 0 ], stream, CL_TRUE, 0, sizeof(cl_int)*TEST_SIZE, (char *)outputData, 0, NULL, NULL );
155 test_error( error, "Unable to get result data set" );
156
157 int errorsThisTime = 0;
158 /* Verify all of the data now */
159 for( i = 0; i < TEST_SIZE; i++ )
160 {
161 if( expectedResults[ i ] != outputData[ i ] )
162 {
163 log_error( "ERROR: Sample data did not verify for queue %d on device %ld (sample %d, expected %d, got %d)\n",
164 q, q % deviceCount, (int)i, expectedResults[ i ], outputData[ i ] );
165 for (size_t j=0; j<deviceCount; j++) {
166 if (expectedResultsOneDevice[j][i] == outputData[i])
167 log_info("Sample consistent with only device %ld having modified the data.\n", j);
168 }
169 errorsThisTime++;
170 break;
171 }
172 }
173 if (errorsThisTime)
174 errors++;
175 }
176
177 /* All done now! */
178 if (errors) return -1;
179 return 0;
180 }
181
test_two_devices(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)182 int test_two_devices(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
183 {
184 cl_platform_id platform;
185 cl_device_id devices[2];
186 int err;
187 cl_uint numDevices;
188
189 err = clGetPlatformIDs(1, &platform, NULL);
190 test_error( err, "Unable to get platform" );
191
192 /* Get some devices */
193 err = clGetDeviceIDs(platform, CL_DEVICE_TYPE_ALL, 2, devices, &numDevices );
194 test_error( err, "Unable to get 2 devices" );
195
196 if( numDevices < 2 )
197 {
198 log_info( "WARNING: two device test unable to get two devices via CL_DEVICE_TYPE_ALL (got %d devices). Skipping test...\n", (int)numDevices );
199 return 0;
200 }
201 else if (numDevices > 2)
202 {
203 log_info("Note: got %d devices, using just the first two.\n", (int)numDevices);
204 }
205
206 /* Run test */
207 return test_device_set( 2, 2, devices, num_elements );
208 }
209
test_max_devices(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)210 int test_max_devices(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
211 {
212 cl_platform_id platform;
213 cl_device_id devices[MAX_DEVICES];
214 cl_uint deviceCount;
215 int err;
216
217 err = clGetPlatformIDs(1, &platform, NULL);
218 test_error( err, "Unable to get platform" );
219
220 /* Get some devices */
221 err = clGetDeviceIDs(platform, CL_DEVICE_TYPE_ALL, MAX_DEVICES, devices, &deviceCount );
222 test_error( err, "Unable to get multiple devices" );
223
224 log_info("Testing with %d devices.", deviceCount);
225
226 /* Run test */
227 return test_device_set( deviceCount, deviceCount, devices, num_elements );
228 }
229
test_hundred_queues(cl_device_id device,cl_context contextIgnore,cl_command_queue queueIgnore,int num_elements)230 int test_hundred_queues(cl_device_id device, cl_context contextIgnore, cl_command_queue queueIgnore, int num_elements)
231 {
232 return test_device_set( 1, 100, &device, num_elements );
233 }
234
235