1 //
2 // Copyright 2025 The ANGLE Project Authors. All rights reserved.
3 // Use of this source code is governed by a BSD-style license that can be
4 // found in the LICENSE file.
5 //
6 // FrameCaptureCL.cpp:
7 // ANGLE CL Frame capture implementation.
8 //
9 #include "libANGLE/capture/FrameCapture.h"
10
11 #include "common/angle_version_info.h"
12 #include "common/frame_capture_utils.h"
13 #include "common/serializer/JsonSerializer.h"
14 #include "common/string_utils.h"
15 #include "common/system_utils.h"
16 #include "libANGLE/CLBuffer.h"
17 #include "libANGLE/CLCommandQueue.h"
18 #include "libANGLE/CLContext.h"
19 #include "libANGLE/CLImage.h"
20 #include "libANGLE/CLProgram.h"
21 #include "libANGLE/capture/capture_cl_autogen.h"
22 #include "libANGLE/capture/serialize.h"
23 #include "libANGLE/cl_utils.h"
24 #include "libGLESv2/cl_stubs_autogen.h"
25
26 #if !ANGLE_CAPTURE_ENABLED
27 # error Frame capture must be enabled to include this file.
28 #endif // !ANGLE_CAPTURE_ENABLED
29
30 #ifndef ANGLE_ENABLE_CL
31 # error OpenCL must be enabled to include this file.
32 #endif // !ANGLE_ENABLE_CL
33
34 namespace angle
35 {
36
37 // Some replay functions can get quite large. If over a certain size, this method breaks up the
38 // function into parts to avoid overflowing the stack and causing slow compilation.
WriteCppReplayFunctionWithPartsCL(ReplayFunc replayFunc,ReplayWriter & replayWriter,uint32_t frameIndex,std::vector<uint8_t> * binaryData,const std::vector<CallCapture> & calls,std::stringstream & header,std::stringstream & out)39 void WriteCppReplayFunctionWithPartsCL(ReplayFunc replayFunc,
40 ReplayWriter &replayWriter,
41 uint32_t frameIndex,
42 std::vector<uint8_t> *binaryData,
43 const std::vector<CallCapture> &calls,
44 std::stringstream &header,
45 std::stringstream &out)
46 {
47 out << "void "
48 << FmtFunction(replayFunc, kNoContextId, FuncUsage::Definition, frameIndex, kNoPartId)
49 << "\n"
50 << "{\n";
51
52 for (const CallCapture &call : calls)
53 {
54 // Process active calls for Setup and inactive calls for SetupInactive
55 if ((call.isActive && replayFunc != ReplayFunc::SetupInactive) ||
56 (!call.isActive && replayFunc == ReplayFunc::SetupInactive))
57 {
58 out << " ";
59 WriteCppReplayForCallCL(call, replayWriter, out, header, binaryData);
60 out << ";\n";
61 }
62 }
63 out << "}\n";
64 }
65
WriteCppReplayForCallCL(const CallCapture & call,ReplayWriter & replayWriter,std::ostream & out,std::ostream & header,std::vector<uint8_t> * binaryData)66 void WriteCppReplayForCallCL(const CallCapture &call,
67 ReplayWriter &replayWriter,
68 std::ostream &out,
69 std::ostream &header,
70 std::vector<uint8_t> *binaryData)
71 {
72 if (call.customFunctionName == "Comment")
73 {
74 // Just write it directly to the file and move on
75 WriteComment(out, call);
76 return;
77 }
78
79 std::ostringstream callOut;
80 std::ostringstream postCallAdditions;
81
82 const ParamCapture &returnValue = call.params.getReturnValue();
83 switch (returnValue.type)
84 {
85 case ParamType::Tcl_context:
86 callOut << "clContextsMap["
87 << cl::Platform::GetDefault()->getFrameCaptureShared()->getIndex(
88 &returnValue.value.cl_contextVal)
89 << "] = ";
90 break;
91 case ParamType::Tcl_command_queue:
92 callOut << "clCommandQueuesMap["
93 << cl::Platform::GetDefault()->getFrameCaptureShared()->getIndex(
94 &returnValue.value.cl_command_queueVal)
95 << "] = ";
96 break;
97 case ParamType::Tcl_mem:
98 callOut << "clMemMap["
99 << cl::Platform::GetDefault()->getFrameCaptureShared()->getIndex(
100 &returnValue.value.cl_memVal)
101 << "] = ";
102 break;
103 case ParamType::Tcl_sampler:
104 callOut << "clSamplerMap["
105 << cl::Platform::GetDefault()->getFrameCaptureShared()->getIndex(
106 &returnValue.value.cl_samplerVal)
107 << "] = ";
108 break;
109 case ParamType::Tcl_program:
110 callOut << "clProgramsMap["
111 << cl::Platform::GetDefault()->getFrameCaptureShared()->getIndex(
112 &returnValue.value.cl_programVal)
113 << "] = ";
114 break;
115 case ParamType::Tcl_kernel:
116 callOut << "clKernelsMap["
117 << cl::Platform::GetDefault()->getFrameCaptureShared()->getIndex(
118 &returnValue.value.cl_kernelVal)
119 << "] = ";
120 break;
121 case ParamType::Tcl_event:
122 callOut << "clEventsMap["
123 << cl::Platform::GetDefault()->getFrameCaptureShared()->getIndex(
124 &returnValue.value.cl_eventVal)
125 << "] = ";
126 break;
127 case ParamType::TvoidPointer:
128 if (cl::Platform::GetDefault()->getFrameCaptureShared()->getCLVoidIndex(
129 returnValue.value.voidPointerVal) != SIZE_MAX)
130 {
131 callOut << "clVoidMap["
132 << cl::Platform::GetDefault()->getFrameCaptureShared()->getCLVoidIndex(
133 returnValue.value.voidPointerVal)
134 << "] = ";
135 }
136 break;
137 default:
138 break;
139 }
140
141 callOut << call.name() << "(";
142
143 bool first = true;
144 for (const ParamCapture ¶m : call.params.getParamCaptures())
145 {
146 if (!first)
147 {
148 callOut << ", ";
149 }
150
151 if (param.arrayClientPointerIndex != -1 && param.value.voidConstPointerVal != nullptr)
152 {
153 callOut << "gClientArrays[" << param.arrayClientPointerIndex << "]";
154 }
155 else if (param.readBufferSizeBytes > 0)
156 {
157 callOut << "(" << ParamTypeToString(param.type) << ")gReadBuffer";
158 }
159 else if (param.data.empty())
160 {
161 if (param.type == ParamType::Tcl_platform_idPointer &&
162 param.value.cl_platform_idPointerVal)
163 {
164 callOut << "clPlatformsMap";
165 }
166 else if (param.type == ParamType::Tcl_platform_id)
167 {
168 callOut << "clPlatformsMap["
169 << cl::Platform::GetDefault()->getFrameCaptureShared()->getIndex(
170 ¶m.value.cl_platform_idVal)
171 << "]";
172 }
173 else if (param.type == ParamType::Tcl_device_idPointer &&
174 param.value.cl_device_idPointerVal)
175 {
176 std::vector<size_t> tempDeviceIndices =
177 cl::Platform::GetDefault()->getFrameCaptureShared()->getCLObjVector(¶m);
178
179 cl_uint numDevices = call.params.getParamCaptures()[2].value.cl_uintVal;
180 out << "temporaryDevicesList.clear();\n temporaryDevicesList.resize("
181 << numDevices << ");\n ";
182 callOut << "temporaryDevicesList.data()";
183 for (cl_uint i = 0; i < numDevices; ++i)
184 {
185 postCallAdditions << ";\n clDevicesMap[" << tempDeviceIndices[i]
186 << "] = temporaryDevicesList[" << i << "]";
187 }
188 }
189 else if (param.type == ParamType::Tcl_device_id)
190 {
191 callOut << "clDevicesMap["
192 << cl::Platform::GetDefault()->getFrameCaptureShared()->getIndex(
193 ¶m.value.cl_device_idVal)
194 << "]";
195 }
196 else if (param.type == ParamType::Tcl_context)
197 {
198 callOut << "clContextsMap["
199 << cl::Platform::GetDefault()->getFrameCaptureShared()->getIndex(
200 ¶m.value.cl_contextVal)
201 << "]";
202 }
203 else if (param.type == ParamType::Tcl_command_queue)
204 {
205 callOut << "clCommandQueuesMap["
206 << cl::Platform::GetDefault()->getFrameCaptureShared()->getIndex(
207 ¶m.value.cl_command_queueVal)
208 << "]";
209 }
210 else if (param.type == ParamType::Tcl_mem)
211 {
212 callOut << "clMemMap["
213 << cl::Platform::GetDefault()->getFrameCaptureShared()->getIndex(
214 ¶m.value.cl_memVal)
215 << "]";
216 }
217 else if (param.type == ParamType::Tcl_sampler)
218 {
219 callOut << "clSamplerMap["
220 << cl::Platform::GetDefault()->getFrameCaptureShared()->getIndex(
221 ¶m.value.cl_samplerVal)
222 << "]";
223 }
224 else if (param.type == ParamType::Tcl_program)
225 {
226 callOut << "clProgramsMap["
227 << cl::Platform::GetDefault()->getFrameCaptureShared()->getIndex(
228 ¶m.value.cl_programVal)
229 << "]";
230 }
231 else if (param.type == ParamType::Tcl_kernel)
232 {
233 callOut << "clKernelsMap["
234 << cl::Platform::GetDefault()->getFrameCaptureShared()->getIndex(
235 ¶m.value.cl_kernelVal)
236 << "]";
237 }
238 else if (param.type == ParamType::Tcl_event)
239 {
240 callOut << "clEventsMap["
241 << cl::Platform::GetDefault()->getFrameCaptureShared()->getIndex(
242 ¶m.value.cl_eventVal)
243 << "]";
244 }
245 else if (param.type == ParamType::Tcl_eventPointer)
246 {
247 if (param.value.cl_eventPointerVal)
248 {
249 callOut << "&clEventsMap["
250 << cl::Platform::GetDefault()->getFrameCaptureShared()->getIndex(
251 ¶m.value.cl_eventVal)
252 << "]";
253 }
254 else
255 {
256 callOut << "NULL";
257 }
258 }
259 else if (param.type == ParamType::TvoidConstPointer)
260 {
261 if (cl::Platform::GetDefault()->getFrameCaptureShared()->getIndex(
262 ¶m.value.cl_memVal) != SIZE_MAX)
263 {
264 callOut << "(const void *)" << "&clMemMap["
265 << cl::Platform::GetDefault()->getFrameCaptureShared()->getIndex(
266 ¶m.value.cl_memVal)
267 << "]";
268 }
269 else if (cl::Platform::GetDefault()->getFrameCaptureShared()->getIndex(
270 ¶m.value.cl_samplerVal) != SIZE_MAX)
271 {
272 callOut << "(const void *)" << "&clSamplerMap["
273 << cl::Platform::GetDefault()->getFrameCaptureShared()->getIndex(
274 ¶m.value.cl_samplerVal)
275 << "]";
276 }
277 else if (cl::Platform::GetDefault()->getFrameCaptureShared()->getIndex(
278 ¶m.value.cl_command_queueVal) != SIZE_MAX)
279 {
280 callOut << "(const void *)" << "&clCommandQueuesMap["
281 << cl::Platform::GetDefault()->getFrameCaptureShared()->getIndex(
282 ¶m.value.cl_command_queueVal)
283 << "]";
284 }
285 else if (cl::Platform::GetDefault()->getFrameCaptureShared()->getCLVoidIndex(
286 param.value.voidConstPointerVal) != SIZE_MAX)
287 {
288 callOut << "clVoidMap["
289 << cl::Platform::GetDefault()->getFrameCaptureShared()->getCLVoidIndex(
290 param.value.voidConstPointerVal)
291 << "]";
292 }
293 else
294 {
295 WriteParamCaptureReplay(callOut, call, param);
296 }
297 }
298 else if (param.type == ParamType::TvoidPointer)
299 {
300 if (cl::Platform::GetDefault()->getFrameCaptureShared()->getCLVoidIndex(
301 param.value.voidPointerVal) != SIZE_MAX)
302 {
303 callOut << "clVoidMap["
304 << cl::Platform::GetDefault()->getFrameCaptureShared()->getCLVoidIndex(
305 param.value.voidPointerVal)
306 << "]";
307 }
308 else
309 {
310 WriteParamCaptureReplay(callOut, call, param);
311 }
312 }
313 else if (param.type == ParamType::Tcl_mem_destructor_func_type ||
314 param.type == ParamType::Tcl_callback_func_type ||
315 param.type == ParamType::Tcl_svm_free_callback_func_type ||
316 param.type == ParamType::Tcl_program_func_type ||
317 param.type == ParamType::Tcl_context_destructor_func_type ||
318 param.type == ParamType::Tcl_context_func_type ||
319 param.type == ParamType::Tcl_void_func_type)
320 {
321 callOut << "NULL";
322 }
323 else if (param.type == ParamType::Tcl_memConstPointer && cl::Platform::GetDefault()
324 ->getFrameCaptureShared()
325 ->getCLObjVector(¶m)
326 .size())
327 {
328 std::vector<size_t> tempBufferIndices =
329 cl::Platform::GetDefault()->getFrameCaptureShared()->getCLObjVector(¶m);
330 out << "temporaryBuffersList = {";
331 for (uint32_t i = 0; i < tempBufferIndices.size(); ++i)
332 {
333 out << (i != 0 ? ", " : "") << "clMemMap[" << tempBufferIndices.at(i) << "]";
334 }
335 out << "};\n ";
336 callOut << "temporaryBuffersList.data()";
337 }
338 else if (param.type == ParamType::Tcl_eventConstPointer)
339 {
340 std::vector<size_t> tempEventIndices =
341 cl::Platform::GetDefault()->getFrameCaptureShared()->getCLObjVector(¶m);
342 if (tempEventIndices.empty())
343 {
344 callOut << "NULL";
345 }
346 else
347 {
348 out << "temporaryEventsList = {";
349 for (uint32_t i = 0; i < tempEventIndices.size(); ++i)
350 {
351 out << (i != 0 ? ", " : "") << "clEventsMap[" << tempEventIndices.at(i)
352 << "]";
353 }
354 out << "};\n ";
355 callOut << "temporaryEventsList.data()";
356 }
357 }
358 else if (param.type == ParamType::Tcl_device_idConstPointer)
359 {
360 std::vector<size_t> tempDeviceIndices =
361 cl::Platform::GetDefault()->getFrameCaptureShared()->getCLObjVector(¶m);
362 if (tempDeviceIndices.empty())
363 {
364 callOut << "NULL";
365 }
366 else
367 {
368 out << "temporaryDevicesList = {";
369 for (uint32_t i = 0; i < tempDeviceIndices.size(); ++i)
370 {
371 if (i != 0)
372 {
373 out << ", ";
374 }
375 out << "clDevicesMap[" << tempDeviceIndices.at(i) << "]";
376 }
377 out << "};\n ";
378 callOut << "temporaryDevicesList.data()";
379 }
380 }
381 else if (param.type == ParamType::Tcl_kernelPointer)
382 {
383 std::vector<size_t> tempKernelIndices =
384 cl::Platform::GetDefault()->getFrameCaptureShared()->getCLObjVector(¶m);
385 cl_uint numKernels = call.params.getParamCaptures()[1].value.cl_uintVal;
386 out << "temporaryKernelsList.clear();\ntemporaryKernelsList.resize(" << numKernels
387 << ");\n ";
388 callOut << "temporaryKernelsList.data()";
389 for (cl_uint i = 0; i < numKernels; ++i)
390 {
391 postCallAdditions << ";\n clKernelsMap[" << tempKernelIndices[i]
392 << "] = temporaryKernelsList[" << i << "]";
393 }
394 }
395 else if (param.type == ParamType::TvoidConstPointerPointer &&
396 cl::Platform::GetDefault()
397 ->getFrameCaptureShared()
398 ->getCLObjVector(¶m)
399 .size())
400 {
401 std::vector<size_t> offsets =
402 cl::Platform::GetDefault()->getFrameCaptureShared()->getCLObjVector(¶m);
403 out << "temporaryVoidPtrList = {";
404 for (size_t i = 0; i < offsets.size(); ++i)
405 {
406 out << (i != 0 ? ", " : "") << "&((char*)temporaryVoidPtr)[" << offsets.at(i)
407 << "]";
408 }
409 out << "};\n ";
410 callOut << "temporaryVoidPtrList.data()";
411 }
412 else if (param.type == ParamType::TvoidPointerPointer ||
413 param.type == ParamType::TvoidConstPointerPointer)
414 {
415 std::vector<size_t> tempVoidIndices =
416 cl::Platform::GetDefault()->getFrameCaptureShared()->getCLObjVector(¶m);
417 out << "temporaryVoidPtrList = {";
418 for (uint32_t i = 0; i < tempVoidIndices.size(); ++i)
419 {
420 out << (i != 0 ? ", " : "") << "clVoidMap[" << tempVoidIndices.at(i) << "]";
421 }
422 out << "};\n ";
423 callOut << "temporaryVoidPtrList.data()";
424 }
425 else if (param.type == ParamType::Tcl_programConstPointer && param.value.size_tVal)
426 {
427 std::vector<size_t> tempProgramIndices =
428 cl::Platform::GetDefault()->getFrameCaptureShared()->getCLObjVector(¶m);
429 out << "temporaryProgramsList = {";
430 for (uint32_t i = 0; i < tempProgramIndices.size(); ++i)
431 {
432 out << (i != 0 ? ", " : "") << "clProgramsMap[" << tempProgramIndices.at(i)
433 << "]";
434 }
435 out << "};\n ";
436 callOut << "temporaryProgramsList.data()";
437 }
438 else if (param.type == ParamType::Tcl_context_propertiesConstPointer)
439 {
440 if (param.value.cl_context_propertiesConstPointerVal)
441 {
442 callOut << "temporaryContextProps.data()";
443 }
444 else
445 {
446 WriteParamCaptureReplay(callOut, call, param);
447 }
448 }
449 else
450 {
451 WriteParamCaptureReplay(callOut, call, param);
452 }
453 }
454 else
455 {
456 switch (param.type)
457 {
458 case ParamType::TcharConstPointerPointer:
459 WriteStringPointerParamReplay(replayWriter, callOut, header, call, param);
460 break;
461 case ParamType::Tcl_device_idPointer:
462 callOut << "clDevicesMap";
463 break;
464 case ParamType::TcharUnsignedConstPointerPointer:
465 {
466 std::string tempStructureName = "temporaryCharPointerList";
467 std::string tempStructureType = "(const char *)";
468 if (param.type == ParamType::TcharUnsignedConstPointerPointer)
469 {
470 tempStructureName = "temporaryUnsignedCharPointerList";
471 tempStructureType = "(const unsigned char *)";
472 }
473 const std::vector<uint8_t> *data;
474 out << tempStructureName << " = {";
475 for (size_t i = 0; i < param.data.size(); ++i)
476 {
477 if (i != 0)
478 {
479 out << ", ";
480 }
481 data = ¶m.data[i];
482 size_t offset = rx::roundUpPow2(binaryData->size(), kBinaryAlignment);
483 binaryData->resize(offset + data->size());
484 memcpy(binaryData->data() + offset, data->data(), data->size());
485 out << tempStructureType << "&gBinaryData[" << offset << "]";
486 }
487 out << "};\n ";
488 callOut << tempStructureName << ".data()";
489 break;
490 }
491 case ParamType::Tcl_image_descConstPointer:
492 cl_image_desc tempImageDesc;
493 std::memcpy(&tempImageDesc, param.data[0].data(), sizeof(cl_image_desc));
494 if (tempImageDesc.mem_object)
495 {
496 out << " std::memcpy(&temporaryImageDesc, ";
497 WriteBinaryParamReplay(replayWriter, out, header, call, param, binaryData);
498 out << ", sizeof(cl_image_desc));\ntemporaryImageDesc.mem_object = "
499 "clMemMap["
500 << cl::Platform::GetDefault()->getFrameCaptureShared()->getIndex(
501 &tempImageDesc.mem_object)
502 << "];\n ";
503 callOut << "&temporaryImageDesc";
504 }
505 else
506 {
507 WriteBinaryParamReplay(replayWriter, callOut, header, call, param,
508 binaryData);
509 }
510 break;
511 case ParamType::TvoidPointer:
512 {
513 // For clEnqueueNativeKernel
514 if (call.entryPoint == EntryPoint::CLEnqueueNativeKernel)
515 {
516 std::vector<size_t> bufferIndices =
517 cl::Platform::GetDefault()->getFrameCaptureShared()->getCLObjVector(
518 ¶m);
519 size_t totalSize = call.params.getParamCaptures()[3].value.size_tVal;
520 out << "temporaryVoidPtr = (void *)std::malloc(" << totalSize
521 << ");\nstd::memcpy(&temporaryVoidPtr, ";
522 WriteBinaryParamReplay(replayWriter, out, header, call, param, binaryData);
523 out << ", " << totalSize << ");\n ";
524 callOut << "temporaryVoidPtr";
525 }
526 else
527 {
528 WriteBinaryParamReplay(replayWriter, callOut, header, call, param,
529 binaryData);
530 }
531 break;
532 }
533 default:
534 WriteBinaryParamReplay(replayWriter, callOut, header, call, param, binaryData);
535 break;
536 }
537 }
538
539 first = false;
540 }
541
542 callOut << ")";
543
544 out << callOut.str() << postCallAdditions.str();
545 }
546
WriteInitReplayCallCL(bool compression,std::ostream & out,const std::string & captureLabel,size_t maxClientArraySize,size_t readBufferSize,const std::map<ParamType,uint32_t> & maxCLParamsSize)547 void WriteInitReplayCallCL(bool compression,
548 std::ostream &out,
549 const std::string &captureLabel,
550 size_t maxClientArraySize,
551 size_t readBufferSize,
552 const std::map<ParamType, uint32_t> &maxCLParamsSize)
553 {
554 std::string binaryDataFileName = GetBinaryDataFilePath(compression, captureLabel);
555 out << " // binaryDataFileName = " << binaryDataFileName << "\n";
556 out << " // maxClientArraySize = " << maxClientArraySize << "\n";
557 out << " // readBufferSize = " << readBufferSize << "\n";
558 out << " // clPlatformsMapSize = " << maxCLParamsSize.at(ParamType::Tcl_platform_idPointer)
559 << "\n";
560 out << " // clDevicesMapSize = " << maxCLParamsSize.at(ParamType::Tcl_device_idPointer)
561 << "\n";
562 out << " // clContextsMapSize = " << maxCLParamsSize.at(ParamType::Tcl_context) << "\n";
563 out << " // clCommandQueuesMapSize = " << maxCLParamsSize.at(ParamType::Tcl_command_queue)
564 << "\n";
565 out << " // clMemMapSize = " << maxCLParamsSize.at(ParamType::Tcl_mem) << "\n";
566 out << " // clEventsMapSize = " << maxCLParamsSize.at(ParamType::Tcl_eventPointer) << "\n";
567 out << " // clProgramsMapSize = " << maxCLParamsSize.at(ParamType::Tcl_program) << "\n";
568 out << " // clKernelsMapSize = " << maxCLParamsSize.at(ParamType::Tcl_kernel) << "\n";
569 out << " // clSamplerMapSize = " << maxCLParamsSize.at(ParamType::Tcl_sampler) << "\n";
570 out << " // clVoidMapSize = " << maxCLParamsSize.at(ParamType::TvoidPointer) << "\n";
571 out << " InitializeReplayCL(\"" << binaryDataFileName << "\", " << maxClientArraySize << ", "
572 << readBufferSize << ", " << maxCLParamsSize.at(ParamType::Tcl_platform_idPointer) << ", "
573 << maxCLParamsSize.at(ParamType::Tcl_device_idPointer) << ", "
574 << maxCLParamsSize.at(ParamType::Tcl_context) << ", "
575 << maxCLParamsSize.at(ParamType::Tcl_command_queue) << ", "
576 << maxCLParamsSize.at(ParamType::Tcl_mem) << ", "
577 << maxCLParamsSize.at(ParamType::Tcl_eventPointer) << ", "
578 << maxCLParamsSize.at(ParamType::Tcl_program) << ", "
579 << maxCLParamsSize.at(ParamType::Tcl_kernel) << ", "
580 << maxCLParamsSize.at(ParamType::Tcl_sampler) << ", "
581 << maxCLParamsSize.at(ParamType::TvoidPointer) << ");\n";
582 }
583
trackCLMemUpdate(const cl_mem * mem,bool referenced)584 void FrameCaptureShared::trackCLMemUpdate(const cl_mem *mem, bool referenced)
585 {
586 std::vector<cl_mem> &mCLDirtyMem = mResourceTrackerCL.mCLDirtyMem;
587 // retained or created cl mem object
588 if (referenced)
589 {
590 // Potentially mark as dirty
591 auto it = std::find(mCLDirtyMem.begin(), mCLDirtyMem.end(), *mem);
592 if (it == mCLDirtyMem.end())
593 {
594 mCLDirtyMem.push_back(*mem);
595 }
596 }
597 else
598 {
599 std::unordered_map<cl_mem, cl_mem> &mCLSubBufferToParent =
600 mResourceTrackerCL.mCLSubBufferToParent;
601 if ((*mem)->cast<cl::Memory>().getRefCount() == 1)
602 {
603 auto it = std::find(mCLDirtyMem.begin(), mCLDirtyMem.end(), *mem);
604 if (it != mCLDirtyMem.end())
605 {
606 mCLDirtyMem.erase(it);
607 }
608
609 if (removeUnneededOpenCLCalls)
610 {
611 removeCLMemOccurrences(mem, &mFrameCalls);
612 }
613 mCLSubBufferToParent.erase(*mem);
614 (*mem)->cast<cl::Memory>().release();
615 }
616 if (mCLSubBufferToParent.find(*mem) != mCLSubBufferToParent.end())
617 {
618 trackCLMemUpdate(&mCLSubBufferToParent[*mem], false);
619 }
620 }
621 }
622
trackCLProgramUpdate(const cl_program * program,bool referenced,cl_uint numLinkedPrograms,const cl_program * linkedPrograms)623 void FrameCaptureShared::trackCLProgramUpdate(const cl_program *program,
624 bool referenced,
625 cl_uint numLinkedPrograms,
626 const cl_program *linkedPrograms)
627 {
628 std::unordered_map<cl_program, cl_uint> &mCLProgramLinkCounter =
629 mResourceTrackerCL.mCLProgramLinkCounter;
630 std::unordered_map<cl_program, std::vector<cl_program>> &mCLLinkedPrograms =
631 mResourceTrackerCL.mCLLinkedPrograms;
632 // retained or created cl program object
633 if (referenced)
634 {
635 // Increment link count for this program
636 if (mCLProgramLinkCounter.find(*program) == mCLProgramLinkCounter.end())
637 {
638 mCLProgramLinkCounter[*program] = 0;
639 }
640 ++mCLProgramLinkCounter[*program];
641
642 // Setup the linked programs if this call is from capturing clCompileProgram or
643 // clLinkProgram
644 if (numLinkedPrograms)
645 {
646 mCLLinkedPrograms[*program] = std::vector<cl_program>();
647 for (cl_uint i = 0; i < numLinkedPrograms; ++i)
648 {
649 mCLLinkedPrograms[*program].push_back(linkedPrograms[i]);
650 }
651 }
652
653 // Go through the linked programs and increment their link counts
654 for (size_t i = 0; mCLLinkedPrograms.find(*program) != mCLLinkedPrograms.end() &&
655 i < mCLLinkedPrograms[*program].size();
656 ++i)
657 {
658 trackCLProgramUpdate(&mCLLinkedPrograms[*program].at(i), true, 0, nullptr);
659 }
660 }
661 else
662 {
663 // Decrement link count for this program and the linked programs
664 --mCLProgramLinkCounter[*program];
665 for (size_t i = 0; mCLLinkedPrograms.find(*program) != mCLLinkedPrograms.end() &&
666 i < mCLLinkedPrograms[*program].size();
667 ++i)
668 {
669 trackCLProgramUpdate(&mCLLinkedPrograms[*program].at(i), false, 0, nullptr);
670 }
671
672 // Remove the calls containing this object if the link count is 0
673 if (mCLProgramLinkCounter[*program] == 0)
674 {
675 mCLProgramLinkCounter.erase(*program);
676 if (mCLLinkedPrograms.find(*program) != mCLLinkedPrograms.end())
677 {
678 mCLLinkedPrograms.erase(*program);
679 }
680
681 if (removeUnneededOpenCLCalls)
682 {
683 removeCLProgramOccurrences(program, &mFrameCalls);
684 }
685 }
686 }
687 }
688
injectMemcpy(void * src,void * dest,size_t size,std::vector<CallCapture> * calls)689 void FrameCaptureShared::injectMemcpy(void *src,
690 void *dest,
691 size_t size,
692 std::vector<CallCapture> *calls)
693 {
694 // Inject memcpy call before unmap
695
696 // Create param buffer
697 ParamBuffer paramBuffer;
698
699 // Create dest parameter
700 ParamCapture destParam("dest", ParamType::TvoidConstPointer);
701 InitParamValue(ParamType::TvoidPointer, dest, &destParam.value);
702 paramBuffer.addParam(std::move(destParam));
703
704 // Create src param
705 ParamCapture updateMemory("src", ParamType::TvoidConstPointer);
706 CaptureMemory(src, size, &updateMemory);
707 paramBuffer.addParam(std::move(updateMemory));
708
709 paramBuffer.addValueParam<size_t>("size", ParamType::Tsize_t, size);
710
711 calls->emplace(calls->end() - 1, "std::memcpy", std::move(paramBuffer));
712 }
713
captureUpdateCLObjs(std::vector<CallCapture> * calls)714 void FrameCaptureShared::captureUpdateCLObjs(std::vector<CallCapture> *calls)
715 {
716 std::vector<cl_mem> &mCLDirtyMem = mResourceTrackerCL.mCLDirtyMem;
717 std::vector<void *> &mCLDirtySVM = mResourceTrackerCL.mCLDirtySVM;
718 cl_command_queue &mCLCurrentCommandQueue = mResourceTrackerCL.mCLCurrentCommandQueue;
719 for (uint32_t i = 0; i < mCLDirtyMem.size(); ++i)
720 {
721 cl_mem_object_type memType;
722 if (IsError(mCLDirtyMem.at(i)->cast<cl::Memory>().getInfo(
723 cl::MemInfo::Type, sizeof(cl_mem_object_type), &memType, nullptr)))
724 {
725 continue;
726 }
727 if (memType == CL_MEM_OBJECT_BUFFER)
728 {
729 void *ptr;
730
731 if (calls->back().entryPoint == EntryPoint::CLEnqueueUnmapMemObject)
732 {
733 CallCapture *mapCall = &mResourceTrackerCL.mCLMapCall.at(
734 calls->back()
735 .params.getParam("mapped_ptr", ParamType::TvoidPointer, 2)
736 .value.voidPointerVal);
737 size_t offset =
738 mapCall->params.getParam("offset", ParamType::Tsize_t, 4).value.size_tVal;
739 size_t size =
740 mapCall->params.getParam("size", ParamType::Tsize_t, 5).value.size_tVal;
741 ptr = malloc(size);
742
743 // Call clEnqueueReadBuffer to get the current data in the buffer
744 EnqueueReadBuffer(mCLCurrentCommandQueue, mCLDirtyMem.at(i), true, offset, size,
745 ptr, 0, nullptr, nullptr);
746
747 // Inject memcpy call BEFORE unmap
748 injectMemcpy(ptr,
749 calls->back()
750 .params.getParam("mapped_ptr", ParamType::TvoidPointer, 2)
751 .value.voidPointerVal,
752 size, calls);
753 }
754 else
755 {
756 size_t bufferSize = mCLDirtyMem.at(i)->cast<cl::Buffer>().getSize();
757 ptr = malloc(bufferSize);
758
759 // Call clEnqueueReadBuffer to get the current data in the buffer
760 EnqueueReadBuffer(mCLCurrentCommandQueue, mCLDirtyMem.at(i), true, 0, bufferSize,
761 ptr, 0, nullptr, nullptr);
762
763 // Pretend that a "clEnqueueWriteBuffer" was called with the above data retrieved
764 calls->push_back(CaptureEnqueueWriteBuffer(true, mCLCurrentCommandQueue,
765 mCLDirtyMem.at(i), true, 0, bufferSize,
766 ptr, 0, nullptr, nullptr, CL_SUCCESS));
767
768 // Implicit release, so going into the starting frame the buffer has the correct
769 // reference count
770 mCLDirtyMem.at(i)->cast<cl::Memory>().release();
771 }
772 free(ptr);
773 }
774 else if (memType == CL_MEM_OBJECT_PIPE)
775 {
776 UNIMPLEMENTED();
777 }
778 else
779 {
780 cl::Image *clImg = &mCLDirtyMem.at(i)->cast<cl::Image>();
781 void *ptr;
782
783 if (calls->back().entryPoint == EntryPoint::CLEnqueueUnmapMemObject)
784 {
785 CallCapture *mapCall = &mResourceTrackerCL.mCLMapCall.at(
786 calls->back()
787 .params.getParam("mapped_ptr", ParamType::TvoidPointer, 2)
788 .value.voidPointerVal);
789 const size_t *origin = (const size_t *)mapCall->params
790 .getParam("origin", ParamType::Tsize_tConstPointer, 4)
791 .data.back()
792 .data();
793 const size_t *region = (const size_t *)mapCall->params
794 .getParam("region", ParamType::Tsize_tConstPointer, 5)
795 .data.back()
796 .data();
797
798 size_t rowPitch = mapCall->params.getParam("image_row_pitch", ParamType::Tsize_t, 6)
799 .value.size_tVal;
800 size_t slicePitch =
801 mapCall->params.getParam("image_slice_pitch", ParamType::Tsize_t, 7)
802 .value.size_tVal;
803
804 // Get the image size to allocate the size of ptr
805 size_t totalSize = (region[2] - 1) * slicePitch + (region[1] - 1) * rowPitch +
806 region[0] * clImg->getElementSize();
807 ptr = malloc(totalSize);
808
809 // Call clEnqueueReadBuffer to get the current data in the image
810 EnqueueReadImage(mCLCurrentCommandQueue, mCLDirtyMem.at(i), true, origin, region,
811 rowPitch, slicePitch, ptr, 0, nullptr, nullptr);
812
813 // Inject memcpy call BEFORE unmap
814 injectMemcpy(ptr,
815 calls->back()
816 .params.getParam("mapped_ptr", ParamType::TvoidPointer, 2)
817 .value.voidPointerVal,
818 totalSize, calls);
819 }
820 else
821 {
822 ptr = malloc(clImg->getSize());
823 size_t origin[3] = {0, 0, 0};
824 size_t region[3] = {clImg->getWidth(), clImg->getHeight(), clImg->getDepth()};
825
826 // Call clEnqueueReadBuffer to get the current data in the image
827 EnqueueReadImage(mCLCurrentCommandQueue, mCLDirtyMem.at(i), true, origin, region,
828 clImg->getRowSize(), clImg->getSliceSize(), ptr, 0, nullptr,
829 nullptr);
830
831 // Pretend that a "clEnqueueWriteImage" was called with the above data retrieved
832 calls->push_back(CaptureEnqueueWriteImage(
833 true, mCLCurrentCommandQueue, mCLDirtyMem.at(i), true, origin, region,
834 clImg->getRowSize(), clImg->getSliceSize(), ptr, 0, nullptr, nullptr,
835 CL_SUCCESS));
836
837 // Implicit release, so going into the starting frame the buffer has the correct
838 // reference count
839 mCLDirtyMem.at(i)->cast<cl::Memory>().release();
840 }
841
842 free(ptr);
843 }
844 }
845 for (uint32_t i = 0; i < mCLDirtySVM.size(); ++i)
846 {
847 size_t SVMSize = mResourceTrackerCL.SVMToSize[mCLDirtySVM.at(i)];
848
849 // Call clEnqueueSVMMap to get the current data in the SVM pointer
850 cl::MemFlags flags;
851 flags.set(CL_MAP_READ);
852 EnqueueSVMMap(mCLCurrentCommandQueue, true, flags, mCLDirtySVM.at(i), SVMSize, 0, nullptr,
853 nullptr);
854
855 // Pretend that a "clEnqueueSVMMemcpy" was called with the above data retrieved
856 calls->push_back(CaptureEnqueueSVMMemcpy(true, mCLCurrentCommandQueue, true,
857 mCLDirtySVM.at(i), mCLDirtySVM.at(i), SVMSize, 0,
858 nullptr, nullptr, CL_SUCCESS));
859
860 // Call clEnqueueSVMUnmap to get the current data in the SVM pointer
861 EnqueueSVMUnmap(mCLCurrentCommandQueue, mResourceTrackerCL.mCLDirtySVM.at(i), 0, nullptr,
862 nullptr);
863 }
864 mCLDirtyMem.clear();
865 mCLDirtySVM.clear();
866 }
867
removeCLMemOccurrences(const cl_mem * mem,std::vector<CallCapture> * calls)868 void FrameCaptureShared::removeCLMemOccurrences(const cl_mem *mem, std::vector<CallCapture> *calls)
869 {
870 // This function gets called when it captures a clReleaseMemObj prior to the starting frame
871 // that sets the reference count to 0, meaning that this cl_mem object isn't necessary for
872 // the wanted frames. So, we can remove the calls that use it.
873
874 for (size_t i = 0; i < calls->size(); ++i)
875 {
876 CallCapture *call = &calls->at(i);
877 cl_mem foundMem;
878 switch (call->entryPoint)
879 {
880 case EntryPoint::CLCreateBuffer:
881 case EntryPoint::CLCreateBufferWithProperties:
882 case EntryPoint::CLCreateImage:
883 case EntryPoint::CLCreateImageWithProperties:
884 case EntryPoint::CLCreateImage2D:
885 case EntryPoint::CLCreateImage3D:
886 case EntryPoint::CLCreatePipe:
887 {
888 foundMem = call->params.getReturnValue().value.cl_memVal;
889 break;
890 }
891 case EntryPoint::CLCreateSubBuffer:
892 {
893 foundMem = call->params.getReturnValue().value.cl_memVal;
894 if (foundMem != *mem)
895 {
896 foundMem =
897 call->params.getParam("buffer", ParamType::Tcl_mem, 0).value.cl_memVal;
898 }
899 break;
900 }
901 case EntryPoint::CLEnqueueReadBuffer:
902 case EntryPoint::CLEnqueueWriteBuffer:
903 case EntryPoint::CLEnqueueReadBufferRect:
904 case EntryPoint::CLEnqueueWriteBufferRect:
905 case EntryPoint::CLEnqueueMapBuffer:
906 {
907 // Can get rid of these calls because the buffer is no longer needed
908 foundMem = call->params.getParam("buffer", ParamType::Tcl_mem, 1).value.cl_memVal;
909 break;
910 }
911 case EntryPoint::CLEnqueueReadImage:
912 case EntryPoint::CLEnqueueWriteImage:
913 case EntryPoint::CLEnqueueMapImage:
914 {
915 // Can get rid of these calls because the image is no longer needed
916 foundMem = call->params.getParam("image", ParamType::Tcl_mem, 1).value.cl_memVal;
917 break;
918 }
919 case EntryPoint::CLEnqueueCopyBuffer:
920 case EntryPoint::CLEnqueueCopyBufferRect:
921 case EntryPoint::CLEnqueueCopyImage:
922 case EntryPoint::CLEnqueueCopyBufferToImage:
923 case EntryPoint::CLEnqueueCopyImageToBuffer:
924 {
925 // Can get rid of these calls because the obj is no longer needed
926 std::string srcType = "src_";
927 srcType += ((call->entryPoint == EntryPoint::CLEnqueueCopyImage ||
928 call->entryPoint == EntryPoint::CLEnqueueCopyImageToBuffer)
929 ? "image"
930 : "buffer");
931 std::string dstType = "dst_";
932 dstType += ((call->entryPoint == EntryPoint::CLEnqueueCopyImage ||
933 call->entryPoint == EntryPoint::CLEnqueueCopyBufferToImage)
934 ? "image"
935 : "buffer");
936 foundMem =
937 call->params.getParam(srcType.c_str(), ParamType::Tcl_mem, 1).value.cl_memVal;
938 if (foundMem != *mem)
939 {
940 foundMem = call->params.getParam(dstType.c_str(), ParamType::Tcl_mem, 2)
941 .value.cl_memVal;
942 }
943 break;
944 }
945 case EntryPoint::CLReleaseMemObject:
946 case EntryPoint::CLRetainMemObject:
947 case EntryPoint::CLGetMemObjectInfo:
948 case EntryPoint::CLSetMemObjectDestructorCallback:
949 case EntryPoint::CLEnqueueUnmapMemObject:
950 {
951 foundMem =
952 call->params
953 .getParam("memobj", ParamType::Tcl_mem,
954 call->entryPoint == EntryPoint::CLEnqueueUnmapMemObject ? 1 : 0)
955 .value.cl_memVal;
956 break;
957 }
958 case EntryPoint::CLGetImageInfo:
959 {
960 foundMem = call->params.getParam("image", ParamType::Tcl_mem, 0).value.cl_memVal;
961 break;
962 }
963 case EntryPoint::CLSetKernelArg:
964 {
965 foundMem = call->params.getParam("arg_value", ParamType::TvoidConstPointer, 3)
966 .value.cl_memVal;
967 if (cl::Platform::GetDefault()->getFrameCaptureShared()->getIndex(&foundMem) ==
968 SIZE_MAX)
969 {
970 continue;
971 }
972 break;
973 }
974 // Leave commented until external memory is upstream
975 // case EntryPoint::CLEnqueueAcquireExternalMemObjectsKHR:
976 // case EntryPoint::CLEnqueueReleaseExternalMemObjectsKHR:
977 case EntryPoint::CLEnqueueMigrateMemObjects:
978 {
979 const cl_mem *memObjs =
980 call->params.getParam("mem_objects", ParamType::Tcl_memConstPointer, 2)
981 .value.cl_memConstPointerVal;
982 cl_uint numMemObjs =
983 call->params.getParam("num_mem_objects", ParamType::Tcl_uint, 1)
984 .value.cl_uintVal;
985
986 std::vector<cl_mem> newMemObjs;
987 for (cl_uint memObjIndex = 0; i < numMemObjs; ++i)
988 {
989 if (memObjs[memObjIndex] != *mem)
990 {
991 newMemObjs.push_back(memObjs[memObjIndex]);
992 }
993 }
994
995 // If all the mem objects used in this array are released, I can remove this call
996 if (newMemObjs.empty())
997 {
998 foundMem = *mem;
999 }
1000 else
1001 {
1002 call->params.setValueParamAtIndex("num_mem_objects", ParamType::Tcl_uint,
1003 newMemObjs.size(), 1);
1004 setCLObjVectorMap(
1005 newMemObjs.data(), newMemObjs.size(),
1006 &call->params.getParam("mem_objects", ParamType::Tcl_memConstPointer, 2),
1007 &angle::FrameCaptureShared::getIndex);
1008 continue;
1009 }
1010 break;
1011 }
1012 default:
1013 continue;
1014 }
1015
1016 if (foundMem == *mem)
1017 {
1018 removeCLCall(calls, i);
1019 --i;
1020 }
1021 }
1022 }
1023
removeCLKernelOccurrences(const cl_kernel * kernel,std::vector<CallCapture> * calls)1024 void FrameCaptureShared::removeCLKernelOccurrences(const cl_kernel *kernel,
1025 std::vector<CallCapture> *calls)
1026 {
1027 // This function gets called when it captures a clReleaseProgram prior to the starting frame
1028 // that sets the program's reference count to 0. This ensures that the kernels in that program
1029 // are/should be released as well, meaning that this cl_kernel object isn't necessary for
1030 // the wanted frames. So, we can remove the calls that use it.
1031 // We cannot remove cl_kernel occurrences at the time of clReleaseKernel because the kernel may
1032 // be an input to clCloneKernel and clCreateKernelsInProgram.
1033
1034 for (size_t i = 0; i < calls->size(); ++i)
1035 {
1036 CallCapture *call = &calls->at(i);
1037 cl_kernel foundKernel;
1038 switch (call->entryPoint)
1039 {
1040 case EntryPoint::CLCreateKernel:
1041 {
1042 foundKernel = call->params.getReturnValue().value.cl_kernelVal;
1043 break;
1044 }
1045 case EntryPoint::CLCloneKernel:
1046 {
1047 foundKernel = call->params.getReturnValue().value.cl_kernelVal;
1048 if (foundKernel != *kernel)
1049 {
1050 foundKernel = call->params.getParam("source_kernel", ParamType::Tcl_kernel, 0)
1051 .value.cl_kernelVal;
1052 }
1053 break;
1054 }
1055 case EntryPoint::CLRetainKernel:
1056 case EntryPoint::CLReleaseKernel:
1057 case EntryPoint::CLSetKernelArg:
1058 case EntryPoint::CLSetKernelArgSVMPointer:
1059 case EntryPoint::CLSetKernelExecInfo:
1060 case EntryPoint::CLGetKernelInfo:
1061 case EntryPoint::CLGetKernelArgInfo:
1062 case EntryPoint::CLGetKernelWorkGroupInfo:
1063 case EntryPoint::CLGetKernelSubGroupInfo:
1064 {
1065 foundKernel =
1066 call->params.getParam("kernel", ParamType::Tcl_kernel, 0).value.cl_kernelVal;
1067 break;
1068 }
1069 case EntryPoint::CLEnqueueNDRangeKernel:
1070 case EntryPoint::CLEnqueueTask:
1071 {
1072 foundKernel =
1073 call->params.getParam("kernel", ParamType::Tcl_kernel, 1).value.cl_kernelVal;
1074 break;
1075 }
1076
1077 default:
1078 continue;
1079 }
1080
1081 if (foundKernel == *kernel)
1082 {
1083 removeCLCall(calls, i);
1084 --i;
1085 }
1086 }
1087 }
1088
removeCLProgramOccurrences(const cl_program * program,std::vector<CallCapture> * calls)1089 void FrameCaptureShared::removeCLProgramOccurrences(const cl_program *program,
1090 std::vector<CallCapture> *calls)
1091 {
1092 // This function gets called when it captures a clReleaseMemObj prior to the starting frame
1093 // that sets the reference count to 0, and the program is not linked to any other program,
1094 // meaning that this cl_mem object isn't necessary for the wanted frames. So, we can
1095 // remove the calls that use it.
1096
1097 for (size_t i = 0; i < calls->size(); ++i)
1098 {
1099 CallCapture *call = &calls->at(i);
1100 cl_program foundProgram;
1101 switch (call->entryPoint)
1102 {
1103 case EntryPoint::CLCreateProgramWithSource:
1104 case EntryPoint::CLCreateProgramWithBinary:
1105 case EntryPoint::CLCreateProgramWithBuiltInKernels:
1106 case EntryPoint::CLCreateProgramWithIL:
1107 case EntryPoint::CLLinkProgram:
1108 {
1109 foundProgram = call->params.getReturnValue().value.cl_programVal;
1110 break;
1111 }
1112 case EntryPoint::CLRetainProgram:
1113 case EntryPoint::CLReleaseProgram:
1114 case EntryPoint::CLBuildProgram:
1115 case EntryPoint::CLGetProgramInfo:
1116 case EntryPoint::CLGetProgramBuildInfo:
1117 case EntryPoint::CLCreateKernel:
1118 case EntryPoint::CLCreateKernelsInProgram:
1119 case EntryPoint::CLUnloadPlatformCompiler:
1120 case EntryPoint::CLCompileProgram:
1121 {
1122 uint8_t programIndex = call->entryPoint == EntryPoint::CLCompileProgram ? 1 : 0;
1123 foundProgram =
1124 call->params.getParam("program", ParamType::Tcl_program, programIndex)
1125 .value.cl_programVal;
1126 break;
1127 }
1128 default:
1129 continue;
1130 }
1131
1132 if (foundProgram == *program)
1133 {
1134 removeCLCall(calls, i);
1135 --i;
1136 }
1137 }
1138
1139 if (mResourceTrackerCL.mCLProgramToKernels.find(*program) !=
1140 mResourceTrackerCL.mCLProgramToKernels.end())
1141 {
1142 for (size_t i = 0; i < mResourceTrackerCL.mCLProgramToKernels[*program].size(); ++i)
1143 {
1144 removeCLKernelOccurrences(&mResourceTrackerCL.mCLProgramToKernels[*program].at(i),
1145 calls);
1146 }
1147 mResourceTrackerCL.mCLProgramToKernels.erase(*program);
1148 }
1149 }
1150
removeCLCall(std::vector<CallCapture> * callVector,size_t & callIndex)1151 void FrameCaptureShared::removeCLCall(std::vector<CallCapture> *callVector, size_t &callIndex)
1152 {
1153 CallCapture *call = &callVector->at(callIndex);
1154 const std::vector<ParamCapture> *params = &call->params.getParamCaptures();
1155 cl_context context = nullptr;
1156
1157 // Checks if there is an event that is implicitly created during the deleted call.
1158 // If there is, need to inject a clCreateUserEvent call and a clSetUserEventStatus call.
1159 for (auto ¶m : *params)
1160 {
1161 if (param.type == ParamType::Tcl_context)
1162 {
1163 context = param.value.cl_contextVal;
1164 }
1165 else if (param.type == ParamType::Tcl_command_queue)
1166 {
1167 context =
1168 param.value.cl_command_queueVal->cast<cl::CommandQueue>().getContext().getNative();
1169 }
1170 else if (param.type == ParamType::Tcl_eventPointer && param.value.cl_eventVal != nullptr &&
1171 context)
1172 {
1173 // Capture the creation of a successful event if the CL call being removed created an
1174 // event (ex: clEnqueueReadBuffer)
1175 cl_event event = param.value.cl_eventVal;
1176 callVector->insert(callVector->begin() + callIndex,
1177 CaptureSetUserEventStatus(true, event, CL_COMPLETE, CL_SUCCESS));
1178 callVector->insert(callVector->begin() + callIndex,
1179 CaptureCreateUserEvent(true, context, nullptr, event));
1180 callIndex += 2;
1181 break;
1182 }
1183 }
1184 callVector->erase(callVector->begin() + callIndex);
1185 }
1186
maybeCapturePreCallUpdatesCL(CallCapture & call)1187 void FrameCaptureShared::maybeCapturePreCallUpdatesCL(CallCapture &call)
1188 {
1189 switch (call.entryPoint)
1190 {
1191 case EntryPoint::CLGetExtensionFunctionAddress:
1192 case EntryPoint::CLGetExtensionFunctionAddressForPlatform:
1193 {
1194 uint32_t index = call.entryPoint == EntryPoint::CLGetExtensionFunctionAddress ? 0 : 1;
1195 std::string funcName =
1196 (const char *)call.params.getParam("func_name", ParamType::TcharConstPointer, index)
1197 .value.charConstPointerPointerVal;
1198 call.customFunctionName =
1199 funcName + " = (" + funcName + "_fn)" + GetEntryPointName(call.entryPoint);
1200
1201 if (std::find(mExtFuncsAdded.begin(), mExtFuncsAdded.end(), funcName) ==
1202 mExtFuncsAdded.end())
1203 {
1204 mExtFuncsAdded.push_back(funcName);
1205 }
1206 break;
1207 }
1208 case EntryPoint::CLCreateContext:
1209 case EntryPoint::CLCreateContextFromType:
1210 {
1211 if (call.params.getParam("properties", ParamType::Tcl_context_propertiesConstPointer, 0)
1212 .value.cl_context_propertiesConstPointerVal)
1213 {
1214 size_t propSize = 0;
1215 size_t platformIDIndex = 0;
1216 const cl_context_properties *propertiesData =
1217 (cl_context_properties *)call.params
1218 .getParam("properties", ParamType::Tcl_context_propertiesConstPointer, 0)
1219 .data[0]
1220 .data();
1221 while (propertiesData[propSize] != 0)
1222 {
1223 if (propertiesData[propSize] == CL_CONTEXT_PLATFORM)
1224 {
1225 // "Each property name is immediately followed by the corresponding desired
1226 // value"
1227 platformIDIndex = propSize + 1;
1228 }
1229 ++propSize;
1230 }
1231 ++propSize;
1232
1233 if (platformIDIndex == 0)
1234 {
1235 ParamBuffer params;
1236
1237 params.addValueParam("propSize", ParamType::Tsize_t, propSize);
1238
1239 ParamCapture propertiesParam("propData",
1240 ParamType::Tcl_context_propertiesConstPointer);
1241 InitParamValue(ParamType::Tcl_context_propertiesConstPointer, propertiesData,
1242 &propertiesParam.value);
1243 CaptureMemory(propertiesData, propSize * sizeof(cl_context_properties),
1244 &propertiesParam);
1245 params.addParam(std::move(propertiesParam));
1246 mFrameCalls.emplace_back(
1247 CallCapture("UpdateCLContextPropertiesNoPlatform", std::move(params)));
1248
1249 call.params
1250 .getParam("properties", ParamType::Tcl_context_propertiesConstPointer, 0)
1251 .data.clear();
1252 break;
1253 }
1254
1255 // Create call to UpdateCLContextProperties
1256 ParamBuffer params;
1257
1258 params.addValueParam("propSize", ParamType::Tsize_t, propSize);
1259
1260 ParamCapture propertiesParam("propData",
1261 ParamType::Tcl_context_propertiesConstPointer);
1262 InitParamValue(ParamType::Tcl_context_propertiesConstPointer, propertiesData,
1263 &propertiesParam.value);
1264 CaptureMemory(propertiesData, propSize * sizeof(cl_context_properties),
1265 &propertiesParam);
1266 params.addParam(std::move(propertiesParam));
1267
1268 params.addValueParam("platformIdxInProps", ParamType::Tsize_t, platformIDIndex);
1269 params.addValueParam("platformIdxInMap", ParamType::Tsize_t,
1270 getIndex((cl_platform_id *)&propertiesData[platformIDIndex]));
1271
1272 call.params.getParam("properties", ParamType::Tcl_context_propertiesConstPointer, 0)
1273 .data.clear();
1274
1275 mFrameCalls.emplace_back(
1276 CallCapture("UpdateCLContextPropertiesWithPlatform", std::move(params)));
1277 }
1278 break;
1279 }
1280 default:
1281 break;
1282 }
1283
1284 updateReadBufferSize(call.params.getReadBufferSize());
1285 }
1286
addCLResetObj(const ParamCapture & param)1287 void FrameCaptureShared::addCLResetObj(const ParamCapture ¶m)
1288 {
1289 mResourceTrackerCL.mCLResetObjs.push_back(angle::ParamCapture("resetObj", param.type));
1290 auto paramValue =
1291 &mResourceTrackerCL.mCLResetObjs.at(mResourceTrackerCL.mCLResetObjs.size() - 1).value;
1292 switch (param.type)
1293 {
1294 case ParamType::Tcl_device_id:
1295 InitParamValue(param.type, param.value.cl_device_idVal, paramValue);
1296 break;
1297 case ParamType::Tcl_mem:
1298 InitParamValue(param.type, param.value.cl_memVal, paramValue);
1299 break;
1300 case ParamType::Tcl_kernel:
1301 InitParamValue(param.type, param.value.cl_kernelVal, paramValue);
1302 break;
1303 case ParamType::Tcl_program:
1304 InitParamValue(param.type, param.value.cl_programVal, paramValue);
1305 break;
1306 case ParamType::Tcl_command_queue:
1307 InitParamValue(param.type, param.value.cl_command_queueVal, paramValue);
1308 break;
1309 case ParamType::Tcl_context:
1310 InitParamValue(param.type, param.value.cl_contextVal, paramValue);
1311 break;
1312 case ParamType::Tcl_sampler:
1313 InitParamValue(param.type, param.value.cl_samplerVal, paramValue);
1314 break;
1315 case ParamType::Tcl_event:
1316 InitParamValue(param.type, param.value.cl_eventVal, paramValue);
1317 break;
1318 default:
1319 break;
1320 }
1321 }
1322
removeCLResetObj(const ParamCapture & param)1323 void FrameCaptureShared::removeCLResetObj(const ParamCapture ¶m)
1324 {
1325 std::vector<ParamCapture> &mCLResetObjs = mResourceTrackerCL.mCLResetObjs;
1326 for (size_t i = 0; i < mCLResetObjs.size(); ++i)
1327 {
1328 bool foundCLObj =
1329 param.type == mCLResetObjs.at(i).type &&
1330 ((param.type == ParamType::Tcl_device_id &&
1331 param.value.cl_device_idVal == mCLResetObjs.at(i).value.cl_device_idVal) ||
1332 (param.type == ParamType::Tcl_mem &&
1333 param.value.cl_memVal == mCLResetObjs.at(i).value.cl_memVal) ||
1334 (param.type == ParamType::Tcl_kernel &&
1335 param.value.cl_kernelVal == mCLResetObjs.at(i).value.cl_kernelVal) ||
1336 (param.type == ParamType::Tcl_program &&
1337 param.value.cl_programVal == mCLResetObjs.at(i).value.cl_programVal) ||
1338 (param.type == ParamType::Tcl_command_queue &&
1339 param.value.cl_command_queueVal == mCLResetObjs.at(i).value.cl_command_queueVal) ||
1340 (param.type == ParamType::Tcl_context &&
1341 param.value.cl_contextVal == mCLResetObjs.at(i).value.cl_contextVal) ||
1342 (param.type == ParamType::Tcl_sampler &&
1343 param.value.cl_samplerVal == mCLResetObjs.at(i).value.cl_samplerVal) ||
1344 (param.type == ParamType::Tcl_event &&
1345 param.value.cl_eventVal == mCLResetObjs.at(i).value.cl_eventVal));
1346
1347 if (foundCLObj)
1348 {
1349 mCLResetObjs.erase(mCLResetObjs.begin() + i);
1350 break;
1351 }
1352 }
1353 }
1354
printCLResetObjs(std::stringstream & stream)1355 void FrameCaptureShared::printCLResetObjs(std::stringstream &stream)
1356 {
1357 std::vector<ParamCapture> &mCLResetObjs = mResourceTrackerCL.mCLResetObjs;
1358 for (size_t i = 0; i < mCLResetObjs.size(); ++i)
1359 {
1360 stream << " ";
1361 switch (mCLResetObjs.at(i).type)
1362 {
1363 case ParamType::Tcl_device_id:
1364 stream << "clReleaseDevice(clDevicesMap["
1365 << std::to_string(getIndex(&mCLResetObjs.at(i).value.cl_device_idVal))
1366 << "]);";
1367 break;
1368 case ParamType::Tcl_mem:
1369 stream << "clReleaseMemObject(clMemMap["
1370 << std::to_string(getIndex(&mCLResetObjs.at(i).value.cl_memVal)) << "]);";
1371 break;
1372 case ParamType::Tcl_kernel:
1373 stream << "clReleaseKernel(clKernelsMap["
1374 << std::to_string(getIndex(&mCLResetObjs.at(i).value.cl_kernelVal)) << "]);";
1375 break;
1376 case ParamType::Tcl_program:
1377 stream << "clReleaseProgram(clProgramsMap["
1378 << std::to_string(getIndex(&mCLResetObjs.at(i).value.cl_programVal))
1379 << "]);";
1380 break;
1381 case ParamType::Tcl_command_queue:
1382 stream << "clReleaseCommandQueue(clCommandQueuesMap["
1383 << std::to_string(getIndex(&mCLResetObjs.at(i).value.cl_command_queueVal))
1384 << "]);";
1385 break;
1386 case ParamType::Tcl_context:
1387 stream << "clReleaseContext(clContextsMap["
1388 << std::to_string(getIndex(&mCLResetObjs.at(i).value.cl_contextVal))
1389 << "]);";
1390 break;
1391 case ParamType::Tcl_sampler:
1392 stream << "clReleaseSampler(clSamplersMap["
1393 << std::to_string(getIndex(&mCLResetObjs.at(i).value.cl_samplerVal))
1394 << "]);";
1395 break;
1396 case ParamType::Tcl_event:
1397 stream << "clReleaseEvent(clEventsMap["
1398 << std::to_string(getIndex(&mCLResetObjs.at(i).value.cl_eventVal)) << "]);";
1399 break;
1400 default:
1401 break;
1402 }
1403 stream << "\n";
1404 }
1405 }
1406
updateResourceCountsFromParamCaptureCL(const ParamCapture & param,const CallCapture & call)1407 void FrameCaptureShared::updateResourceCountsFromParamCaptureCL(const ParamCapture ¶m,
1408 const CallCapture &call)
1409 {
1410 switch (param.type)
1411 {
1412 case ParamType::Tcl_platform_idPointer:
1413 if (call.entryPoint == EntryPoint::CLIcdGetPlatformIDsKHR ||
1414 call.entryPoint == EntryPoint::CLGetPlatformIDs)
1415 {
1416 mMaxCLParamsSize[param.type] +=
1417 sizeof(cl_platform_id) * ((call.params.getParamCaptures()[0]).value.cl_uintVal);
1418 }
1419 break;
1420 case ParamType::Tcl_device_idPointer:
1421 if (call.entryPoint == EntryPoint::CLGetDeviceIDs)
1422 {
1423 mMaxCLParamsSize[param.type] +=
1424 sizeof(cl_device_id) * ((call.params.getParamCaptures()[2]).value.cl_uintVal);
1425 }
1426 break;
1427 case ParamType::Tcl_context:
1428 if (call.entryPoint == EntryPoint::CLCreateContext ||
1429 call.entryPoint == EntryPoint::CLCreateContextFromType)
1430 {
1431 if ((getIndex(¶m.value.cl_contextVal) + 1) * sizeof(cl_context) >
1432 mMaxCLParamsSize[param.type])
1433 {
1434 mMaxCLParamsSize[param.type] =
1435 (uint32_t)((getIndex(¶m.value.cl_contextVal) + 1) * sizeof(cl_context));
1436 }
1437 addCLResetObj(param);
1438 }
1439 break;
1440 case ParamType::Tcl_command_queue:
1441 if (call.entryPoint == EntryPoint::CLCreateCommandQueueWithProperties ||
1442 call.entryPoint == EntryPoint::CLCreateCommandQueue)
1443 {
1444 if ((getIndex(¶m.value.cl_command_queueVal) + 1) * sizeof(cl_command_queue) >
1445 mMaxCLParamsSize[param.type])
1446 {
1447 mMaxCLParamsSize[param.type] =
1448 (uint32_t)((getIndex(¶m.value.cl_command_queueVal) + 1) *
1449 sizeof(cl_command_queue));
1450 }
1451 addCLResetObj(param);
1452 }
1453 break;
1454 case ParamType::Tcl_mem:
1455 if (call.entryPoint == EntryPoint::CLCreateBufferWithProperties ||
1456 call.entryPoint == EntryPoint::CLCreateBuffer ||
1457 call.entryPoint == EntryPoint::CLCreateSubBuffer ||
1458 call.entryPoint == EntryPoint::CLCreateImageWithProperties ||
1459 call.entryPoint == EntryPoint::CLCreateImage ||
1460 call.entryPoint == EntryPoint::CLCreateImage2D ||
1461 call.entryPoint == EntryPoint::CLCreateImage3D)
1462 {
1463 if ((getIndex(¶m.value.cl_memVal) + 1) * sizeof(cl_mem) >
1464 mMaxCLParamsSize[param.type])
1465 {
1466 mMaxCLParamsSize[param.type] =
1467 (uint32_t)((getIndex(¶m.value.cl_memVal) + 1) * sizeof(cl_mem));
1468 }
1469 addCLResetObj(param);
1470 }
1471 break;
1472 case ParamType::Tcl_eventPointer:
1473 {
1474 if (param.value.cl_eventVal)
1475 {
1476 if ((getIndex(¶m.value.cl_eventVal) + 1) * sizeof(cl_event) >
1477 mMaxCLParamsSize[param.type])
1478 {
1479 mMaxCLParamsSize[param.type] =
1480 (uint32_t)((getIndex(¶m.value.cl_eventVal) + 1) * sizeof(cl_event));
1481 }
1482 angle::ParamCapture eventParam("event", angle::ParamType::Tcl_event);
1483 InitParamValue(angle::ParamType::Tcl_event, param.value.cl_eventVal,
1484 &eventParam.value);
1485 addCLResetObj(eventParam);
1486 }
1487 break;
1488 }
1489 case ParamType::Tcl_program:
1490 if (call.entryPoint == EntryPoint::CLCreateProgramWithSource ||
1491 call.entryPoint == EntryPoint::CLCreateProgramWithBinary ||
1492 call.entryPoint == EntryPoint::CLCreateProgramWithBuiltInKernels ||
1493 call.entryPoint == EntryPoint::CLLinkProgram ||
1494 call.entryPoint == EntryPoint::CLCreateProgramWithIL)
1495 {
1496 if ((getIndex(¶m.value.cl_programVal) + 1) * sizeof(cl_program) >
1497 mMaxCLParamsSize[param.type])
1498 {
1499 mMaxCLParamsSize[param.type] =
1500 (uint32_t)((getIndex(¶m.value.cl_programVal) + 1) * sizeof(cl_program));
1501 }
1502 addCLResetObj(param);
1503 }
1504 break;
1505 case ParamType::Tcl_kernel:
1506 if (call.entryPoint == EntryPoint::CLCreateKernel ||
1507 call.entryPoint == EntryPoint::CLCloneKernel)
1508 {
1509 if ((getIndex(¶m.value.cl_kernelVal) + 1) * sizeof(cl_kernel) >
1510 mMaxCLParamsSize[param.type])
1511 {
1512 mMaxCLParamsSize[param.type] =
1513 (uint32_t)((getIndex(¶m.value.cl_kernelVal) + 1) * sizeof(cl_kernel));
1514 }
1515 addCLResetObj(param);
1516 }
1517 break;
1518 case ParamType::Tcl_sampler:
1519 if (call.entryPoint == EntryPoint::CLCreateSampler ||
1520 call.entryPoint == EntryPoint::CLCreateSamplerWithProperties)
1521 {
1522 if ((getIndex(¶m.value.cl_samplerVal) + 1) * sizeof(cl_sampler) >
1523 mMaxCLParamsSize[param.type])
1524 {
1525 mMaxCLParamsSize[param.type] =
1526 (uint32_t)((getIndex(¶m.value.cl_samplerVal) + 1) * sizeof(cl_sampler));
1527 }
1528 addCLResetObj(param);
1529 }
1530 break;
1531 case ParamType::TvoidPointer:
1532 if (call.entryPoint == EntryPoint::CLEnqueueMapImage ||
1533 call.entryPoint == EntryPoint::CLEnqueueMapBuffer)
1534 {
1535 mMaxCLParamsSize[param.type] += sizeof(void *);
1536 }
1537 break;
1538 default:
1539 break;
1540 }
1541 }
1542
updateResourceCountsFromCallCaptureCL(const CallCapture & call)1543 void FrameCaptureShared::updateResourceCountsFromCallCaptureCL(const CallCapture &call)
1544 {
1545 for (const ParamCapture ¶m : call.params.getParamCaptures())
1546 {
1547 updateResourceCountsFromParamCaptureCL(param, call);
1548 }
1549
1550 // Update resource IDs in the return value.
1551 switch (call.entryPoint)
1552 {
1553 case EntryPoint::CLCreateContext:
1554 case EntryPoint::CLCreateContextFromType:
1555 setIndex(&call.params.getReturnValue().value.cl_contextVal);
1556 updateResourceCountsFromParamCaptureCL(call.params.getReturnValue(), call);
1557 break;
1558 case EntryPoint::CLCreateBuffer:
1559 case EntryPoint::CLCreateBufferWithProperties:
1560 case EntryPoint::CLCreateSubBuffer:
1561 case EntryPoint::CLCreateImageWithProperties:
1562 case EntryPoint::CLCreateImage:
1563 case EntryPoint::CLCreateImage2D:
1564 case EntryPoint::CLCreateImage3D:
1565 case EntryPoint::CLCreatePipe:
1566 setIndex(&call.params.getReturnValue().value.cl_memVal);
1567 updateResourceCountsFromParamCaptureCL(call.params.getReturnValue(), call);
1568 break;
1569 case EntryPoint::CLCreateSampler:
1570 case EntryPoint::CLCreateSamplerWithProperties:
1571 setIndex(&call.params.getReturnValue().value.cl_samplerVal);
1572 updateResourceCountsFromParamCaptureCL(call.params.getReturnValue(), call);
1573 break;
1574 case EntryPoint::CLCreateCommandQueue:
1575 case EntryPoint::CLCreateCommandQueueWithProperties:
1576 setIndex(&call.params.getReturnValue().value.cl_command_queueVal);
1577 updateResourceCountsFromParamCaptureCL(call.params.getReturnValue(), call);
1578 break;
1579 case EntryPoint::CLCreateProgramWithSource:
1580 case EntryPoint::CLCreateProgramWithBinary:
1581 case EntryPoint::CLCreateProgramWithBuiltInKernels:
1582 case EntryPoint::CLLinkProgram:
1583 case EntryPoint::CLCreateProgramWithIL:
1584 setIndex(&call.params.getReturnValue().value.cl_programVal);
1585 updateResourceCountsFromParamCaptureCL(call.params.getReturnValue(), call);
1586 break;
1587 case EntryPoint::CLCreateKernel:
1588 case EntryPoint::CLCloneKernel:
1589 setIndex(&call.params.getReturnValue().value.cl_kernelVal);
1590 updateResourceCountsFromParamCaptureCL(call.params.getReturnValue(), call);
1591 break;
1592 case EntryPoint::CLEnqueueMapBuffer:
1593 case EntryPoint::CLEnqueueMapImage:
1594 case EntryPoint::CLSVMAlloc:
1595 if (call.params.getReturnValue().value.voidPointerVal)
1596 {
1597 setCLVoidIndex(call.params.getReturnValue().value.voidPointerVal);
1598 updateResourceCountsFromParamCaptureCL(call.params.getReturnValue(), call);
1599 }
1600 break;
1601 case EntryPoint::CLCreateUserEvent:
1602 setIndex(&call.params.getReturnValue().value.cl_eventVal);
1603 break;
1604 case EntryPoint::CLReleaseDevice:
1605 case EntryPoint::CLReleaseCommandQueue:
1606 case EntryPoint::CLReleaseContext:
1607 case EntryPoint::CLReleaseEvent:
1608 case EntryPoint::CLReleaseKernel:
1609 case EntryPoint::CLReleaseMemObject:
1610 case EntryPoint::CLReleaseProgram:
1611 case EntryPoint::CLReleaseSampler:
1612 removeCLResetObj(call.params.getParamCaptures()[0]);
1613 break;
1614 default:
1615 break;
1616 }
1617 }
1618
captureCLCall(CallCapture && inCall,bool isCallValid)1619 void FrameCaptureShared::captureCLCall(CallCapture &&inCall, bool isCallValid)
1620 {
1621 if (!mCallCaptured)
1622 {
1623 mReplayWriter.captureAPI = CaptureAPI::CL;
1624 mBinaryData.clear();
1625 mCallCaptured = true;
1626 std::atexit(onCLProgramEnd);
1627 }
1628
1629 if (mFrameIndex <= mCaptureEndFrame)
1630 {
1631
1632 // Keep track of return values from OpenCL calls
1633 updateResourceCountsFromCallCaptureCL(inCall);
1634
1635 // Set to true if the call signifies the end of a frame
1636 // ex: clEnqueueNDRangeKernel
1637 bool frameEnd = false;
1638
1639 // Covers pre call updates, like updating the read buffer size
1640 maybeCapturePreCallUpdatesCL(inCall);
1641
1642 // If it's an unnecessary call for replay (ex: clGetDeviceInfo)
1643 if (mCLOptionalCalls.find(inCall.entryPoint) == mCLOptionalCalls.end())
1644 {
1645 if (mCLEndFrameCalls.find(inCall.entryPoint) != mCLEndFrameCalls.end())
1646 {
1647 frameEnd = true;
1648 }
1649
1650 mFrameCalls.emplace_back(std::move(inCall));
1651 }
1652 else
1653 {
1654 saveCLGetInfo(inCall);
1655 return;
1656 }
1657
1658 // For kernel argument memory snapshots
1659 maybeCapturePostCallUpdatesCL();
1660 if (mFrameIndex >= mCaptureStartFrame ||
1661 (mFrameIndex + 1 == mCaptureStartFrame && frameEnd))
1662 {
1663 // Maybe add clEnqueueWrite* or memcpy for memory snapshots
1664 captureUpdateCLObjs(&mFrameCalls);
1665 }
1666
1667 if (frameEnd && mFrameIndex >= mCaptureStartFrame)
1668 {
1669 mActiveFrameIndices.push_back(mFrameIndex);
1670 writeMainContextCppReplayCL();
1671 if (mFrameIndex == mCaptureEndFrame)
1672 {
1673 writeCppReplayIndexFilesCL();
1674 SaveBinaryData(mCompression, mOutDirectory, kNoContextId, mCaptureLabel,
1675 mBinaryData);
1676 }
1677 reset();
1678 }
1679
1680 if (frameEnd)
1681 {
1682 if (mFrameIndex == (mCaptureStartFrame == 0 ? 0 : mCaptureStartFrame - 1))
1683 {
1684 mCLSetupCalls = std::move(mFrameCalls);
1685 }
1686 ++mFrameIndex;
1687 }
1688 }
1689 }
1690
maybeCapturePostCallUpdatesCL()1691 void FrameCaptureShared::maybeCapturePostCallUpdatesCL()
1692 {
1693 CallCapture &lastCall = mFrameCalls.back();
1694 switch (lastCall.entryPoint)
1695 {
1696 case EntryPoint::CLEnqueueMapBuffer:
1697 {
1698 // Recreate the map call to store in the mCLMapCall unordered_map
1699 // so later upon the unmap call, the original map data will be available
1700 cl_command_queue command_queue =
1701 lastCall.params.getParam("command_queue", ParamType::Tcl_command_queue, 0)
1702 .value.cl_command_queueVal;
1703 cl_mem buffer =
1704 lastCall.params.getParam("buffer", ParamType::Tcl_mem, 1).value.cl_memVal;
1705 cl_bool blocking_map =
1706 lastCall.params.getParam("blocking_map", ParamType::Tcl_bool, 2).value.cl_boolVal;
1707 cl::MapFlags map_flags =
1708 lastCall.params.getParam("map_flagsPacked", ParamType::TMapFlags, 3)
1709 .value.MapFlagsVal;
1710 size_t offset =
1711 lastCall.params.getParam("offset", ParamType::Tsize_t, 4).value.size_tVal;
1712 size_t size = lastCall.params.getParam("size", ParamType::Tsize_t, 5).value.size_tVal;
1713
1714 mResourceTrackerCL.mCLMapCall.emplace(
1715 lastCall.params.getReturnValue().value.voidPointerVal,
1716 CaptureEnqueueMapBuffer(true, command_queue, buffer, blocking_map, map_flags,
1717 offset, size, 0, nullptr, nullptr, nullptr, nullptr));
1718 break;
1719 }
1720 case EntryPoint::CLEnqueueMapImage:
1721 {
1722 // Recreate the map call to store in the mCLMapCall unordered_map
1723 // so later upon the unmap call, the original map data will be available
1724 cl_command_queue command_queue =
1725 lastCall.params.getParam("command_queue", ParamType::Tcl_command_queue, 0)
1726 .value.cl_command_queueVal;
1727 cl_mem image = lastCall.params.getParam("image", ParamType::Tcl_mem, 1).value.cl_memVal;
1728 cl_bool blocking_map =
1729 lastCall.params.getParam("blocking_map", ParamType::Tcl_bool, 2).value.cl_boolVal;
1730 cl::MapFlags map_flags =
1731 lastCall.params.getParam("map_flagsPacked", ParamType::TMapFlags, 3)
1732 .value.MapFlagsVal;
1733 const size_t *origin =
1734 lastCall.params.getParam("origin", ParamType::Tsize_tConstPointer, 4)
1735 .value.size_tConstPointerVal;
1736 const size_t *region =
1737 lastCall.params.getParam("region", ParamType::Tsize_tConstPointer, 5)
1738 .value.size_tConstPointerVal;
1739 size_t *image_row_pitch =
1740 lastCall.params.getParam("image_row_pitch", ParamType::Tsize_tPointer, 6)
1741 .value.size_tPointerVal;
1742 size_t *image_slice_pitch =
1743 lastCall.params.getParam("image_slice_pitch", ParamType::Tsize_tPointer, 7)
1744 .value.size_tPointerVal;
1745
1746 mResourceTrackerCL.mCLMapCall.emplace(
1747 lastCall.params.getReturnValue().value.voidPointerVal,
1748 CaptureEnqueueMapImage(true, command_queue, image, blocking_map, map_flags, origin,
1749 region, image_row_pitch, image_slice_pitch, 0, nullptr,
1750 nullptr, nullptr, nullptr));
1751 mResourceTrackerCL.mCLMapCall.at(lastCall.params.getReturnValue().value.voidPointerVal)
1752 .params.setValueParamAtIndex("image_row_pitch", ParamType::Tsize_t,
1753 *image_row_pitch, 6);
1754 mResourceTrackerCL.mCLMapCall.at(lastCall.params.getReturnValue().value.voidPointerVal)
1755 .params.setValueParamAtIndex("image_slice_pitch", ParamType::Tsize_t,
1756 image_slice_pitch == nullptr ? 0 : *image_slice_pitch,
1757 7);
1758 break;
1759 }
1760 case EntryPoint::CLEnqueueUnmapMemObject:
1761 {
1762 if (mFrameIndex >= mCaptureStartFrame)
1763 {
1764 // Mark as dirty
1765 cl_mem *mem =
1766 &lastCall.params.getParam("memobj", ParamType::Tcl_mem, 1).value.cl_memVal;
1767 mResourceTrackerCL.mCLCurrentCommandQueue =
1768 lastCall.params.getParam("command_queue", ParamType::Tcl_command_queue, 0)
1769 .value.cl_command_queueVal;
1770 CallCapture *mapCall = &mResourceTrackerCL.mCLMapCall.at(
1771 lastCall.params.getParam("mapped_ptr", ParamType::TvoidPointer, 2)
1772 .value.voidPointerVal);
1773 auto it = std::find(mResourceTrackerCL.mCLDirtyMem.begin(),
1774 mResourceTrackerCL.mCLDirtyMem.end(), *mem);
1775 if (it == mResourceTrackerCL.mCLDirtyMem.end() &&
1776 mapCall->params.getParam("map_flagsPacked", ParamType::TMapFlags, 3)
1777 .value.MapFlagsVal.mask(CL_MAP_WRITE |
1778 CL_MAP_WRITE_INVALIDATE_REGION) != 0u)
1779 {
1780 mResourceTrackerCL.mCLDirtyMem.push_back(*mem);
1781 }
1782 }
1783 break;
1784 }
1785 case EntryPoint::CLEnqueueSVMUnmap:
1786 {
1787 // Mark as dirty
1788 void *svm = &lastCall.params.getParam("svm_ptr", ParamType::TvoidPointer, 1)
1789 .value.voidPointerVal;
1790 mResourceTrackerCL.mCLCurrentCommandQueue =
1791 lastCall.params.getParam("command_queue", ParamType::Tcl_command_queue, 0)
1792 .value.cl_command_queueVal;
1793 mResourceTrackerCL.mCLDirtySVM.push_back(svm);
1794 break;
1795 }
1796 default:
1797 break;
1798 }
1799
1800 // OpenCL calls that come before the starting frame
1801 if (mFrameIndex < mCaptureStartFrame)
1802 {
1803 std::unordered_map<cl_program, std::vector<cl_kernel>> &mCLProgramToKernels =
1804 mResourceTrackerCL.mCLProgramToKernels;
1805 switch (lastCall.entryPoint)
1806 {
1807 // There should be no unnecessary enqueue functions prior to the starting frame.
1808 // captureUpdateCLObjs accounts for it by dynamically adding
1809 // CLEnqueueWriteBuffer/CLEnqueueWriteImage to ensure the cl_mem objects
1810 // have the needed info upon replay
1811 case EntryPoint::CLEnqueueNDRangeKernel:
1812 case EntryPoint::CLEnqueueNativeKernel:
1813 case EntryPoint::CLEnqueueTask:
1814 case EntryPoint::CLEnqueueReadBuffer:
1815 case EntryPoint::CLEnqueueWriteBuffer:
1816 case EntryPoint::CLEnqueueReadBufferRect:
1817 case EntryPoint::CLEnqueueWriteBufferRect:
1818 case EntryPoint::CLEnqueueReadImage:
1819 case EntryPoint::CLEnqueueWriteImage:
1820 case EntryPoint::CLEnqueueCopyBuffer:
1821 case EntryPoint::CLEnqueueCopyBufferRect:
1822 case EntryPoint::CLEnqueueCopyImage:
1823 case EntryPoint::CLEnqueueCopyBufferToImage:
1824 case EntryPoint::CLEnqueueCopyImageToBuffer:
1825 case EntryPoint::CLEnqueueFillBuffer:
1826 case EntryPoint::CLEnqueueFillImage:
1827 case EntryPoint::CLEnqueueWaitForEvents:
1828 case EntryPoint::CLEnqueueMarkerWithWaitList:
1829 case EntryPoint::CLEnqueueBarrierWithWaitList:
1830 case EntryPoint::CLEnqueueBarrier:
1831 case EntryPoint::CLEnqueueMarker:
1832 case EntryPoint::CLEnqueueMigrateMemObjects:
1833 case EntryPoint::CLEnqueueSVMMemcpy:
1834 case EntryPoint::CLEnqueueSVMMemFill:
1835 case EntryPoint::CLEnqueueSVMMigrateMem:
1836 {
1837 size_t index = mFrameCalls.size() - 1;
1838 removeCLCall(&mFrameCalls, index);
1839 break;
1840 }
1841 case EntryPoint::CLCreateBuffer:
1842 case EntryPoint::CLCreateBufferWithProperties:
1843 case EntryPoint::CLCreateImage:
1844 case EntryPoint::CLCreateImageWithProperties:
1845 case EntryPoint::CLCreateImage2D:
1846 case EntryPoint::CLCreateImage3D:
1847 case EntryPoint::CLCreatePipe:
1848 case EntryPoint::CLCreateSubBuffer:
1849 {
1850 const cl_mem *newBuff = &lastCall.params.getReturnValue().value.cl_memVal;
1851
1852 // Set the parent
1853 if (lastCall.entryPoint == EntryPoint::CLCreateSubBuffer)
1854 {
1855 cl_mem parent =
1856 lastCall.params.getParam("buffer", ParamType::Tcl_mem, 0).value.cl_memVal;
1857 mResourceTrackerCL.mCLSubBufferToParent[*newBuff] = parent;
1858 }
1859
1860 // Implicit retain
1861 (*newBuff)->cast<cl::Memory>().retain();
1862
1863 // Add buffer as tracked
1864 trackCLMemUpdate(newBuff, true);
1865 break;
1866 }
1867 case EntryPoint::CLReleaseMemObject:
1868 {
1869 // Potentially remove buffer/image (and potentially parents) as tracked
1870 trackCLMemUpdate(
1871 &lastCall.params.getParam("memobj", ParamType::Tcl_mem, 0).value.cl_memVal,
1872 false);
1873 break;
1874 }
1875 case EntryPoint::CLCreateCommandQueue:
1876 case EntryPoint::CLCreateCommandQueueWithProperties:
1877 {
1878 mResourceTrackerCL.mCLCurrentCommandQueue =
1879 lastCall.params.getReturnValue().value.cl_command_queueVal;
1880 break;
1881 }
1882 case EntryPoint::CLCreateProgramWithSource:
1883 case EntryPoint::CLCreateProgramWithBinary:
1884 case EntryPoint::CLCreateProgramWithBuiltInKernels:
1885 case EntryPoint::CLCreateProgramWithIL:
1886 {
1887 mCLProgramToKernels[lastCall.params.getReturnValue().value.cl_programVal] =
1888 std::vector<cl_kernel>();
1889 trackCLProgramUpdate(&lastCall.params.getReturnValue().value.cl_programVal, true, 0,
1890 nullptr);
1891 break;
1892 }
1893 case EntryPoint::CLRetainProgram:
1894 {
1895 trackCLProgramUpdate(&lastCall.params.getParam("program", ParamType::Tcl_program, 0)
1896 .value.cl_programVal,
1897 true, 0, nullptr);
1898 break;
1899 }
1900 case EntryPoint::CLCompileProgram:
1901 {
1902 const cl_program *program =
1903 &lastCall.params.getParam("program", ParamType::Tcl_program, 0)
1904 .value.cl_programVal;
1905 trackCLProgramUpdate(
1906 program, true,
1907 lastCall.params.getParam("num_input_headers", ParamType::Tcl_uint, 4)
1908 .value.cl_uintVal,
1909 lastCall.params.getParam("input_headers", ParamType::Tcl_programConstPointer, 5)
1910 .value.cl_programConstPointerVal);
1911 break;
1912 }
1913 case EntryPoint::CLLinkProgram:
1914 {
1915 const cl_program *program = &lastCall.params.getReturnValue().value.cl_programVal;
1916 mCLProgramToKernels[*program] = std::vector<cl_kernel>();
1917 trackCLProgramUpdate(
1918 program, true,
1919 lastCall.params.getParam("num_input_programs", ParamType::Tcl_uint, 4)
1920 .value.cl_uintVal,
1921 lastCall.params
1922 .getParam("input_programs", ParamType::Tcl_programConstPointer, 5)
1923 .value.cl_programConstPointerVal);
1924 break;
1925 }
1926 case EntryPoint::CLReleaseProgram:
1927 {
1928 trackCLProgramUpdate(&lastCall.params.getParam("program", ParamType::Tcl_program, 0)
1929 .value.cl_programVal,
1930 false, 0, nullptr);
1931 break;
1932 }
1933 case EntryPoint::CLCreateKernel:
1934 {
1935 cl_program *program =
1936 &lastCall.params.getParam("program", ParamType::Tcl_program, 0)
1937 .value.cl_programVal;
1938 const cl_kernel *kernel = &lastCall.params.getReturnValue().value.cl_kernelVal;
1939 mCLProgramToKernels[*program].push_back(*kernel);
1940 mResourceTrackerCL.mCLKernelToProgram[*kernel] = *program;
1941 break;
1942 }
1943 case EntryPoint::CLCloneKernel:
1944 {
1945 cl_program *program =
1946 &mResourceTrackerCL.mCLKernelToProgram[lastCall.params
1947 .getParam("source_kernel",
1948 ParamType::Tcl_kernel, 0)
1949 .value.cl_kernelVal];
1950 const cl_kernel *kernel = &lastCall.params.getReturnValue().value.cl_kernelVal;
1951 mCLProgramToKernels[*program].push_back(*kernel);
1952 mResourceTrackerCL.mCLKernelToProgram[*kernel] = *program;
1953 break;
1954 }
1955 case EntryPoint::CLSVMAlloc:
1956 {
1957 void *svm = lastCall.params.getReturnValue().value.voidPointerVal;
1958
1959 // Potentially mark as dirty
1960 auto it = std::find(mResourceTrackerCL.mCLDirtySVM.begin(),
1961 mResourceTrackerCL.mCLDirtySVM.end(), svm);
1962 if (it == mResourceTrackerCL.mCLDirtySVM.end())
1963 {
1964 mResourceTrackerCL.mCLDirtySVM.push_back(svm);
1965 }
1966 break;
1967 }
1968 case EntryPoint::CLSVMFree:
1969 {
1970 void *svm = lastCall.params.getParam("svm_pointer", ParamType::TvoidPointer, 1)
1971 .value.voidPointerVal;
1972 auto it = std::find(mResourceTrackerCL.mCLDirtySVM.begin(),
1973 mResourceTrackerCL.mCLDirtySVM.end(), svm);
1974 if (it != mResourceTrackerCL.mCLDirtySVM.end())
1975 {
1976 mResourceTrackerCL.mCLDirtySVM.erase(it);
1977 }
1978 break;
1979 }
1980 case EntryPoint::CLEnqueueSVMFree:
1981 {
1982 for (cl_uint svmIndex = 0;
1983 svmIndex < lastCall.params.getParam("num_svm_pointers", ParamType::Tcl_uint, 1)
1984 .value.cl_uintVal;
1985 ++svmIndex)
1986 {
1987 void *svm =
1988 lastCall.params.getParam("svm_pointers", ParamType::TvoidPointerPointer, 1)
1989 .value.voidPointerPointerVal[svmIndex];
1990 auto it = std::find(mResourceTrackerCL.mCLDirtySVM.begin(),
1991 mResourceTrackerCL.mCLDirtySVM.end(), svm);
1992 if (it != mResourceTrackerCL.mCLDirtySVM.end())
1993 {
1994 mResourceTrackerCL.mCLDirtySVM.erase(it);
1995 }
1996 }
1997 break;
1998 }
1999 default:
2000 break;
2001 }
2002 }
2003 }
2004
onCLProgramEnd()2005 void FrameCaptureShared::onCLProgramEnd()
2006 {
2007 if (cl::Platform::GetDefault()->getFrameCaptureShared()->onEndCLCapture())
2008 {
2009 delete cl::Platform::GetDefault()->getFrameCaptureShared();
2010 }
2011 }
2012
onEndCLCapture()2013 bool FrameCaptureShared::onEndCLCapture()
2014 {
2015 if (mFrameIndex >= mCaptureStartFrame && mFrameIndex <= mCaptureEndFrame)
2016 {
2017 mActiveFrameIndices.push_back(mFrameIndex);
2018 mCaptureEndFrame = mFrameIndex;
2019 writeMainContextCppReplayCL();
2020 writeCppReplayIndexFilesCL();
2021 SaveBinaryData(mCompression, mOutDirectory, kNoContextId, mCaptureLabel, mBinaryData);
2022 return true;
2023 }
2024 return false;
2025 }
2026
2027 ResourceTrackerCL::ResourceTrackerCL() = default;
2028
2029 ResourceTrackerCL::~ResourceTrackerCL() = default;
2030
setCLPlatformIndices(cl_platform_id * platforms,size_t numPlatforms)2031 void FrameCaptureShared::setCLPlatformIndices(cl_platform_id *platforms, size_t numPlatforms)
2032 {
2033 for (uint32_t i = 0; i < numPlatforms; ++i)
2034 {
2035 setIndex(&platforms[i]);
2036 }
2037 }
2038
setCLDeviceIndices(cl_device_id * devices,size_t numDevices)2039 void FrameCaptureShared::setCLDeviceIndices(cl_device_id *devices, size_t numDevices)
2040 {
2041 for (uint32_t i = 0; i < numDevices; ++i)
2042 {
2043 setIndex(&devices[i]);
2044 }
2045 }
2046
getCLVoidIndex(const void * v)2047 size_t FrameCaptureShared::getCLVoidIndex(const void *v)
2048 {
2049 if (mResourceTrackerCL.mCLVoidIndices.find(v) == mResourceTrackerCL.mCLVoidIndices.end())
2050 {
2051 return SIZE_MAX;
2052 }
2053 return mResourceTrackerCL.mCLVoidIndices[v];
2054 }
2055
setCLVoidIndex(const void * v)2056 void FrameCaptureShared::setCLVoidIndex(const void *v)
2057 {
2058 if (mResourceTrackerCL.mCLVoidIndices.find(v) == mResourceTrackerCL.mCLVoidIndices.end())
2059 {
2060 size_t tempSize = mResourceTrackerCL.mCLVoidIndices.size();
2061 mResourceTrackerCL.mCLVoidIndices[v] = tempSize;
2062 }
2063 }
2064
setCLVoidVectorIndex(const void * pointers[],size_t numPointers,const angle::ParamCapture * paramCaptureKey)2065 void FrameCaptureShared::setCLVoidVectorIndex(const void *pointers[],
2066 size_t numPointers,
2067 const angle::ParamCapture *paramCaptureKey)
2068 {
2069 mResourceTrackerCL.mCLParamIDToIndexVector[paramCaptureKey->uniqueID] = std::vector<size_t>();
2070 for (size_t i = 0; i < numPointers; ++i)
2071 {
2072 mResourceTrackerCL.mCLParamIDToIndexVector[paramCaptureKey->uniqueID].push_back(
2073 getCLVoidIndex(pointers[i]));
2074 }
2075 }
2076
setOffsetsVector(const void * args,const void ** argsLocations,size_t numLocations,const angle::ParamCapture * paramCaptureKey)2077 void FrameCaptureShared::setOffsetsVector(const void *args,
2078 const void **argsLocations,
2079 size_t numLocations,
2080 const angle::ParamCapture *paramCaptureKey)
2081 {
2082 mResourceTrackerCL.mCLParamIDToIndexVector[paramCaptureKey->uniqueID] = std::vector<size_t>();
2083 for (size_t i = 0; i < numLocations; ++i)
2084 {
2085 mResourceTrackerCL.mCLParamIDToIndexVector[paramCaptureKey->uniqueID].push_back(
2086 (char *)argsLocations[i] - (char *)args);
2087 }
2088 }
2089
getCLObjVector(const angle::ParamCapture * paramCaptureKey)2090 std::vector<size_t> FrameCaptureShared::getCLObjVector(const angle::ParamCapture *paramCaptureKey)
2091 {
2092 if (mResourceTrackerCL.mCLParamIDToIndexVector.find(paramCaptureKey->uniqueID) !=
2093 mResourceTrackerCL.mCLParamIDToIndexVector.end())
2094 {
2095 return mResourceTrackerCL.mCLParamIDToIndexVector[paramCaptureKey->uniqueID];
2096 }
2097 return std::vector<size_t>();
2098 }
2099
2100 template <typename T>
getMap()2101 std::unordered_map<T, size_t> &FrameCaptureShared::getMap()
2102 {
2103 ASSERT(false);
2104 return std::unordered_map<T, size_t>();
2105 }
2106 template <>
getMap()2107 std::unordered_map<cl_platform_id, size_t> &FrameCaptureShared::getMap<cl_platform_id>()
2108 {
2109 return mResourceTrackerCL.mCLPlatformIDIndices;
2110 }
2111 template <>
getMap()2112 std::unordered_map<cl_device_id, size_t> &FrameCaptureShared::getMap<cl_device_id>()
2113 {
2114 return mResourceTrackerCL.mCLDeviceIDIndices;
2115 }
2116 template <>
getMap()2117 std::unordered_map<cl_context, size_t> &FrameCaptureShared::getMap<cl_context>()
2118 {
2119 return mResourceTrackerCL.mCLContextIndices;
2120 }
2121 template <>
getMap()2122 std::unordered_map<cl_event, size_t> &FrameCaptureShared::getMap<cl_event>()
2123 {
2124 return mResourceTrackerCL.mCLEventsIndices;
2125 }
2126 template <>
getMap()2127 std::unordered_map<cl_command_queue, size_t> &FrameCaptureShared::getMap<cl_command_queue>()
2128 {
2129 return mResourceTrackerCL.mCLCommandQueueIndices;
2130 }
2131 template <>
getMap()2132 std::unordered_map<cl_mem, size_t> &FrameCaptureShared::getMap<cl_mem>()
2133 {
2134 return mResourceTrackerCL.mCLMemIndices;
2135 }
2136 template <>
getMap()2137 std::unordered_map<cl_sampler, size_t> &FrameCaptureShared::getMap<cl_sampler>()
2138 {
2139 return mResourceTrackerCL.mCLSamplerIndices;
2140 }
2141 template <>
getMap()2142 std::unordered_map<cl_program, size_t> &FrameCaptureShared::getMap<cl_program>()
2143 {
2144 return mResourceTrackerCL.mCLProgramIndices;
2145 }
2146 template <>
getMap()2147 std::unordered_map<cl_kernel, size_t> &FrameCaptureShared::getMap<cl_kernel>()
2148 {
2149 return mResourceTrackerCL.mCLKernelIndices;
2150 }
2151
writeJSONCL()2152 void FrameCaptureShared::writeJSONCL()
2153 {
2154
2155 JsonSerializer json;
2156 json.startGroup("TraceMetadata");
2157 json.addBool("IsBinaryDataCompressed", mCompression);
2158 json.addScalar("CaptureRevision", GetANGLERevision());
2159 json.addScalar("FrameStart", mCaptureStartFrame);
2160 json.addScalar("FrameEnd", mFrameIndex);
2161 json.addBool("IsOpenCL", true);
2162 json.endGroup();
2163
2164 {
2165 const std::vector<std::string> &traceFiles = mReplayWriter.getAndResetWrittenFiles();
2166 json.addVectorOfStrings("TraceFiles", traceFiles);
2167 }
2168
2169 {
2170 std::stringstream jsonFileNameStream;
2171 jsonFileNameStream << mOutDirectory << FmtCapturePrefix(kNoContextId, mCaptureLabel)
2172 << ".json";
2173 std::string jsonFileName = jsonFileNameStream.str();
2174
2175 SaveFileHelper saveData(jsonFileName);
2176 saveData.write(reinterpret_cast<const uint8_t *>(json.data()), json.length());
2177 }
2178 }
2179
saveCLGetInfo(const CallCapture & call)2180 void FrameCaptureShared::saveCLGetInfo(const CallCapture &call)
2181 {
2182 std::string prevCall = "";
2183 size_t size;
2184 std::ostringstream objectStream;
2185 std::string clObject;
2186 JsonSerializer json;
2187
2188 json.startGroup(call.name());
2189
2190 // Below ONLY for clGetSupportedImageFormats
2191 if (call.entryPoint == EntryPoint::CLGetSupportedImageFormats)
2192 {
2193 const cl_image_format *data =
2194 call.params.getParam("image_formats", ParamType::Tcl_image_formatPointer, 4)
2195 .value.cl_image_formatPointerVal;
2196 if (!data)
2197 {
2198 return;
2199 }
2200 size_t *sizePointer =
2201 call.params.getParam("num_image_formats", ParamType::Tcl_uintPointer, 5)
2202 .value.size_tPointerVal;
2203 if (!sizePointer)
2204 {
2205 size =
2206 call.params.getParam("param_value_size", ParamType::Tcl_uint, 3).value.cl_uintVal;
2207 }
2208 else
2209 {
2210 size = *sizePointer;
2211 }
2212
2213 cl_context context = call.params.getParamCaptures()[0].value.cl_contextVal;
2214 objectStream << static_cast<void *>(context);
2215 clObject = objectStream.str();
2216 json.addString("context", clObject);
2217 json.addScalar("flags", call.params.getParamCaptures()[1].value.MemFlagsVal.get());
2218
2219 cl::MemObjectType imageType =
2220 call.params.getParam("image_typePacked", ParamType::TMemObjectType, 2)
2221 .value.MemObjectTypeVal;
2222 std::ostringstream oss;
2223 oss << imageType;
2224 std::string infoString = oss.str();
2225 json.startGroup(infoString);
2226 for (size_t j = 0; j < size; ++j)
2227 {
2228 std::ostringstream temp;
2229 temp << (j + 1);
2230 json.addScalar("image_channel_order" + temp.str(), data[j].image_channel_order);
2231 json.addScalar("image_channel_data_type" + temp.str(), data[j].image_channel_data_type);
2232 }
2233
2234 json.endGroup();
2235 json.endGroup();
2236 return;
2237 }
2238
2239 // Get the param_value and size
2240 bool offsetData = 0;
2241 switch (call.entryPoint)
2242 {
2243 case EntryPoint::CLGetProgramBuildInfo:
2244 case EntryPoint::CLGetKernelArgInfo:
2245 case EntryPoint::CLGetKernelWorkGroupInfo:
2246 {
2247 offsetData = 1;
2248 break;
2249 }
2250 default:
2251 break;
2252 }
2253
2254 const void *data = call.params.getParam("param_value", ParamType::TvoidPointer, 3 + offsetData)
2255 .value.voidPointerVal;
2256 if (!data)
2257 {
2258 return;
2259 }
2260 size_t *sizePointer =
2261 call.params.getParam("param_value_size_ret", ParamType::Tsize_tPointer, 4 + offsetData)
2262 .value.size_tPointerVal;
2263 if (!sizePointer)
2264 {
2265 size = call.params.getParam("param_value_size", ParamType::Tsize_t, 2 + offsetData)
2266 .value.size_tVal;
2267 }
2268 else
2269 {
2270 size = *sizePointer;
2271 }
2272
2273 // Get string representation of OpenCL object specified
2274 switch (call.entryPoint)
2275 {
2276 case EntryPoint::CLGetPlatformInfo:
2277 {
2278 cl_platform_id platform = call.params.getParamCaptures()[0].value.cl_platform_idVal;
2279 objectStream << static_cast<void *>(platform);
2280 break;
2281 }
2282 case EntryPoint::CLGetDeviceInfo:
2283 {
2284 cl_device_id device = call.params.getParamCaptures()[0].value.cl_device_idVal;
2285 objectStream << static_cast<void *>(device);
2286 break;
2287 }
2288 case EntryPoint::CLGetContextInfo:
2289 {
2290 cl_context context = call.params.getParamCaptures()[0].value.cl_contextVal;
2291 objectStream << static_cast<void *>(context);
2292 break;
2293 }
2294 case EntryPoint::CLGetCommandQueueInfo:
2295 {
2296 cl_command_queue commandQueue =
2297 call.params.getParamCaptures()[0].value.cl_command_queueVal;
2298 objectStream << static_cast<void *>(commandQueue);
2299 break;
2300 }
2301 case EntryPoint::CLGetProgramInfo:
2302 case EntryPoint::CLGetProgramBuildInfo:
2303 {
2304 cl_program program = call.params.getParamCaptures()[0].value.cl_programVal;
2305 objectStream << static_cast<void *>(program);
2306 break;
2307 }
2308 case EntryPoint::CLGetKernelInfo:
2309 case EntryPoint::CLGetKernelArgInfo:
2310 case EntryPoint::CLGetKernelWorkGroupInfo:
2311 {
2312 cl_kernel kernel = call.params.getParamCaptures()[0].value.cl_kernelVal;
2313 objectStream << static_cast<void *>(kernel);
2314 break;
2315 }
2316 case EntryPoint::CLGetEventInfo:
2317 case EntryPoint::CLGetEventProfilingInfo:
2318 {
2319 cl_event event = call.params.getParamCaptures()[0].value.cl_eventVal;
2320 objectStream << static_cast<void *>(event);
2321 break;
2322 }
2323 case EntryPoint::CLGetMemObjectInfo:
2324 case EntryPoint::CLGetImageInfo:
2325 {
2326 cl_mem mem = call.params.getParamCaptures()[0].value.cl_memVal;
2327 objectStream << static_cast<void *>(mem);
2328 break;
2329 }
2330 case EntryPoint::CLGetSamplerInfo:
2331 {
2332 cl_sampler sampler = call.params.getParamCaptures()[0].value.cl_samplerVal;
2333 objectStream << static_cast<void *>(sampler);
2334 break;
2335 }
2336 default:
2337 break;
2338 }
2339 clObject = objectStream.str();
2340
2341 // Go through the param_name options
2342 switch (call.entryPoint)
2343 {
2344 case EntryPoint::CLGetPlatformInfo:
2345 {
2346 cl::PlatformInfo info = call.params.getParamCaptures()[1].value.PlatformInfoVal;
2347 std::ostringstream oss;
2348 oss << info;
2349 std::string infoString = oss.str();
2350 json.addString("platform", clObject);
2351
2352 switch (ToCLenum(info))
2353 {
2354 case CL_PLATFORM_PROFILE:
2355 case CL_PLATFORM_VERSION:
2356 case CL_PLATFORM_NAME:
2357 case CL_PLATFORM_VENDOR:
2358 case CL_PLATFORM_EXTENSIONS:
2359 case CL_PLATFORM_ICD_SUFFIX_KHR:
2360 {
2361 json.addCString(infoString, static_cast<const char *>(data));
2362 break;
2363 }
2364 case CL_PLATFORM_EXTENSIONS_WITH_VERSION:
2365 {
2366 const cl_name_version *nameVersion = static_cast<const cl_name_version *>(data);
2367 json.startGroup(infoString);
2368 for (size_t j = 0; j < size / sizeof(cl_name_version); ++j)
2369 {
2370 json.addScalar(nameVersion[j].name, nameVersion[j].version);
2371 }
2372 json.endGroup();
2373 break;
2374 }
2375 case CL_PLATFORM_NUMERIC_VERSION:
2376 json.addScalar(infoString, *static_cast<const cl_uint *>(data));
2377 break;
2378 case CL_PLATFORM_HOST_TIMER_RESOLUTION:
2379 case CL_PLATFORM_COMMAND_BUFFER_CAPABILITIES_KHR:
2380 json.addScalar(infoString, *static_cast<const cl_ulong *>(data));
2381 break;
2382 case CL_PLATFORM_EXTERNAL_MEMORY_IMPORT_HANDLE_TYPES_KHR:
2383 case CL_PLATFORM_SEMAPHORE_TYPES_KHR:
2384 case CL_PLATFORM_SEMAPHORE_IMPORT_HANDLE_TYPES_KHR:
2385 case CL_PLATFORM_SEMAPHORE_EXPORT_HANDLE_TYPES_KHR:
2386 json.addVector(infoString,
2387 std::vector<cl_uint>((cl_uint *)data,
2388 (cl_uint *)data + size / sizeof(cl_uint)));
2389 break;
2390 default:
2391 // Not supported or cannot add to JSON file
2392 break;
2393 }
2394
2395 break;
2396 }
2397 case EntryPoint::CLGetDeviceInfo:
2398 {
2399 cl::DeviceInfo info = call.params.getParamCaptures()[1].value.DeviceInfoVal;
2400 std::ostringstream oss;
2401 oss << info;
2402 std::string infoString = oss.str();
2403 json.addString("device", clObject);
2404 switch (ToCLenum(info))
2405 {
2406 case CL_DEVICE_IL_VERSION:
2407 case CL_DEVICE_LATEST_CONFORMANCE_VERSION_PASSED:
2408 case CL_DEVICE_OPENCL_C_VERSION:
2409 case CL_DEVICE_EXTENSIONS:
2410 case CL_DEVICE_VERSION:
2411 case CL_DEVICE_PROFILE:
2412 case CL_DRIVER_VERSION:
2413 case CL_DEVICE_VENDOR:
2414 case CL_DEVICE_NAME:
2415 json.addCString(infoString, static_cast<const char *>(data));
2416 break;
2417 case CL_DEVICE_TYPE:
2418 case CL_DEVICE_MAX_MEM_ALLOC_SIZE:
2419 case CL_DEVICE_LOCAL_MEM_SIZE:
2420 case CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE:
2421 case CL_DEVICE_GLOBAL_MEM_SIZE:
2422 case CL_DEVICE_GLOBAL_MEM_CACHE_SIZE:
2423 case CL_DEVICE_HALF_FP_CONFIG:
2424 case CL_DEVICE_SINGLE_FP_CONFIG:
2425 case CL_DEVICE_DEVICE_ENQUEUE_CAPABILITIES:
2426 case CL_DEVICE_ATOMIC_FENCE_CAPABILITIES:
2427 case CL_DEVICE_ATOMIC_MEMORY_CAPABILITIES:
2428 case CL_DEVICE_SVM_CAPABILITIES:
2429 case CL_DEVICE_QUEUE_ON_DEVICE_PROPERTIES:
2430 case CL_DEVICE_PARTITION_AFFINITY_DOMAIN:
2431 case CL_DEVICE_DOUBLE_FP_CONFIG:
2432 case CL_DEVICE_QUEUE_ON_HOST_PROPERTIES:
2433 case CL_DEVICE_EXECUTION_CAPABILITIES:
2434 // cl_ulong and cl_bitfield
2435 json.addScalar(infoString, *static_cast<const cl_ulong *>(data));
2436 break;
2437 case CL_DEVICE_VENDOR_ID:
2438 case CL_DEVICE_MAX_COMPUTE_UNITS:
2439 case CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS:
2440 case CL_DEVICE_PREFERRED_VECTOR_WIDTH_CHAR:
2441 case CL_DEVICE_PREFERRED_VECTOR_WIDTH_SHORT:
2442 case CL_DEVICE_PREFERRED_VECTOR_WIDTH_INT:
2443 case CL_DEVICE_PREFERRED_VECTOR_WIDTH_LONG:
2444 case CL_DEVICE_PREFERRED_VECTOR_WIDTH_FLOAT:
2445 case CL_DEVICE_PREFERRED_VECTOR_WIDTH_DOUBLE:
2446 case CL_DEVICE_PREFERRED_VECTOR_WIDTH_HALF:
2447 case CL_DEVICE_NATIVE_VECTOR_WIDTH_CHAR:
2448 case CL_DEVICE_NATIVE_VECTOR_WIDTH_SHORT:
2449 case CL_DEVICE_NATIVE_VECTOR_WIDTH_INT:
2450 case CL_DEVICE_NATIVE_VECTOR_WIDTH_LONG:
2451 case CL_DEVICE_NATIVE_VECTOR_WIDTH_FLOAT:
2452 case CL_DEVICE_NATIVE_VECTOR_WIDTH_DOUBLE:
2453 case CL_DEVICE_NATIVE_VECTOR_WIDTH_HALF:
2454 case CL_DEVICE_MAX_CLOCK_FREQUENCY:
2455 case CL_DEVICE_ADDRESS_BITS:
2456 case CL_DEVICE_IMAGE_SUPPORT:
2457 case CL_DEVICE_MAX_READ_IMAGE_ARGS:
2458 case CL_DEVICE_MAX_WRITE_IMAGE_ARGS:
2459 case CL_DEVICE_MAX_READ_WRITE_IMAGE_ARGS:
2460 case CL_DEVICE_PIPE_SUPPORT:
2461 case CL_DEVICE_GENERIC_ADDRESS_SPACE_SUPPORT:
2462 case CL_DEVICE_WORK_GROUP_COLLECTIVE_FUNCTIONS_SUPPORT:
2463 case CL_DEVICE_NON_UNIFORM_WORK_GROUP_SUPPORT:
2464 case CL_DEVICE_NUMERIC_VERSION:
2465 case CL_DEVICE_SUB_GROUP_INDEPENDENT_FORWARD_PROGRESS:
2466 case CL_DEVICE_MAX_NUM_SUB_GROUPS:
2467 case CL_DEVICE_PREFERRED_LOCAL_ATOMIC_ALIGNMENT:
2468 case CL_DEVICE_PREFERRED_GLOBAL_ATOMIC_ALIGNMENT:
2469 case CL_DEVICE_PREFERRED_PLATFORM_ATOMIC_ALIGNMENT:
2470 case CL_DEVICE_PIPE_MAX_PACKET_SIZE:
2471 case CL_DEVICE_PIPE_MAX_ACTIVE_RESERVATIONS:
2472 case CL_DEVICE_MAX_PIPE_ARGS:
2473 case CL_DEVICE_MAX_ON_DEVICE_EVENTS:
2474 case CL_DEVICE_MAX_ON_DEVICE_QUEUES:
2475 case CL_DEVICE_QUEUE_ON_DEVICE_MAX_SIZE:
2476 case CL_DEVICE_QUEUE_ON_DEVICE_PREFERRED_SIZE:
2477 case CL_DEVICE_IMAGE_BASE_ADDRESS_ALIGNMENT:
2478 case CL_DEVICE_IMAGE_PITCH_ALIGNMENT:
2479 case CL_DEVICE_PREFERRED_INTEROP_USER_SYNC:
2480 case CL_DEVICE_REFERENCE_COUNT:
2481 case CL_DEVICE_PARTITION_MAX_SUB_DEVICES:
2482 case CL_DEVICE_LINKER_AVAILABLE:
2483 case CL_DEVICE_HOST_UNIFIED_MEMORY:
2484 case CL_DEVICE_COMPILER_AVAILABLE:
2485 case CL_DEVICE_AVAILABLE:
2486 case CL_DEVICE_ENDIAN_LITTLE:
2487 case CL_DEVICE_ERROR_CORRECTION_SUPPORT:
2488 case CL_DEVICE_LOCAL_MEM_TYPE:
2489 case CL_DEVICE_MAX_CONSTANT_ARGS:
2490 case CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE:
2491 case CL_DEVICE_MIN_DATA_TYPE_ALIGN_SIZE:
2492 case CL_DEVICE_MEM_BASE_ADDR_ALIGN:
2493 case CL_DEVICE_MAX_SAMPLERS:
2494 json.addScalar(infoString, *static_cast<const cl_uint *>(data));
2495 break;
2496 case CL_DEVICE_MAX_WORK_GROUP_SIZE:
2497 case CL_DEVICE_IMAGE2D_MAX_WIDTH:
2498 case CL_DEVICE_IMAGE2D_MAX_HEIGHT:
2499 case CL_DEVICE_IMAGE3D_MAX_WIDTH:
2500 case CL_DEVICE_IMAGE3D_MAX_HEIGHT:
2501 case CL_DEVICE_IMAGE3D_MAX_DEPTH:
2502 case CL_DEVICE_IMAGE_MAX_BUFFER_SIZE:
2503 case CL_DEVICE_IMAGE_MAX_ARRAY_SIZE:
2504 case CL_DEVICE_PREFERRED_WORK_GROUP_SIZE_MULTIPLE:
2505 case CL_DEVICE_GLOBAL_VARIABLE_PREFERRED_TOTAL_SIZE:
2506 case CL_DEVICE_MAX_GLOBAL_VARIABLE_SIZE:
2507 case CL_DEVICE_PRINTF_BUFFER_SIZE:
2508 case CL_DEVICE_PROFILING_TIMER_RESOLUTION:
2509 case CL_DEVICE_MAX_PARAMETER_SIZE:
2510 json.addScalar(infoString, *static_cast<const size_t *>(data));
2511 break;
2512 case CL_DEVICE_MAX_WORK_ITEM_SIZES:
2513 json.addVector(infoString,
2514 std::vector<size_t>((size_t *)data,
2515 (size_t *)data + size / sizeof(size_t)));
2516 break;
2517 case CL_DEVICE_PARTITION_TYPE:
2518 case CL_DEVICE_PARTITION_PROPERTIES:
2519 json.addVector(infoString, std::vector<cl_ulong>(
2520 (cl_ulong *)data,
2521 (cl_ulong *)data + size / sizeof(cl_ulong)));
2522 break;
2523 case CL_DEVICE_PARENT_DEVICE:
2524 {
2525 std::ostringstream voidStream;
2526 voidStream << static_cast<void *>(*(cl_device_id *)data);
2527 std::string memLoc = voidStream.str();
2528 json.addCString(infoString, memLoc.c_str());
2529 break;
2530 }
2531 case CL_DEVICE_PLATFORM:
2532 {
2533 std::ostringstream voidStream;
2534 voidStream << static_cast<void *>(*(cl_platform_id *)data);
2535 std::string memLoc = voidStream.str();
2536 json.addCString(infoString, memLoc.c_str());
2537 break;
2538 }
2539 case CL_DEVICE_ILS_WITH_VERSION:
2540 case CL_DEVICE_OPENCL_C_FEATURES:
2541 case CL_DEVICE_OPENCL_C_ALL_VERSIONS:
2542 case CL_DEVICE_BUILT_IN_KERNELS_WITH_VERSION:
2543 case CL_DEVICE_EXTENSIONS_WITH_VERSION:
2544 {
2545 const cl_name_version *nameVersion = static_cast<const cl_name_version *>(data);
2546 json.startGroup(infoString);
2547 for (size_t j = 0; j < size / sizeof(cl_name_version); ++j)
2548 {
2549 json.addScalar(nameVersion[j].name, nameVersion[j].version);
2550 }
2551 json.endGroup();
2552 break;
2553 }
2554
2555 default:
2556 // Not supported or cannot add to JSON file
2557 break;
2558 }
2559
2560 break;
2561 }
2562 case EntryPoint::CLGetContextInfo:
2563 {
2564 cl::ContextInfo info = call.params.getParamCaptures()[1].value.ContextInfoVal;
2565 std::ostringstream oss;
2566 oss << info;
2567 std::string infoString = oss.str();
2568
2569 json.addString("context", clObject);
2570 switch (ToCLenum(info))
2571 {
2572 case CL_PLATFORM_PROFILE:
2573 case CL_PLATFORM_VERSION:
2574 case CL_PLATFORM_NAME:
2575 case CL_PLATFORM_VENDOR:
2576 case CL_PLATFORM_EXTENSIONS:
2577 case CL_PLATFORM_ICD_SUFFIX_KHR:
2578 {
2579 json.addCString(infoString, static_cast<const char *>(data));
2580 break;
2581 }
2582 case CL_CONTEXT_REFERENCE_COUNT:
2583 case CL_CONTEXT_NUM_DEVICES:
2584 json.addScalar(infoString, *static_cast<const cl_uint *>(data));
2585 break;
2586 case CL_CONTEXT_PROPERTIES:
2587 json.addVector(infoString, std::vector<cl_ulong>(
2588 (cl_ulong *)data,
2589 (cl_ulong *)data + size / sizeof(cl_ulong)));
2590 break;
2591 case CL_CONTEXT_DEVICES:
2592 {
2593 const cl_device_id *devices = static_cast<const cl_device_id *>(data);
2594 std::vector<std::string> devicesStrings;
2595 for (size_t j = 0; j < size / sizeof(cl_device_id); ++j)
2596 {
2597 std::ostringstream voidStream;
2598 voidStream << static_cast<void *>(devices[j]);
2599 devicesStrings.push_back(voidStream.str());
2600 }
2601 json.addVectorOfStrings(infoString, devicesStrings);
2602 break;
2603 }
2604 default:
2605 // Not supported or cannot add to JSON file
2606 break;
2607 }
2608
2609 break;
2610 }
2611 case EntryPoint::CLGetCommandQueueInfo:
2612 {
2613 cl::CommandQueueInfo info = call.params.getParamCaptures()[1].value.CommandQueueInfoVal;
2614 std::ostringstream oss;
2615 oss << info;
2616 std::string infoString = oss.str();
2617
2618 json.addString("command_queue", clObject);
2619 switch (ToCLenum(info))
2620 {
2621 case CL_QUEUE_DEVICE_DEFAULT:
2622 {
2623 std::ostringstream voidStream;
2624 voidStream << static_cast<void *>(*(cl_command_queue *)data);
2625 std::string memLoc = voidStream.str();
2626 json.addCString(infoString, memLoc.c_str());
2627 break;
2628 }
2629 case CL_QUEUE_CONTEXT:
2630 {
2631 std::ostringstream voidStream;
2632 voidStream << static_cast<void *>(*(cl_context *)data);
2633 std::string memLoc = voidStream.str();
2634 json.addCString(infoString, memLoc.c_str());
2635 break;
2636 }
2637 case CL_QUEUE_DEVICE:
2638 {
2639 std::ostringstream voidStream;
2640 voidStream << static_cast<void *>(*(cl_device_id *)data);
2641 std::string memLoc = voidStream.str();
2642 json.addCString(infoString, memLoc.c_str());
2643 break;
2644 }
2645 case CL_QUEUE_REFERENCE_COUNT:
2646 case CL_QUEUE_SIZE:
2647 {
2648 json.addScalar(infoString, *static_cast<const cl_uint *>(data));
2649 break;
2650 }
2651 case CL_QUEUE_PROPERTIES:
2652 {
2653 json.addScalar(infoString, *static_cast<const cl_ulong *>(data));
2654 break;
2655 }
2656 case CL_QUEUE_PROPERTIES_ARRAY:
2657 {
2658 json.addVector(infoString, std::vector<cl_ulong>(
2659 (cl_ulong *)data,
2660 (cl_ulong *)data + size / sizeof(cl_ulong)));
2661 break;
2662 }
2663 default:
2664 // Not supported or cannot add to JSON file
2665 break;
2666 }
2667
2668 break;
2669 }
2670 case EntryPoint::CLGetProgramInfo:
2671 {
2672 cl::ProgramInfo info = call.params.getParamCaptures()[1].value.ProgramInfoVal;
2673 std::ostringstream oss;
2674 oss << info;
2675 std::string infoString = oss.str();
2676
2677 json.addString("program", clObject);
2678 switch (ToCLenum(info))
2679 {
2680 case CL_PROGRAM_SOURCE:
2681 case CL_PROGRAM_IL:
2682 case CL_PROGRAM_KERNEL_NAMES:
2683 {
2684 json.addCString(infoString, static_cast<const char *>(data));
2685 break;
2686 }
2687 case CL_PROGRAM_CONTEXT:
2688 {
2689 std::ostringstream voidStream;
2690 voidStream << static_cast<void *>(*(cl_context *)data);
2691 std::string memLoc = voidStream.str();
2692 json.addCString(infoString, memLoc.c_str());
2693 break;
2694 }
2695 case CL_PROGRAM_REFERENCE_COUNT:
2696 case CL_PROGRAM_NUM_DEVICES:
2697 case CL_PROGRAM_SCOPE_GLOBAL_CTORS_PRESENT:
2698 case CL_PROGRAM_SCOPE_GLOBAL_DTORS_PRESENT:
2699 {
2700 json.addScalar(infoString, *static_cast<const cl_uint *>(data));
2701 break;
2702 }
2703 case CL_PROGRAM_DEVICES:
2704 {
2705 const cl_device_id *devices = static_cast<const cl_device_id *>(data);
2706 std::vector<std::string> devicesStrings;
2707 for (size_t j = 0; j < size / sizeof(cl_device_id); ++j)
2708 {
2709 std::ostringstream voidStream;
2710 voidStream << static_cast<void *>(devices[j]);
2711 devicesStrings.push_back(voidStream.str());
2712 }
2713 json.addVectorOfStrings(infoString, devicesStrings);
2714 break;
2715 }
2716 case CL_PROGRAM_NUM_KERNELS:
2717 {
2718 json.addScalar(infoString, *static_cast<const size_t *>(data));
2719 break;
2720 }
2721 case CL_PROGRAM_BINARY_SIZES:
2722 json.addVector(infoString,
2723 std::vector<size_t>((size_t *)data,
2724 (size_t *)data + size / sizeof(size_t)));
2725 break;
2726 case CL_PROGRAM_BINARIES:
2727 json.addVector(infoString,
2728 std::vector<unsigned char>(
2729 (unsigned char *)data,
2730 (unsigned char *)data + size / sizeof(unsigned char)));
2731 break;
2732 default:
2733 // Not supported or cannot add to JSON file
2734 break;
2735 }
2736
2737 break;
2738 }
2739 case EntryPoint::CLGetProgramBuildInfo:
2740 {
2741 cl::ProgramBuildInfo info = call.params.getParamCaptures()[2].value.ProgramBuildInfoVal;
2742 std::ostringstream oss;
2743 oss << info;
2744 std::string infoString = oss.str();
2745
2746 json.addString("program", clObject);
2747 cl_device_id device = call.params.getParamCaptures()[1].value.cl_device_idVal;
2748 std::ostringstream objectStream2;
2749 objectStream2 << static_cast<void *>(device);
2750 clObject = objectStream2.str();
2751 json.addString("device", clObject);
2752 switch (ToCLenum(info))
2753 {
2754 case CL_PROGRAM_BUILD_OPTIONS:
2755 case CL_PROGRAM_BUILD_LOG:
2756 {
2757 json.addCString(infoString, static_cast<const char *>(data));
2758 break;
2759 }
2760 case CL_PROGRAM_BINARY_TYPE:
2761 {
2762 json.addScalar(infoString, *static_cast<const cl_uint *>(data));
2763 break;
2764 }
2765 case CL_PROGRAM_BUILD_STATUS:
2766 {
2767 json.addScalar(infoString, *static_cast<const cl_int *>(data));
2768 break;
2769 }
2770 case CL_PROGRAM_BUILD_GLOBAL_VARIABLE_TOTAL_SIZE:
2771 {
2772 json.addScalar(infoString, *static_cast<const size_t *>(data));
2773 break;
2774 }
2775 default:
2776 // Not supported or cannot add to JSON file
2777 break;
2778 }
2779 break;
2780 }
2781 case EntryPoint::CLGetKernelInfo:
2782 {
2783 cl::KernelInfo info = call.params.getParamCaptures()[1].value.KernelInfoVal;
2784 std::ostringstream oss;
2785 oss << info;
2786 std::string infoString = oss.str();
2787
2788 json.addString("kernel", clObject);
2789 switch (ToCLenum(info))
2790 {
2791 case CL_KERNEL_FUNCTION_NAME:
2792 case CL_KERNEL_ATTRIBUTES:
2793 {
2794 json.addCString(infoString, static_cast<const char *>(data));
2795 break;
2796 }
2797 case CL_KERNEL_NUM_ARGS:
2798 case CL_KERNEL_REFERENCE_COUNT:
2799 {
2800 json.addScalar(infoString, *static_cast<const cl_uint *>(data));
2801 break;
2802 }
2803 case CL_KERNEL_CONTEXT:
2804 {
2805 std::ostringstream voidStream;
2806 voidStream << static_cast<void *>(*(cl_context *)data);
2807 std::string memLoc = voidStream.str();
2808 json.addCString(infoString, memLoc.c_str());
2809 break;
2810 }
2811 case CL_KERNEL_PROGRAM:
2812 {
2813 std::ostringstream voidStream;
2814 voidStream << static_cast<void *>(*(cl_program *)data);
2815 std::string memLoc = voidStream.str();
2816 json.addCString(infoString, memLoc.c_str());
2817 break;
2818 }
2819 default:
2820 // Not supported or cannot add to JSON file
2821 break;
2822 }
2823 break;
2824 }
2825 case EntryPoint::CLGetKernelArgInfo:
2826 {
2827 cl::KernelArgInfo info = call.params.getParamCaptures()[2].value.KernelArgInfoVal;
2828 std::ostringstream oss;
2829 oss << info;
2830 std::string infoString = oss.str();
2831
2832 json.addString("kernel", clObject);
2833 cl_uint index = call.params.getParamCaptures()[1].value.cl_uintVal;
2834 json.addScalar("arg_index", index);
2835 switch (ToCLenum(info))
2836 {
2837 case CL_KERNEL_ARG_TYPE_NAME:
2838 case CL_KERNEL_ARG_NAME:
2839 {
2840 json.addCString(infoString, static_cast<const char *>(data));
2841 break;
2842 }
2843 case CL_KERNEL_ARG_ADDRESS_QUALIFIER:
2844 case CL_KERNEL_ARG_ACCESS_QUALIFIER:
2845 {
2846 json.addScalar(infoString, *static_cast<const cl_uint *>(data));
2847 break;
2848 }
2849 case CL_KERNEL_ARG_TYPE_QUALIFIER:
2850 {
2851 json.addScalar(infoString, *static_cast<const cl_ulong *>(data));
2852 break;
2853 }
2854 default:
2855 // Not supported or cannot add to JSON file
2856 break;
2857 }
2858 break;
2859 }
2860 case EntryPoint::CLGetKernelWorkGroupInfo:
2861 {
2862 cl::KernelWorkGroupInfo info =
2863 call.params.getParamCaptures()[2].value.KernelWorkGroupInfoVal;
2864 std::ostringstream oss;
2865 oss << info;
2866 std::string infoString = oss.str();
2867
2868 json.addString("kernel", clObject);
2869 cl_device_id device = call.params.getParamCaptures()[1].value.cl_device_idVal;
2870 std::ostringstream objectStream2;
2871 objectStream2 << static_cast<void *>(device);
2872 clObject = objectStream2.str();
2873 json.addString("device", clObject);
2874 switch (ToCLenum(info))
2875 {
2876 case CL_KERNEL_LOCAL_MEM_SIZE:
2877 case CL_KERNEL_PRIVATE_MEM_SIZE:
2878 {
2879 json.addScalar(infoString, *static_cast<const cl_ulong *>(data));
2880 break;
2881 }
2882 case CL_KERNEL_WORK_GROUP_SIZE:
2883 case CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE:
2884 {
2885 json.addScalar(infoString, *static_cast<const size_t *>(data));
2886 break;
2887 }
2888 case CL_KERNEL_GLOBAL_WORK_SIZE:
2889 case CL_KERNEL_COMPILE_WORK_GROUP_SIZE:
2890 {
2891 json.addVector(infoString,
2892 std::vector<size_t>((size_t *)data, (size_t *)data + 3));
2893 break;
2894 }
2895 default:
2896 // Not supported or cannot add to JSON file
2897 break;
2898 }
2899 break;
2900 }
2901 case EntryPoint::CLGetEventInfo:
2902 {
2903 cl::EventInfo info = call.params.getParamCaptures()[1].value.EventInfoVal;
2904 std::ostringstream oss;
2905 oss << info;
2906 std::string infoString = oss.str();
2907
2908 json.addString("event", clObject);
2909 switch (ToCLenum(info))
2910 {
2911 case CL_EVENT_REFERENCE_COUNT:
2912 case CL_EVENT_COMMAND_TYPE:
2913 {
2914 json.addScalar(infoString, *static_cast<const cl_uint *>(data));
2915 break;
2916 }
2917 case CL_EVENT_COMMAND_EXECUTION_STATUS:
2918 {
2919 json.addScalar(infoString, *static_cast<const cl_int *>(data));
2920 break;
2921 }
2922 case CL_EVENT_CONTEXT:
2923 {
2924 std::ostringstream voidStream;
2925 voidStream << static_cast<void *>(*(cl_context *)data);
2926 std::string memLoc = voidStream.str();
2927 json.addCString(infoString, memLoc.c_str());
2928 break;
2929 }
2930 case CL_EVENT_COMMAND_QUEUE:
2931 {
2932 std::ostringstream voidStream;
2933 voidStream << static_cast<void *>(*(cl_command_queue *)data);
2934 std::string memLoc = voidStream.str();
2935 json.addCString(infoString, memLoc.c_str());
2936 break;
2937 }
2938 default:
2939 // Not supported or cannot add to JSON file
2940 break;
2941 }
2942 break;
2943 }
2944 case EntryPoint::CLGetEventProfilingInfo:
2945 {
2946 cl::ProfilingInfo info = call.params.getParamCaptures()[1].value.ProfilingInfoVal;
2947 std::ostringstream oss;
2948 oss << info;
2949 std::string infoString = oss.str();
2950
2951 json.addString("event", clObject);
2952 switch (ToCLenum(info))
2953 {
2954 case CL_PROFILING_COMMAND_QUEUED:
2955 case CL_PROFILING_COMMAND_SUBMIT:
2956 case CL_PROFILING_COMMAND_START:
2957 case CL_PROFILING_COMMAND_END:
2958 case CL_PROFILING_COMMAND_COMPLETE:
2959 {
2960 json.addScalar(infoString, *static_cast<const cl_ulong *>(data));
2961 break;
2962 }
2963 default:
2964 // Not supported or cannot add to JSON file
2965 break;
2966 }
2967 break;
2968 }
2969 case EntryPoint::CLGetMemObjectInfo:
2970 {
2971 cl::MemInfo info = call.params.getParamCaptures()[1].value.MemInfoVal;
2972 std::ostringstream oss;
2973 oss << info;
2974 std::string infoString = oss.str();
2975
2976 json.addString("memObj", clObject);
2977 switch (ToCLenum(info))
2978 {
2979 case CL_MEM_TYPE:
2980 case CL_MEM_MAP_COUNT:
2981 case CL_MEM_REFERENCE_COUNT:
2982 case CL_MEM_USES_SVM_POINTER:
2983 {
2984 json.addScalar(infoString, *static_cast<const cl_uint *>(data));
2985 break;
2986 }
2987 case CL_MEM_FLAGS:
2988 {
2989 json.addScalar(infoString, *static_cast<const cl_ulong *>(data));
2990 break;
2991 }
2992 case CL_MEM_SIZE:
2993 case CL_MEM_OFFSET:
2994 {
2995 json.addScalar(infoString, *static_cast<const size_t *>(data));
2996 break;
2997 }
2998 case CL_MEM_HOST_PTR:
2999 {
3000 std::ostringstream voidStream;
3001 voidStream << data;
3002 std::string memLoc = voidStream.str();
3003 json.addCString(infoString, memLoc.c_str());
3004 break;
3005 }
3006 case CL_MEM_CONTEXT:
3007 {
3008 std::ostringstream voidStream;
3009 voidStream << static_cast<void *>(*(cl_context *)data);
3010 std::string memLoc = voidStream.str();
3011 json.addCString(infoString, memLoc.c_str());
3012 break;
3013 }
3014 case CL_MEM_ASSOCIATED_MEMOBJECT:
3015 {
3016 std::ostringstream voidStream;
3017 voidStream << static_cast<void *>(*(cl_mem *)data);
3018 std::string memLoc = voidStream.str();
3019 json.addCString(infoString, memLoc.c_str());
3020 break;
3021 }
3022 case CL_MEM_PROPERTIES:
3023 {
3024 json.addVector(infoString, std::vector<cl_mem_properties>(
3025 (cl_mem_properties *)data,
3026 (cl_mem_properties *)data +
3027 size / sizeof(cl_mem_properties)));
3028 break;
3029 }
3030 default:
3031 // Not supported or cannot add to JSON file
3032 break;
3033 }
3034 break;
3035 }
3036 case EntryPoint::CLGetImageInfo:
3037 {
3038 cl::ImageInfo info = call.params.getParamCaptures()[1].value.ImageInfoVal;
3039 std::ostringstream oss;
3040 oss << info;
3041 std::string infoString = oss.str();
3042
3043 json.addString("image", clObject);
3044 switch (ToCLenum(info))
3045 {
3046 case CL_IMAGE_FORMAT:
3047 {
3048 json.startGroup(infoString);
3049 json.addScalar("image_channel_order",
3050 static_cast<const cl_image_format *>(data)->image_channel_order);
3051 json.addScalar(
3052 "image_channel_data_type",
3053 static_cast<const cl_image_format *>(data)->image_channel_data_type);
3054 json.endGroup();
3055 break;
3056 }
3057 case CL_IMAGE_NUM_MIP_LEVELS:
3058 case CL_IMAGE_NUM_SAMPLES:
3059 {
3060 json.addScalar(infoString, *static_cast<const cl_uint *>(data));
3061 break;
3062 }
3063 case CL_IMAGE_ELEMENT_SIZE:
3064 case CL_IMAGE_ROW_PITCH:
3065 case CL_IMAGE_SLICE_PITCH:
3066 case CL_IMAGE_WIDTH:
3067 case CL_IMAGE_HEIGHT:
3068 case CL_IMAGE_DEPTH:
3069 case CL_IMAGE_ARRAY_SIZE:
3070 {
3071 json.addScalar(infoString, *static_cast<const size_t *>(data));
3072 break;
3073 }
3074 case CL_IMAGE_BUFFER:
3075 {
3076 std::ostringstream voidStream;
3077 voidStream << static_cast<void *>(*(cl_mem *)data);
3078 std::string memLoc = voidStream.str();
3079 json.addCString(infoString, memLoc.c_str());
3080 break;
3081 }
3082 default:
3083 // Not supported or cannot add to JSON file
3084 break;
3085 }
3086 break;
3087 }
3088 case EntryPoint::CLGetSamplerInfo:
3089 {
3090 cl::SamplerInfo info = call.params.getParamCaptures()[1].value.SamplerInfoVal;
3091 std::ostringstream oss;
3092 oss << info;
3093 std::string infoString = oss.str();
3094
3095 json.addString("image", clObject);
3096 switch (ToCLenum(info))
3097 {
3098 case CL_SAMPLER_REFERENCE_COUNT:
3099 case CL_SAMPLER_NORMALIZED_COORDS:
3100 case CL_SAMPLER_ADDRESSING_MODE:
3101 case CL_SAMPLER_FILTER_MODE:
3102 {
3103 json.addScalar(infoString, *static_cast<const cl_uint *>(data));
3104 break;
3105 }
3106 case CL_SAMPLER_PROPERTIES:
3107 {
3108 json.addVector(infoString, std::vector<cl_sampler_properties>(
3109 (cl_sampler_properties *)data,
3110 (cl_sampler_properties *)data +
3111 size / sizeof(cl_sampler_properties)));
3112 break;
3113 }
3114 case CL_SAMPLER_CONTEXT:
3115 {
3116 std::ostringstream voidStream;
3117 voidStream << static_cast<void *>(*(cl_context *)data);
3118 std::string memLoc = voidStream.str();
3119 json.addCString(infoString, memLoc.c_str());
3120 break;
3121 }
3122 default:
3123 // Not supported or cannot add to JSON file
3124 break;
3125 }
3126 break;
3127 }
3128 default:
3129 break;
3130 }
3131
3132 json.endGroup();
3133
3134 mCLInfoJson += std::string(json.data()) + ",\n";
3135 }
3136
writeJSONCLGetInfo()3137 void FrameCaptureShared::writeJSONCLGetInfo()
3138 {
3139 std::stringstream jsonFileNameStream;
3140 jsonFileNameStream << mOutDirectory << FmtCapturePrefix(kNoContextId, mCaptureLabel)
3141 << "_OpenCL_info.json";
3142 std::string jsonFileName = jsonFileNameStream.str();
3143
3144 SaveFileHelper saveData(jsonFileName);
3145
3146 saveData.write(reinterpret_cast<const uint8_t *>(mCLInfoJson.c_str()), mCLInfoJson.length());
3147 }
3148
writeCppReplayIndexFilesCL()3149 void FrameCaptureShared::writeCppReplayIndexFilesCL()
3150 {
3151 // Ensure the last frame is written. This will no-op if the frame is already written.
3152 mReplayWriter.saveFrame();
3153
3154 {
3155 std::stringstream header;
3156
3157 header << "#pragma once\n";
3158 header << "\n";
3159 header << "#define CL_NO_EXTENSION_PROTOTYPES\n";
3160 header << "#include <angle_cl.h>\n";
3161 header << "#include <stdint.h>\n";
3162 header << "#include \"trace_fixture_cl.h\"\n";
3163
3164 std::string includes = header.str();
3165 mReplayWriter.setHeaderPrologue(includes);
3166 }
3167
3168 {
3169 std::stringstream source;
3170
3171 source << "#include \"" << FmtCapturePrefix(kNoContextId, mCaptureLabel) << ".h\"\n";
3172 source << "#include \"trace_fixture_cl.h\"\n";
3173
3174 std::string sourcePrologue = source.str();
3175 mReplayWriter.setSourcePrologue(sourcePrologue);
3176 }
3177
3178 {
3179 std::string proto = "void InitReplay(void)";
3180
3181 std::stringstream source;
3182 source << proto << "\n";
3183 source << "{\n";
3184 WriteInitReplayCallCL(mCompression, source, mCaptureLabel, 0, mReadBufferSize,
3185 mMaxCLParamsSize);
3186 source << "}\n";
3187
3188 mReplayWriter.addPrivateFunction(proto, std::stringstream(), source);
3189 }
3190
3191 {
3192 std::string proto = "void ReplayFrame(uint32_t frameIndex)";
3193
3194 std::stringstream source;
3195
3196 source << proto << "\n";
3197 source << "{\n";
3198 source << " switch (frameIndex)\n";
3199 source << " {\n";
3200 for (uint32_t frameIndex : mActiveFrameIndices)
3201 {
3202 source << " case " << frameIndex << ":\n";
3203 source << " " << FmtReplayFunction(kNoContextId, FuncUsage::Call, frameIndex)
3204 << ";\n";
3205 source << " break;\n";
3206 }
3207 source << " default:\n";
3208 source << " break;\n";
3209 source << " }\n";
3210 source << "}\n";
3211
3212 mReplayWriter.addPublicFunction(proto, std::stringstream(), source);
3213 }
3214
3215 for (auto extFuncName : mExtFuncsAdded)
3216 {
3217 mReplayWriter.addStaticVariable(extFuncName + "_fn", extFuncName);
3218 }
3219
3220 std::stringstream protoSetupStream;
3221 protoSetupStream << "void SetupFirstFrame(void)";
3222 std::string protoSetup = protoSetupStream.str();
3223 std::stringstream headerStreamSetup;
3224 std::stringstream bodyStreamSetup;
3225 WriteCppReplayFunctionWithPartsCL(ReplayFunc::SetupFirstFrame, mReplayWriter,
3226 mCaptureStartFrame, &mBinaryData, mCLSetupCalls,
3227 headerStreamSetup, bodyStreamSetup);
3228 mReplayWriter.addPublicFunction(protoSetup, headerStreamSetup, bodyStreamSetup);
3229
3230 {
3231 std::string proto = "void ResetReplay(void)";
3232 std::stringstream source;
3233 source << proto << "\n" << "{\n";
3234 printCLResetObjs(source);
3235 source << "}\n";
3236 mReplayWriter.addPublicFunction(proto, std::stringstream(), source);
3237 }
3238
3239 {
3240 std::stringstream fnameStream;
3241 fnameStream << mOutDirectory << FmtCapturePrefix(kNoContextId, mCaptureLabel);
3242 std::string fnamePattern = fnameStream.str();
3243
3244 mReplayWriter.setFilenamePattern(fnamePattern);
3245 }
3246
3247 mReplayWriter.saveIndexFilesAndHeader();
3248
3249 writeJSONCL();
3250 writeJSONCLGetInfo();
3251 }
3252
writeMainContextCppReplayCL()3253 void FrameCaptureShared::writeMainContextCppReplayCL()
3254 {
3255 {
3256 std::stringstream header;
3257
3258 header << "#include \"" << FmtCapturePrefix(kNoContextId, mCaptureLabel) << ".h\"\n";
3259 header << "#include \"trace_fixture_cl.h\"\n";
3260
3261 std::string headerString = header.str();
3262 mReplayWriter.setSourcePrologue(headerString);
3263 }
3264
3265 uint32_t frameIndex = getReplayFrameIndex();
3266
3267 if (frameIndex == 1)
3268 {
3269 {
3270 std::string proto = "void SetupReplay(void)";
3271
3272 std::stringstream out;
3273
3274 out << proto << "\n";
3275 out << "{\n";
3276
3277 // Setup all of the shared objects.
3278 out << " InitReplay();\n";
3279
3280 out << "}\n";
3281
3282 mReplayWriter.addPublicFunction(proto, std::stringstream(), out);
3283 }
3284 }
3285
3286 if (!mFrameCalls.empty())
3287 {
3288 std::stringstream protoStream;
3289 protoStream << "void "
3290 << FmtReplayFunction(kNoContextId, FuncUsage::Prototype, mFrameIndex);
3291 std::string proto = protoStream.str();
3292 std::stringstream headerStream;
3293 std::stringstream bodyStream;
3294
3295 WriteCppReplayFunctionWithPartsCL(ReplayFunc::Replay, mReplayWriter, mFrameIndex,
3296 &mBinaryData, mFrameCalls, headerStream, bodyStream);
3297
3298 mReplayWriter.addPrivateFunction(proto, headerStream, bodyStream);
3299 }
3300
3301 {
3302 std::stringstream fnamePatternStream;
3303 fnamePatternStream << mOutDirectory << FmtCapturePrefix(kNoContextId, mCaptureLabel);
3304 std::string fnamePattern = fnamePatternStream.str();
3305
3306 mReplayWriter.setFilenamePattern(fnamePattern);
3307 }
3308
3309 if (mFrameIndex == mCaptureEndFrame)
3310 {
3311 mReplayWriter.saveFrame();
3312 }
3313 }
3314
3315 } // namespace angle
3316