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