1 /*
2 * Copyright (c) 2024, Oracle and/or its affiliates. All rights reserved.
3 * DO NOT ALTER OR REMOVE COPYRIGHT NOTICES OR THIS FILE HEADER.
4 *
5 * This code is free software; you can redistribute it and/or modify it
6 * under the terms of the GNU General Public License version 2 only, as
7 * published by the Free Software Foundation. Oracle designates this
8 * particular file as subject to the "Classpath" exception as provided
9 * by Oracle in the LICENSE file that accompanied this code.
10 *
11 * This code is distributed in the hope that it will be useful, but WITHOUT
12 * ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
13 * FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License
14 * version 2 for more details (a copy is included in the LICENSE file that
15 * accompanied this code).
16 *
17 * You should have received a copy of the GNU General Public License version
18 * 2 along with this work; if not, write to the Free Software Foundation,
19 * Inc., 51 Franklin St, Fifth Floor, Boston, MA 02110-1301 USA.
20 *
21 * Please contact Oracle, 500 Oracle Parkway, Redwood Shores, CA 94065 USA
22 * or visit www.oracle.com if you need additional information or have any
23 * questions.
24 */
25 #include "opencl_backend.h"
26
27 /*
28 While based on OpenCL's event list, I think we need to use a MOD eventMax queue.
29
30 So
31 */
32 OpenCLBackend::OpenCLQueue::OpenCLQueue(Backend *backend)
33 : ProfilableQueue(backend, 10000),
34 command_queue(),
35 events(new cl_event[eventMax]) {
36 }
37
38 cl_event *OpenCLBackend::OpenCLQueue::eventListPtr() const {
39 return (eventc == 0) ? nullptr : events;
40 }
41
42 cl_event *OpenCLBackend::OpenCLQueue::nextEventPtr() const {
43 return &events[eventc];
44 }
45
46 void OpenCLBackend::OpenCLQueue::showEvents(const int width) {
47 constexpr int SAMPLE_TYPES = 4;
48 auto *samples = new cl_ulong[SAMPLE_TYPES * eventc]; // queued, submit, start, end, complete
49 int sample = 0;
50 cl_ulong min = CL_LONG_MAX;
51 cl_ulong max = CL_LONG_MIN;
52
53 for (int event = 0; event < eventc; event++) {
54 for (int type = 0; type < SAMPLE_TYPES; type++) {
55 cl_profiling_info profiling_info_arr[] = {
56 CL_PROFILING_COMMAND_QUEUED,CL_PROFILING_COMMAND_SUBMIT,CL_PROFILING_COMMAND_START,
57 CL_PROFILING_COMMAND_END
58 };
59 if ((clGetEventProfilingInfo(events[event], profiling_info_arr[type], sizeof(samples[sample]),
60 &samples[sample], NULL)) !=
61 CL_SUCCESS) {
62 const char *profiling_info_name_arr[] = {
63 "CL_PROFILING_COMMAND_QUEUED", "CL_PROFILING_COMMAND_SUBMIT", "CL_PROFILING_COMMAND_START",
64 "CL_PROFILING_COMMAND_END"
65 };
66 std::cerr << "failed to get profile info " << profiling_info_name_arr[type] << std::endl;
67 }
68 if (sample == 0) {
69 if (type == 0) {
70 min = max = samples[sample];
71 }
72 } else {
73 if (samples[sample] < min) {
74 min = samples[sample];
75 }
76 if (samples[sample] > max) {
77 max = samples[sample];
78 }
79 }
80 sample++;
81 }
82 }
83 sample = 0;
84 const cl_ulong range = (max - min);
85 const cl_ulong scale = range / width; // range per char
86 std::cout << "Range: " << min << "-" << max << "(" << range << "ns)"
87 << " (" << scale << "ns) per char"
88 << " +:submitted, .:started, =:end " << std::endl;
89
90 for (int event = 0; event < eventc; event++) {
91 /* cl_command_type command_type;
92 clGetEventInfo(events[event],CL_EVENT_COMMAND_TYPE,sizeof(command_type), &command_type, nullptr);
93 switch (command_type){
94 case CL_COMMAND_MARKER: std::cout << "marker "; break;
95 case CL_COMMAND_USER: std::cout << " user "; break;
96 case CL_COMMAND_NDRANGE_KERNEL: std::cout << "kernel "; break;
97 case CL_COMMAND_READ_BUFFER: std::cout << " read "; break;
98 case CL_COMMAND_WRITE_BUFFER: std::cout << " write "; break;
99 default: std::cout << " other "; break;
100 } */
101 const int bits = eventInfoBits[event];
102 if ((bits & CopyToDeviceBits) == CopyToDeviceBits) {
103 std::cout << " write " << (bits & 0xffff) << " ";
104 }
105 if ((bits & CopyFromDeviceBits) == CopyFromDeviceBits) {
106 std::cout << " read " << (bits & 0xffff) << " ";
107 }
108 if ((bits & StartComputeBits) == StartComputeBits) {
109 std::cout << " start ";
110 }
111 if ((bits & EndComputeBits) == EndComputeBits) {
112 std::cout << " end ";
113 }
114 if ((bits & NDRangeBits) == NDRangeBits) {
115 std::cout << " kernel ";
116 }
117 if ((bits & EnterKernelDispatchBits) == EnterKernelDispatchBits) {
118 if ((bits & HasConstCharPtrArgBits) == HasConstCharPtrArgBits) {
119 std::cout << eventInfoConstCharPtrArgs[event] << std::endl;
120 }
121 std::cout << " enter{ ";
122 }
123 if ((bits & LeaveKernelDispatchBits) == LeaveKernelDispatchBits) {
124 // std::cout << " leave ";
125 if ((bits & HasConstCharPtrArgBits) == HasConstCharPtrArgBits) {
126 std::cout << eventInfoConstCharPtrArgs[event] << std::endl;
127 }
128 std::cout << " }leave ";
129 }
130
131
132 const cl_ulong queue = (samples[sample++] - min) / scale;
133 const cl_ulong submit = (samples[sample++] - min) / scale;
134 const cl_ulong start = (samples[sample++] - min) / scale;
135 const cl_ulong end = (samples[sample++] - min) / scale;
136
137 std::cout << std::setw(20) << (queue - end) << "(ns) ";
138 for (int c = 0; c < width; c++) {
139 char ch = ' ';
140 if (c >= queue && c <= submit) {
141 ch = '+';
142 } else if (c > submit && c < start) {
143 ch = '.';
144 } else if (c >= start && c < end) {
145 ch = '=';
146 }
147 std::cout << ch;
148 }
149 std::cout << std::endl;
150 }
151 delete[] samples;
152 }
153
154 void OpenCLBackend::OpenCLQueue::wait() {
155 auto backendInstance = dynamic_cast<OpenCLBackend *>(this->backend);
156 if ((isProfileOrOutOfOrder() || backendInstance -> isPlatformOnProfile()) && (eventc > 0)) {
157 OPENCL_CHECK(clWaitForEvents(eventc, events), "clWaitForEvents");
158 } else {
159 OPENCL_CHECK(clFinish(command_queue), "clFinish");
160 }
161 }
162
163 void OpenCLBackend::OpenCLQueue::marker(int bits) {
164 auto backendInstance = dynamic_cast<OpenCLBackend *>(this->backend);
165 if (isProfileOrOutOfOrder() || backendInstance -> isPlatformOnProfile()) {
166 OPENCL_CHECK(clEnqueueMarkerWithWaitList(
167 command_queue,
168 this->eventc,
169 this->eventListPtr(),
170 this->nextEventPtr()
171 ), "clEnqueueMarkerWithWaitList");
172 }
173 inc(bits);
174 }
175
176 void OpenCLBackend::OpenCLQueue::marker(int bits, const char *arg) {
177 auto backendInstance = dynamic_cast<OpenCLBackend *>(this->backend);
178 if (isProfileOrOutOfOrder() || backendInstance -> isPlatformOnProfile()) {
179 OPENCL_CHECK(clEnqueueMarkerWithWaitList(
180 command_queue,
181 this->eventc,
182 this->eventListPtr(),
183 this->nextEventPtr()),
184 "clEnqueueMarkerWithWaitList");
185 }
186 inc(bits, arg);
187 }
188
189 void OpenCLBackend::OpenCLQueue::computeStart() {
190 if (backend -> config -> alwaysCopy) {
191 wait(); // should be no-op
192 release(); // also ;
193 }
194 marker(StartComputeBits);
195 }
196
197 void OpenCLBackend::OpenCLQueue::computeEnd() {
198 marker(EndComputeBits);
199 }
200
201 void OpenCLBackend::OpenCLQueue::inc(const int bits) {
202 if (eventc + 1 >= eventMax) {
203 std::cerr << "OpenCLBackend::OpenCLQueue event list overflowed!!" << std::endl;
204 } else {
205 eventInfoBits[eventc] = bits;
206 }
207 eventc++;
208 }
209
210 void OpenCLBackend::OpenCLQueue::inc(const int bits, const char *arg) {
211 if (eventc + 1 >= eventMax) {
212 std::cerr << "OpenCLBackend::OpenCLQueue event list overflowed!!" << std::endl;
213 } else {
214 eventInfoBits[eventc] = bits | HasConstCharPtrArgBits;
215 eventInfoConstCharPtrArgs[eventc] = arg;
216 }
217 eventc++;
218 }
219
220 void OpenCLBackend::OpenCLQueue::markAsEndComputeAndInc() {
221 inc(EndComputeBits);
222 }
223
224 void OpenCLBackend::OpenCLQueue::markAsStartComputeAndInc() {
225 inc(StartComputeBits);
226 }
227
228 void OpenCLBackend::OpenCLQueue::markAsEnterKernelDispatchAndInc() {
229 inc(EnterKernelDispatchBits);
230 }
231
232 void OpenCLBackend::OpenCLQueue::markAsLeaveKernelDispatchAndInc() {
233 inc(LeaveKernelDispatchBits);
234 }
235
236 void OpenCLBackend::OpenCLQueue::release() {
237 // TODO: possible check ALL events before return from the macro
238 auto backendInstance = dynamic_cast<OpenCLBackend *>(this->backend);
239 if (isProfileOrOutOfOrder() || backendInstance -> isPlatformOnProfile()) {
240 for (int i = 0; i < eventc; i++) {
241 OPENCL_CHECK(clReleaseEvent(events[i]), "clReleaseEvent");
242 }
243 }
244 eventc = 0;
245 }
246
247 OpenCLBackend::OpenCLQueue::~OpenCLQueue() {
248 OPENCL_CHECK(clReleaseCommandQueue(command_queue), "clReleaseCommandQueue");
249 delete []events;
250 }
251
252 bool OpenCLBackend::OpenCLQueue::isProfileOrOutOfOrder() const {
253 return (backend->config -> profile || backend->config -> outOfOrder);
254 }
255
256 void printWarningLocalGroupResized(const size_t local_work_size[]) {
257 std::cout << "[Warning] Thread-Block size got automatically resized: [" << local_work_size[0] << "," << local_work_size[1] << "," << local_work_size[2] << "]" << std::endl;
258 }
259
260 void checkThreadBlockFits(OpenCLBackend *backend, const DispatchContext *dispatchContext, const size_t global_work_size[], size_t *local_work_size) {
261 size_t totalThreads = dispatchContext->lsx * dispatchContext->lsy * dispatchContext->lsz;
262 size_t max_group_size = backend->getPlatformInfo().deviceInfo.maxWorkGroupSize;
263
264 // Adjust depending on the total number of threads in the local-work-group
265 while (totalThreads > max_group_size) {
266 // Here just a simple heuristic, starting with the first dimension, 16, 4, 2 local group size
267 if (local_work_size[0] >= 16) {
268 local_work_size[0] /= 2;
269 } else if (local_work_size[1] >= 4) {
270 local_work_size[1] /= 2;
271 } else if (local_work_size[2] >= 2) {
272 local_work_size[2] /= 2;
273 }
274 totalThreads = local_work_size[0] * local_work_size[1] * local_work_size[2];
275 if (backend->config->info) {
276 printWarningLocalGroupResized(local_work_size);
277 }
278 }
279
280 // Adjust also depending on the global size. We can't launch more threads as local work than global work for
281 // each dimension
282 for (int i = 0; i < 3; i++) {
283 while (local_work_size[i] > global_work_size[i]) {
284 local_work_size[i] /= 2;
285 if (backend->config->info) {
286 printWarningLocalGroupResized(local_work_size);
287 }
288 }
289 }
290 }
291
292 void OpenCLBackend::OpenCLQueue::dispatch(DispatchContext *dispatchContext, CompilationUnit::Kernel *kernel) {
293 size_t numDimensions = dispatchContext->dimensions;
294
295 size_t global_work_size[] {
296 static_cast<size_t>(dispatchContext->gsx),
297 static_cast<size_t>(dispatchContext->gsy),
298 static_cast<size_t>(dispatchContext->gsz)
299 };
300
301 size_t local_work_size[] = {
302 static_cast<size_t>(dispatchContext->lsx),
303 static_cast<size_t>(dispatchContext->lsy),
304 static_cast<size_t>(dispatchContext->lsz),
305 };
306
307 if (dispatchContext->tlx > 0) {
308 global_work_size[0] = ceil_div(global_work_size[0], dispatchContext->tlx);
309 }
310 if (dispatchContext->tly > 0) {
311 global_work_size[1] = ceil_div(global_work_size[1], dispatchContext->tly);
312 }
313 if (dispatchContext->tlz > 0) {
314 global_work_size[2] = ceil_div(global_work_size[2], dispatchContext->tlz);
315 }
316
317 // In the OpenCL backend, we don't currently support warp-sizes to be able to run with OpenCL 1.2 (Apple)
318 // The CUDA backend supports warp-sizes
319
320 // Check the local-sizes fit
321 auto backendInstance = dynamic_cast<OpenCLBackend *>(this->backend);
322 checkThreadBlockFits(backendInstance, dispatchContext, global_work_size, local_work_size);
323
324 if (backend->config->info) {
325 backend->shortDeviceInfo();
326 std::cout << "[INFO] OpenCLBackend::OpenCLQueue::dispatch" << std::endl;
327 std::cout << "[INFO] numDimensions: " << numDimensions << std::endl;
328 std::cout << "[INFO] GLOBAL [" << global_work_size[0] << "," << global_work_size[1] << "," << global_work_size[2] << "]" << std::endl;
329 if (dispatchContext->lsx > 0) {
330 std::cout << "[INFO] LOCAL [" << local_work_size[0] << "," << local_work_size[1] << "," << local_work_size[2] << "]" << std::endl;
331 } else {
332 std::cout << "[INFO] LOCAL [ nullptr ] // The driver will setup a default value" << std::endl;
333 }
334 }
335
336 cl_uint numWaitQueues = 0;
337 cl_event* waitEvents = nullptr;
338 cl_event* nextEvents = nullptr;
339 if (isProfileOrOutOfOrder() || backendInstance -> isPlatformOnProfile()) {
340 numWaitQueues = eventc;
341 waitEvents = eventListPtr();
342 nextEvents = nextEventPtr();
343 }
344
345 const cl_int status = clEnqueueNDRangeKernel(
346 command_queue,
347 dynamic_cast<OpenCLProgram::OpenCLKernel *>(kernel)->kernel,
348 numDimensions,
349 nullptr,
350 global_work_size,
351 dispatchContext->lsx > 0 ? local_work_size : nullptr,
352 numWaitQueues,
353 waitEvents,
354 nextEvents);
355
356 inc(NDRangeBits);
357 // markAsNDRangeAndInc();
358
359 OPENCL_CHECK(status, "clEnqueueNDRangeKernel");
360 if (backend->config->trace | backend->config->traceEnqueues) {
361 std::cout << "enqueued kernel dispatch \"" << kernel->name << "\" globalSize=" << dispatchContext->gsx <<
362 std::endl;
363 }
364 }
365
366 void OpenCLBackend::OpenCLQueue::copyToDevice(Buffer *buffer) {
367 cl_uint numWaitQueues = 0;
368 cl_event* waitEvents = nullptr;
369 cl_event* nextEvents = nullptr;
370 auto backendInstance = dynamic_cast<OpenCLBackend *>(this->backend);
371 if (isProfileOrOutOfOrder() || backendInstance -> isPlatformOnProfile()) {
372 numWaitQueues = eventc;
373 waitEvents = eventListPtr();
374 nextEvents = nextEventPtr();
375 }
376 auto openclBuffer = dynamic_cast<OpenCLBuffer *>(buffer);
377 cl_int status = clEnqueueWriteBuffer(
378 command_queue,
379 openclBuffer->clMem,
380 CL_FALSE,
381 0,
382 buffer->bufferState->length,
383 buffer->bufferState->ptr,
384 numWaitQueues,
385 waitEvents,
386 nextEvents
387 );
388
389 OPENCL_CHECK(status, "clEnqueueWriteBuffer");
390
391 inc(CopyToDeviceBits);
392 // markAsCopyToDeviceAndInc();
393 }
394
395 void OpenCLBackend::OpenCLQueue::copyFromDevice(Buffer *buffer) {
396 cl_uint numWaitQueues = 0;
397 cl_event* waitEvents = nullptr;
398 cl_event* nextEvents = nullptr;
399 auto backendInstance = dynamic_cast<OpenCLBackend *>(this->backend);
400 if (isProfileOrOutOfOrder() || backendInstance -> isPlatformOnProfile()) {
401 numWaitQueues = eventc;
402 waitEvents = eventListPtr();
403 nextEvents = nextEventPtr();
404 }
405
406 auto openclBuffer = dynamic_cast<OpenCLBuffer *>(buffer);
407 cl_int status = clEnqueueReadBuffer(
408 command_queue,
409 openclBuffer->clMem,
410 CL_FALSE,
411 0,
412 buffer->bufferState->length,
413 buffer->bufferState->ptr,
414 numWaitQueues,
415 waitEvents,
416 nextEvents
417 );
418 OPENCL_CHECK(status, "clEnqueueReadBuffer");
419 inc(CopyFromDeviceBits);
420 //markAsCopyFromDeviceAndInc();
421 }