Google Git
Sign in
android / platform / external / angle / HEAD / . / src / libANGLE / capture / FrameCaptureCL.cpp
blob: c56dfd04ba753af45d8548de9a26e9086a9260b3 [file] [log] [blame]
//
// Copyright 2025 The ANGLE Project Authors. All rights reserved.
// Use of this source code is governed by a BSD-style license that can be
// found in the LICENSE file.
//
// FrameCaptureCL.cpp:
// ANGLE CL Frame capture implementation.
//
#include "libANGLE/capture/FrameCapture.h"
#include "common/angle_version_info.h"
#include "common/frame_capture_utils.h"
#include "common/serializer/JsonSerializer.h"
#include "common/string_utils.h"
#include "common/system_utils.h"
#include "libANGLE/CLBuffer.h"
#include "libANGLE/CLCommandQueue.h"
#include "libANGLE/CLContext.h"
#include "libANGLE/CLImage.h"
#include "libANGLE/CLProgram.h"
#include "libANGLE/capture/capture_cl_autogen.h"
#include "libANGLE/capture/serialize.h"
#include "libANGLE/cl_utils.h"
#include "libGLESv2/cl_stubs_autogen.h"
#if !ANGLE_CAPTURE_ENABLED
# error Frame capture must be enabled to include this file.
#endif // !ANGLE_CAPTURE_ENABLED
#ifndef ANGLE_ENABLE_CL
# error OpenCL must be enabled to include this file.
#endif // !ANGLE_ENABLE_CL
namespace angle
{
// Some replay functions can get quite large. If over a certain size, this method breaks up the
// function into parts to avoid overflowing the stack and causing slow compilation.
void WriteCppReplayFunctionWithPartsCL(ReplayFunc replayFunc,
ReplayWriter &replayWriter,
uint32_t frameIndex,
std::vector<uint8_t> *binaryData,
const std::vector<CallCapture> &calls,
std::stringstream &header,
std::stringstream &out)
{
out << "void "
<< FmtFunction(replayFunc, kNoContextId, FuncUsage::Definition, frameIndex, kNoPartId)
<< "\n"
<< "{\n";
for (const CallCapture &call : calls)
{
// Process active calls for Setup and inactive calls for SetupInactive
if ((call.isActive && replayFunc != ReplayFunc::SetupInactive) ||
(!call.isActive && replayFunc == ReplayFunc::SetupInactive))
{
out << " ";
WriteCppReplayForCallCL(call, replayWriter, out, header, binaryData);
out << ";\n";
}
}
out << "}\n";
}
void WriteCppReplayForCallCL(const CallCapture &call,
ReplayWriter &replayWriter,
std::ostream &out,
std::ostream &header,
std::vector<uint8_t> *binaryData)
{
if (call.customFunctionName == "Comment")
{
// Just write it directly to the file and move on
WriteComment(out, call);
return;
}
std::ostringstream callOut;
std::ostringstream postCallAdditions;
const ParamCapture &returnValue = call.params.getReturnValue();
switch (returnValue.type)
{
case ParamType::Tcl_context:
callOut << "clContextsMap["
<< cl::Platform::GetDefault()->getFrameCaptureShared()->getIndex(
&returnValue.value.cl_contextVal)
<< "] = ";
break;
case ParamType::Tcl_command_queue:
callOut << "clCommandQueuesMap["
<< cl::Platform::GetDefault()->getFrameCaptureShared()->getIndex(
&returnValue.value.cl_command_queueVal)
<< "] = ";
break;
case ParamType::Tcl_mem:
callOut << "clMemMap["
<< cl::Platform::GetDefault()->getFrameCaptureShared()->getIndex(
&returnValue.value.cl_memVal)
<< "] = ";
break;
case ParamType::Tcl_sampler:
callOut << "clSamplerMap["
<< cl::Platform::GetDefault()->getFrameCaptureShared()->getIndex(
&returnValue.value.cl_samplerVal)
<< "] = ";
break;
case ParamType::Tcl_program:
callOut << "clProgramsMap["
<< cl::Platform::GetDefault()->getFrameCaptureShared()->getIndex(
&returnValue.value.cl_programVal)
<< "] = ";
break;
case ParamType::Tcl_kernel:
callOut << "clKernelsMap["
<< cl::Platform::GetDefault()->getFrameCaptureShared()->getIndex(
&returnValue.value.cl_kernelVal)
<< "] = ";
break;
case ParamType::Tcl_event:
callOut << "clEventsMap["
<< cl::Platform::GetDefault()->getFrameCaptureShared()->getIndex(
&returnValue.value.cl_eventVal)
<< "] = ";
break;
case ParamType::TvoidPointer:
if (cl::Platform::GetDefault()->getFrameCaptureShared()->getCLVoidIndex(
returnValue.value.voidPointerVal) != SIZE_MAX)
{
callOut << "clVoidMap["
<< cl::Platform::GetDefault()->getFrameCaptureShared()->getCLVoidIndex(
returnValue.value.voidPointerVal)
<< "] = ";
}
break;
default:
break;
}
callOut << call.name() << "(";
bool first = true;
for (const ParamCapture &param : call.params.getParamCaptures())
{
if (!first)
{
callOut << ", ";
}
if (param.arrayClientPointerIndex != -1 && param.value.voidConstPointerVal != nullptr)
{
callOut << "gClientArrays[" << param.arrayClientPointerIndex << "]";
}
else if (param.readBufferSizeBytes > 0)
{
callOut << "(" << ParamTypeToString(param.type) << ")gReadBuffer";
}
else if (param.data.empty())
{
if (param.type == ParamType::Tcl_platform_idPointer &&
param.value.cl_platform_idPointerVal)
{
callOut << "clPlatformsMap";
}
else if (param.type == ParamType::Tcl_platform_id)
{
callOut << "clPlatformsMap["
<< cl::Platform::GetDefault()->getFrameCaptureShared()->getIndex(
&param.value.cl_platform_idVal)
<< "]";
}
else if (param.type == ParamType::Tcl_device_idPointer &&
param.value.cl_device_idPointerVal)
{
std::vector<size_t> tempDeviceIndices =
cl::Platform::GetDefault()->getFrameCaptureShared()->getCLObjVector(&param);
cl_uint numDevices = call.params.getParamCaptures()[2].value.cl_uintVal;
out << "temporaryDevicesList.clear();\n temporaryDevicesList.resize("
<< numDevices << ");\n ";
callOut << "temporaryDevicesList.data()";
for (cl_uint i = 0; i < numDevices; ++i)
{
postCallAdditions << ";\n clDevicesMap[" << tempDeviceIndices[i]
<< "] = temporaryDevicesList[" << i << "]";
}
}
else if (param.type == ParamType::Tcl_device_id)
{
callOut << "clDevicesMap["
<< cl::Platform::GetDefault()->getFrameCaptureShared()->getIndex(
&param.value.cl_device_idVal)
<< "]";
}
else if (param.type == ParamType::Tcl_context)
{
callOut << "clContextsMap["
<< cl::Platform::GetDefault()->getFrameCaptureShared()->getIndex(
&param.value.cl_contextVal)
<< "]";
}
else if (param.type == ParamType::Tcl_command_queue)
{
callOut << "clCommandQueuesMap["
<< cl::Platform::GetDefault()->getFrameCaptureShared()->getIndex(
&param.value.cl_command_queueVal)
<< "]";
}
else if (param.type == ParamType::Tcl_mem)
{
callOut << "clMemMap["
<< cl::Platform::GetDefault()->getFrameCaptureShared()->getIndex(
&param.value.cl_memVal)
<< "]";
}
else if (param.type == ParamType::Tcl_sampler)
{
callOut << "clSamplerMap["
<< cl::Platform::GetDefault()->getFrameCaptureShared()->getIndex(
&param.value.cl_samplerVal)
<< "]";
}
else if (param.type == ParamType::Tcl_program)
{
callOut << "clProgramsMap["
<< cl::Platform::GetDefault()->getFrameCaptureShared()->getIndex(
&param.value.cl_programVal)
<< "]";
}
else if (param.type == ParamType::Tcl_kernel)
{
callOut << "clKernelsMap["
<< cl::Platform::GetDefault()->getFrameCaptureShared()->getIndex(
&param.value.cl_kernelVal)
<< "]";
}
else if (param.type == ParamType::Tcl_event)
{
callOut << "clEventsMap["
<< cl::Platform::GetDefault()->getFrameCaptureShared()->getIndex(
&param.value.cl_eventVal)
<< "]";
}
else if (param.type == ParamType::Tcl_eventPointer)
{
if (param.value.cl_eventPointerVal)
{
callOut << "&clEventsMap["
<< cl::Platform::GetDefault()->getFrameCaptureShared()->getIndex(
&param.value.cl_eventVal)
<< "]";
}
else
{
callOut << "NULL";
}
}
else if (param.type == ParamType::TvoidConstPointer)
{
if (cl::Platform::GetDefault()->getFrameCaptureShared()->getIndex(
&param.value.cl_memVal) != SIZE_MAX)
{
callOut << "(const void *)" << "&clMemMap["
<< cl::Platform::GetDefault()->getFrameCaptureShared()->getIndex(
&param.value.cl_memVal)
<< "]";
}
else if (cl::Platform::GetDefault()->getFrameCaptureShared()->getIndex(
&param.value.cl_samplerVal) != SIZE_MAX)
{
callOut << "(const void *)" << "&clSamplerMap["
<< cl::Platform::GetDefault()->getFrameCaptureShared()->getIndex(
&param.value.cl_samplerVal)
<< "]";
}
else if (cl::Platform::GetDefault()->getFrameCaptureShared()->getIndex(
&param.value.cl_command_queueVal) != SIZE_MAX)
{
callOut << "(const void *)" << "&clCommandQueuesMap["
<< cl::Platform::GetDefault()->getFrameCaptureShared()->getIndex(
&param.value.cl_command_queueVal)
<< "]";
}
else if (cl::Platform::GetDefault()->getFrameCaptureShared()->getCLVoidIndex(
param.value.voidConstPointerVal) != SIZE_MAX)
{
callOut << "clVoidMap["
<< cl::Platform::GetDefault()->getFrameCaptureShared()->getCLVoidIndex(
param.value.voidConstPointerVal)
<< "]";
}
else
{
WriteParamCaptureReplay(callOut, call, param);
}
}
else if (param.type == ParamType::TvoidPointer)
{
if (cl::Platform::GetDefault()->getFrameCaptureShared()->getCLVoidIndex(
param.value.voidPointerVal) != SIZE_MAX)
{
callOut << "clVoidMap["
<< cl::Platform::GetDefault()->getFrameCaptureShared()->getCLVoidIndex(
param.value.voidPointerVal)
<< "]";
}
else
{
WriteParamCaptureReplay(callOut, call, param);
}
}
else if (param.type == ParamType::Tcl_mem_destructor_func_type ||
param.type == ParamType::Tcl_callback_func_type ||
param.type == ParamType::Tcl_svm_free_callback_func_type ||
param.type == ParamType::Tcl_program_func_type ||
param.type == ParamType::Tcl_context_destructor_func_type ||
param.type == ParamType::Tcl_context_func_type ||
param.type == ParamType::Tcl_void_func_type)
{
callOut << "NULL";
}
else if (param.type == ParamType::Tcl_memConstPointer && cl::Platform::GetDefault()
->getFrameCaptureShared()
->getCLObjVector(&param)
.size())
{
std::vector<size_t> tempBufferIndices =
cl::Platform::GetDefault()->getFrameCaptureShared()->getCLObjVector(&param);
out << "temporaryBuffersList = {";
for (uint32_t i = 0; i < tempBufferIndices.size(); ++i)
{
out << (i != 0 ? ", " : "") << "clMemMap[" << tempBufferIndices.at(i) << "]";
}
out << "};\n ";
callOut << "temporaryBuffersList.data()";
}
else if (param.type == ParamType::Tcl_eventConstPointer)
{
std::vector<size_t> tempEventIndices =
cl::Platform::GetDefault()->getFrameCaptureShared()->getCLObjVector(&param);
if (tempEventIndices.empty())
{
callOut << "NULL";
}
else
{
out << "temporaryEventsList = {";
for (uint32_t i = 0; i < tempEventIndices.size(); ++i)
{
out << (i != 0 ? ", " : "") << "clEventsMap[" << tempEventIndices.at(i)
<< "]";
}
out << "};\n ";
callOut << "temporaryEventsList.data()";
}
}
else if (param.type == ParamType::Tcl_device_idConstPointer)
{
std::vector<size_t> tempDeviceIndices =
cl::Platform::GetDefault()->getFrameCaptureShared()->getCLObjVector(&param);
if (tempDeviceIndices.empty())
{
callOut << "NULL";
}
else
{
out << "temporaryDevicesList = {";
for (uint32_t i = 0; i < tempDeviceIndices.size(); ++i)
{
if (i != 0)
{
out << ", ";
}
out << "clDevicesMap[" << tempDeviceIndices.at(i) << "]";
}
out << "};\n ";
callOut << "temporaryDevicesList.data()";
}
}
else if (param.type == ParamType::Tcl_kernelPointer)
{
std::vector<size_t> tempKernelIndices =
cl::Platform::GetDefault()->getFrameCaptureShared()->getCLObjVector(&param);
cl_uint numKernels = call.params.getParamCaptures()[1].value.cl_uintVal;
out << "temporaryKernelsList.clear();\ntemporaryKernelsList.resize(" << numKernels
<< ");\n ";
callOut << "temporaryKernelsList.data()";
for (cl_uint i = 0; i < numKernels; ++i)
{
postCallAdditions << ";\n clKernelsMap[" << tempKernelIndices[i]
<< "] = temporaryKernelsList[" << i << "]";
}
}
else if (param.type == ParamType::TvoidConstPointerPointer &&
cl::Platform::GetDefault()
->getFrameCaptureShared()
->getCLObjVector(&param)
.size())
{
std::vector<size_t> offsets =
cl::Platform::GetDefault()->getFrameCaptureShared()->getCLObjVector(&param);
out << "temporaryVoidPtrList = {";
for (size_t i = 0; i < offsets.size(); ++i)
{
out << (i != 0 ? ", " : "") << "&((char*)temporaryVoidPtr)[" << offsets.at(i)
<< "]";
}
out << "};\n ";
callOut << "temporaryVoidPtrList.data()";
}
else if (param.type == ParamType::TvoidPointerPointer ||
param.type == ParamType::TvoidConstPointerPointer)
{
std::vector<size_t> tempVoidIndices =
cl::Platform::GetDefault()->getFrameCaptureShared()->getCLObjVector(&param);
out << "temporaryVoidPtrList = {";
for (uint32_t i = 0; i < tempVoidIndices.size(); ++i)
{
out << (i != 0 ? ", " : "") << "clVoidMap[" << tempVoidIndices.at(i) << "]";
}
out << "};\n ";
callOut << "temporaryVoidPtrList.data()";
}
else if (param.type == ParamType::Tcl_programConstPointer && param.value.size_tVal)
{
std::vector<size_t> tempProgramIndices =
cl::Platform::GetDefault()->getFrameCaptureShared()->getCLObjVector(&param);
out << "temporaryProgramsList = {";
for (uint32_t i = 0; i < tempProgramIndices.size(); ++i)
{
out << (i != 0 ? ", " : "") << "clProgramsMap[" << tempProgramIndices.at(i)
<< "]";
}
out << "};\n ";
callOut << "temporaryProgramsList.data()";
}
else if (param.type == ParamType::Tcl_context_propertiesConstPointer)
{
if (param.value.cl_context_propertiesConstPointerVal)
{
callOut << "temporaryContextProps.data()";
}
else
{
WriteParamCaptureReplay(callOut, call, param);
}
}
else
{
WriteParamCaptureReplay(callOut, call, param);
}
}
else
{
switch (param.type)
{
case ParamType::TcharConstPointerPointer:
WriteStringPointerParamReplay(replayWriter, callOut, header, call, param);
break;
case ParamType::Tcl_device_idPointer:
callOut << "clDevicesMap";
break;
case ParamType::TcharUnsignedConstPointerPointer:
{
std::string tempStructureName = "temporaryCharPointerList";
std::string tempStructureType = "(const char *)";
if (param.type == ParamType::TcharUnsignedConstPointerPointer)
{
tempStructureName = "temporaryUnsignedCharPointerList";
tempStructureType = "(const unsigned char *)";
}
const std::vector<uint8_t> *data;
out << tempStructureName << " = {";
for (size_t i = 0; i < param.data.size(); ++i)
{
if (i != 0)
{
out << ", ";
}
data = &param.data[i];
size_t offset = rx::roundUpPow2(binaryData->size(), kBinaryAlignment);
binaryData->resize(offset + data->size());
memcpy(binaryData->data() + offset, data->data(), data->size());
out << tempStructureType << "&gBinaryData[" << offset << "]";
}
out << "};\n ";
callOut << tempStructureName << ".data()";
break;
}
case ParamType::Tcl_image_descConstPointer:
cl_image_desc tempImageDesc;
std::memcpy(&tempImageDesc, param.data[0].data(), sizeof(cl_image_desc));
if (tempImageDesc.mem_object)
{
out << " std::memcpy(&temporaryImageDesc, ";
WriteBinaryParamReplay(replayWriter, out, header, call, param, binaryData);
out << ", sizeof(cl_image_desc));\ntemporaryImageDesc.mem_object = "
"clMemMap["
<< cl::Platform::GetDefault()->getFrameCaptureShared()->getIndex(
&tempImageDesc.mem_object)
<< "];\n ";
callOut << "&temporaryImageDesc";
}
else
{
WriteBinaryParamReplay(replayWriter, callOut, header, call, param,
binaryData);
}
break;
case ParamType::TvoidPointer:
{
// For clEnqueueNativeKernel
if (call.entryPoint == EntryPoint::CLEnqueueNativeKernel)
{
std::vector<size_t> bufferIndices =
cl::Platform::GetDefault()->getFrameCaptureShared()->getCLObjVector(
&param);
size_t totalSize = call.params.getParamCaptures()[3].value.size_tVal;
out << "temporaryVoidPtr = (void *)std::malloc(" << totalSize
<< ");\nstd::memcpy(&temporaryVoidPtr, ";
WriteBinaryParamReplay(replayWriter, out, header, call, param, binaryData);
out << ", " << totalSize << ");\n ";
callOut << "temporaryVoidPtr";
}
else
{
WriteBinaryParamReplay(replayWriter, callOut, header, call, param,
binaryData);
}
break;
}
default:
WriteBinaryParamReplay(replayWriter, callOut, header, call, param, binaryData);
break;
}
}
first = false;
}
callOut << ")";
out << callOut.str() << postCallAdditions.str();
}
void WriteInitReplayCallCL(bool compression,
std::ostream &out,
const std::string &captureLabel,
size_t maxClientArraySize,
size_t readBufferSize,
const std::map<ParamType, uint32_t> &maxCLParamsSize)
{
std::string binaryDataFileName = GetBinaryDataFilePath(compression, captureLabel);
out << " // binaryDataFileName = " << binaryDataFileName << "\n";
out << " // maxClientArraySize = " << maxClientArraySize << "\n";
out << " // readBufferSize = " << readBufferSize << "\n";
out << " // clPlatformsMapSize = " << maxCLParamsSize.at(ParamType::Tcl_platform_idPointer)
<< "\n";
out << " // clDevicesMapSize = " << maxCLParamsSize.at(ParamType::Tcl_device_idPointer)
<< "\n";
out << " // clContextsMapSize = " << maxCLParamsSize.at(ParamType::Tcl_context) << "\n";
out << " // clCommandQueuesMapSize = " << maxCLParamsSize.at(ParamType::Tcl_command_queue)
<< "\n";
out << " // clMemMapSize = " << maxCLParamsSize.at(ParamType::Tcl_mem) << "\n";
out << " // clEventsMapSize = " << maxCLParamsSize.at(ParamType::Tcl_eventPointer) << "\n";
out << " // clProgramsMapSize = " << maxCLParamsSize.at(ParamType::Tcl_program) << "\n";
out << " // clKernelsMapSize = " << maxCLParamsSize.at(ParamType::Tcl_kernel) << "\n";
out << " // clSamplerMapSize = " << maxCLParamsSize.at(ParamType::Tcl_sampler) << "\n";
out << " // clVoidMapSize = " << maxCLParamsSize.at(ParamType::TvoidPointer) << "\n";
out << " InitializeReplayCL(\"" << binaryDataFileName << "\", " << maxClientArraySize << ", "
<< readBufferSize << ", " << maxCLParamsSize.at(ParamType::Tcl_platform_idPointer) << ", "
<< maxCLParamsSize.at(ParamType::Tcl_device_idPointer) << ", "
<< maxCLParamsSize.at(ParamType::Tcl_context) << ", "
<< maxCLParamsSize.at(ParamType::Tcl_command_queue) << ", "
<< maxCLParamsSize.at(ParamType::Tcl_mem) << ", "
<< maxCLParamsSize.at(ParamType::Tcl_eventPointer) << ", "
<< maxCLParamsSize.at(ParamType::Tcl_program) << ", "
<< maxCLParamsSize.at(ParamType::Tcl_kernel) << ", "
<< maxCLParamsSize.at(ParamType::Tcl_sampler) << ", "
<< maxCLParamsSize.at(ParamType::TvoidPointer) << ");\n";
}
void FrameCaptureShared::trackCLMemUpdate(const cl_mem *mem, bool referenced)
{
std::vector<cl_mem> &mCLDirtyMem = mResourceTrackerCL.mCLDirtyMem;
// retained or created cl mem object
if (referenced)
{
// Potentially mark as dirty
auto it = std::find(mCLDirtyMem.begin(), mCLDirtyMem.end(), *mem);
if (it == mCLDirtyMem.end())
{
mCLDirtyMem.push_back(*mem);
}
}
else
{
std::unordered_map<cl_mem, cl_mem> &mCLSubBufferToParent =
mResourceTrackerCL.mCLSubBufferToParent;
if ((*mem)->cast<cl::Memory>().getRefCount() == 1)
{
auto it = std::find(mCLDirtyMem.begin(), mCLDirtyMem.end(), *mem);
if (it != mCLDirtyMem.end())
{
mCLDirtyMem.erase(it);
}
if (removeUnneededOpenCLCalls)
{
removeCLMemOccurrences(mem, &mFrameCalls);
}
mCLSubBufferToParent.erase(*mem);
(*mem)->cast<cl::Memory>().release();
}
if (mCLSubBufferToParent.find(*mem) != mCLSubBufferToParent.end())
{
trackCLMemUpdate(&mCLSubBufferToParent[*mem], false);
}
}
}
void FrameCaptureShared::trackCLProgramUpdate(const cl_program *program,
bool referenced,
cl_uint numLinkedPrograms,
const cl_program *linkedPrograms)
{
std::unordered_map<cl_program, cl_uint> &mCLProgramLinkCounter =
mResourceTrackerCL.mCLProgramLinkCounter;
std::unordered_map<cl_program, std::vector<cl_program>> &mCLLinkedPrograms =
mResourceTrackerCL.mCLLinkedPrograms;
// retained or created cl program object
if (referenced)
{
// Increment link count for this program
if (mCLProgramLinkCounter.find(*program) == mCLProgramLinkCounter.end())
{
mCLProgramLinkCounter[*program] = 0;
}
++mCLProgramLinkCounter[*program];
// Setup the linked programs if this call is from capturing clCompileProgram or
// clLinkProgram
if (numLinkedPrograms)
{
mCLLinkedPrograms[*program] = std::vector<cl_program>();
for (cl_uint i = 0; i < numLinkedPrograms; ++i)
{
mCLLinkedPrograms[*program].push_back(linkedPrograms[i]);
}
}
// Go through the linked programs and increment their link counts
for (size_t i = 0; mCLLinkedPrograms.find(*program) != mCLLinkedPrograms.end() &&
i < mCLLinkedPrograms[*program].size();
++i)
{
trackCLProgramUpdate(&mCLLinkedPrograms[*program].at(i), true, 0, nullptr);
}
}
else
{
// Decrement link count for this program and the linked programs
--mCLProgramLinkCounter[*program];
for (size_t i = 0; mCLLinkedPrograms.find(*program) != mCLLinkedPrograms.end() &&
i < mCLLinkedPrograms[*program].size();
++i)
{
trackCLProgramUpdate(&mCLLinkedPrograms[*program].at(i), false, 0, nullptr);
}
// Remove the calls containing this object if the link count is 0
if (mCLProgramLinkCounter[*program] == 0)
{
mCLProgramLinkCounter.erase(*program);
if (mCLLinkedPrograms.find(*program) != mCLLinkedPrograms.end())
{
mCLLinkedPrograms.erase(*program);
}
if (removeUnneededOpenCLCalls)
{
removeCLProgramOccurrences(program, &mFrameCalls);
}
}
}
}
void FrameCaptureShared::injectMemcpy(void *src,
void *dest,
size_t size,
std::vector<CallCapture> *calls)
{
// Inject memcpy call before unmap
// Create param buffer
ParamBuffer paramBuffer;
// Create dest parameter
ParamCapture destParam("dest", ParamType::TvoidConstPointer);
InitParamValue(ParamType::TvoidPointer, dest, &destParam.value);
paramBuffer.addParam(std::move(destParam));
// Create src param
ParamCapture updateMemory("src", ParamType::TvoidConstPointer);
CaptureMemory(src, size, &updateMemory);
paramBuffer.addParam(std::move(updateMemory));
paramBuffer.addValueParam<size_t>("size", ParamType::Tsize_t, size);
calls->emplace(calls->end() - 1, "std::memcpy", std::move(paramBuffer));
}
void FrameCaptureShared::captureUpdateCLObjs(std::vector<CallCapture> *calls)
{
std::vector<cl_mem> &mCLDirtyMem = mResourceTrackerCL.mCLDirtyMem;
std::vector<void *> &mCLDirtySVM = mResourceTrackerCL.mCLDirtySVM;
cl_command_queue &mCLCurrentCommandQueue = mResourceTrackerCL.mCLCurrentCommandQueue;
for (uint32_t i = 0; i < mCLDirtyMem.size(); ++i)
{
cl_mem_object_type memType;
if (IsError(mCLDirtyMem.at(i)->cast<cl::Memory>().getInfo(
cl::MemInfo::Type, sizeof(cl_mem_object_type), &memType, nullptr)))
{
continue;
}
if (memType == CL_MEM_OBJECT_BUFFER)
{
void *ptr;
if (calls->back().entryPoint == EntryPoint::CLEnqueueUnmapMemObject)
{
CallCapture *mapCall = &mResourceTrackerCL.mCLMapCall.at(
calls->back()
.params.getParam("mapped_ptr", ParamType::TvoidPointer, 2)
.value.voidPointerVal);
size_t offset =
mapCall->params.getParam("offset", ParamType::Tsize_t, 4).value.size_tVal;
size_t size =
mapCall->params.getParam("size", ParamType::Tsize_t, 5).value.size_tVal;
ptr = malloc(size);
// Call clEnqueueReadBuffer to get the current data in the buffer
EnqueueReadBuffer(mCLCurrentCommandQueue, mCLDirtyMem.at(i), true, offset, size,
ptr, 0, nullptr, nullptr);
// Inject memcpy call BEFORE unmap
injectMemcpy(ptr,
calls->back()
.params.getParam("mapped_ptr", ParamType::TvoidPointer, 2)
.value.voidPointerVal,
size, calls);
}
else
{
size_t bufferSize = mCLDirtyMem.at(i)->cast<cl::Buffer>().getSize();
ptr = malloc(bufferSize);
// Call clEnqueueReadBuffer to get the current data in the buffer
EnqueueReadBuffer(mCLCurrentCommandQueue, mCLDirtyMem.at(i), true, 0, bufferSize,
ptr, 0, nullptr, nullptr);
// Pretend that a "clEnqueueWriteBuffer" was called with the above data retrieved
calls->push_back(CaptureEnqueueWriteBuffer(true, mCLCurrentCommandQueue,
mCLDirtyMem.at(i), true, 0, bufferSize,
ptr, 0, nullptr, nullptr, CL_SUCCESS));
// Implicit release, so going into the starting frame the buffer has the correct
// reference count
mCLDirtyMem.at(i)->cast<cl::Memory>().release();
}
free(ptr);
}
else if (memType == CL_MEM_OBJECT_PIPE)
{
UNIMPLEMENTED();
}
else
{
cl::Image *clImg = &mCLDirtyMem.at(i)->cast<cl::Image>();
void *ptr;
if (calls->back().entryPoint == EntryPoint::CLEnqueueUnmapMemObject)
{
CallCapture *mapCall = &mResourceTrackerCL.mCLMapCall.at(
calls->back()
.params.getParam("mapped_ptr", ParamType::TvoidPointer, 2)
.value.voidPointerVal);
const size_t *origin = (const size_t *)mapCall->params
.getParam("origin", ParamType::Tsize_tConstPointer, 4)
.data.back()
.data();
const size_t *region = (const size_t *)mapCall->params
.getParam("region", ParamType::Tsize_tConstPointer, 5)
.data.back()
.data();
size_t rowPitch = mapCall->params.getParam("image_row_pitch", ParamType::Tsize_t, 6)
.value.size_tVal;
size_t slicePitch =
mapCall->params.getParam("image_slice_pitch", ParamType::Tsize_t, 7)
.value.size_tVal;
// Get the image size to allocate the size of ptr
size_t totalSize = (region[2] - 1) * slicePitch + (region[1] - 1) * rowPitch +
region[0] * clImg->getElementSize();
ptr = malloc(totalSize);
// Call clEnqueueReadBuffer to get the current data in the image
EnqueueReadImage(mCLCurrentCommandQueue, mCLDirtyMem.at(i), true, origin, region,
rowPitch, slicePitch, ptr, 0, nullptr, nullptr);
// Inject memcpy call BEFORE unmap
injectMemcpy(ptr,
calls->back()
.params.getParam("mapped_ptr", ParamType::TvoidPointer, 2)
.value.voidPointerVal,
totalSize, calls);
}
else
{
ptr = malloc(clImg->getSize());
size_t origin[3] = {0, 0, 0};
size_t region[3] = {clImg->getWidth(), clImg->getHeight(), clImg->getDepth()};
// Call clEnqueueReadBuffer to get the current data in the image
EnqueueReadImage(mCLCurrentCommandQueue, mCLDirtyMem.at(i), true, origin, region,
clImg->getRowSize(), clImg->getSliceSize(), ptr, 0, nullptr,
nullptr);
// Pretend that a "clEnqueueWriteImage" was called with the above data retrieved
calls->push_back(CaptureEnqueueWriteImage(
true, mCLCurrentCommandQueue, mCLDirtyMem.at(i), true, origin, region,
clImg->getRowSize(), clImg->getSliceSize(), ptr, 0, nullptr, nullptr,
CL_SUCCESS));
// Implicit release, so going into the starting frame the buffer has the correct
// reference count
mCLDirtyMem.at(i)->cast<cl::Memory>().release();
}
free(ptr);
}
}
for (uint32_t i = 0; i < mCLDirtySVM.size(); ++i)
{
size_t SVMSize = mResourceTrackerCL.SVMToSize[mCLDirtySVM.at(i)];
// Call clEnqueueSVMMap to get the current data in the SVM pointer
cl::MemFlags flags;
flags.set(CL_MAP_READ);
EnqueueSVMMap(mCLCurrentCommandQueue, true, flags, mCLDirtySVM.at(i), SVMSize, 0, nullptr,
nullptr);
// Pretend that a "clEnqueueSVMMemcpy" was called with the above data retrieved
calls->push_back(CaptureEnqueueSVMMemcpy(true, mCLCurrentCommandQueue, true,
mCLDirtySVM.at(i), mCLDirtySVM.at(i), SVMSize, 0,
nullptr, nullptr, CL_SUCCESS));
// Call clEnqueueSVMUnmap to get the current data in the SVM pointer
EnqueueSVMUnmap(mCLCurrentCommandQueue, mResourceTrackerCL.mCLDirtySVM.at(i), 0, nullptr,
nullptr);
}
mCLDirtyMem.clear();
mCLDirtySVM.clear();
}
void FrameCaptureShared::removeCLMemOccurrences(const cl_mem *mem, std::vector<CallCapture> *calls)
{
// This function gets called when it captures a clReleaseMemObj prior to the starting frame
// that sets the reference count to 0, meaning that this cl_mem object isn't necessary for
// the wanted frames. So, we can remove the calls that use it.
for (size_t i = 0; i < calls->size(); ++i)
{
CallCapture *call = &calls->at(i);
cl_mem foundMem;
switch (call->entryPoint)
{
case EntryPoint::CLCreateBuffer:
case EntryPoint::CLCreateBufferWithProperties:
case EntryPoint::CLCreateImage:
case EntryPoint::CLCreateImageWithProperties:
case EntryPoint::CLCreateImage2D:
case EntryPoint::CLCreateImage3D:
case EntryPoint::CLCreatePipe:
{
foundMem = call->params.getReturnValue().value.cl_memVal;
break;
}
case EntryPoint::CLCreateSubBuffer:
{
foundMem = call->params.getReturnValue().value.cl_memVal;
if (foundMem != *mem)
{
foundMem =
call->params.getParam("buffer", ParamType::Tcl_mem, 0).value.cl_memVal;
}
break;
}
case EntryPoint::CLEnqueueReadBuffer:
case EntryPoint::CLEnqueueWriteBuffer:
case EntryPoint::CLEnqueueReadBufferRect:
case EntryPoint::CLEnqueueWriteBufferRect:
case EntryPoint::CLEnqueueMapBuffer:
{
// Can get rid of these calls because the buffer is no longer needed
foundMem = call->params.getParam("buffer", ParamType::Tcl_mem, 1).value.cl_memVal;
break;
}
case EntryPoint::CLEnqueueReadImage:
case EntryPoint::CLEnqueueWriteImage:
case EntryPoint::CLEnqueueMapImage:
{
// Can get rid of these calls because the image is no longer needed
foundMem = call->params.getParam("image", ParamType::Tcl_mem, 1).value.cl_memVal;
break;
}
case EntryPoint::CLEnqueueCopyBuffer:
case EntryPoint::CLEnqueueCopyBufferRect:
case EntryPoint::CLEnqueueCopyImage:
case EntryPoint::CLEnqueueCopyBufferToImage:
case EntryPoint::CLEnqueueCopyImageToBuffer:
{
// Can get rid of these calls because the obj is no longer needed
std::string srcType = "src_";
srcType += ((call->entryPoint == EntryPoint::CLEnqueueCopyImage ||
call->entryPoint == EntryPoint::CLEnqueueCopyImageToBuffer)
? "image"
: "buffer");
std::string dstType = "dst_";
dstType += ((call->entryPoint == EntryPoint::CLEnqueueCopyImage ||
call->entryPoint == EntryPoint::CLEnqueueCopyBufferToImage)
? "image"
: "buffer");
foundMem =
call->params.getParam(srcType.c_str(), ParamType::Tcl_mem, 1).value.cl_memVal;
if (foundMem != *mem)
{
foundMem = call->params.getParam(dstType.c_str(), ParamType::Tcl_mem, 2)
.value.cl_memVal;
}
break;
}
case EntryPoint::CLReleaseMemObject:
case EntryPoint::CLRetainMemObject:
case EntryPoint::CLGetMemObjectInfo:
case EntryPoint::CLSetMemObjectDestructorCallback:
case EntryPoint::CLEnqueueUnmapMemObject:
{
foundMem =
call->params
.getParam("memobj", ParamType::Tcl_mem,
call->entryPoint == EntryPoint::CLEnqueueUnmapMemObject ? 1 : 0)
.value.cl_memVal;
break;
}
case EntryPoint::CLGetImageInfo:
{
foundMem = call->params.getParam("image", ParamType::Tcl_mem, 0).value.cl_memVal;
break;
}
case EntryPoint::CLSetKernelArg:
{
foundMem = call->params.getParam("arg_value", ParamType::TvoidConstPointer, 3)
.value.cl_memVal;
if (cl::Platform::GetDefault()->getFrameCaptureShared()->getIndex(&foundMem) ==
SIZE_MAX)
{
continue;
}
break;
}
// Leave commented until external memory is upstream
// case EntryPoint::CLEnqueueAcquireExternalMemObjectsKHR:
// case EntryPoint::CLEnqueueReleaseExternalMemObjectsKHR:
case EntryPoint::CLEnqueueMigrateMemObjects:
{
const cl_mem *memObjs =
call->params.getParam("mem_objects", ParamType::Tcl_memConstPointer, 2)
.value.cl_memConstPointerVal;
cl_uint numMemObjs =
call->params.getParam("num_mem_objects", ParamType::Tcl_uint, 1)
.value.cl_uintVal;
std::vector<cl_mem> newMemObjs;
for (cl_uint memObjIndex = 0; i < numMemObjs; ++i)
{
if (memObjs[memObjIndex] != *mem)
{
newMemObjs.push_back(memObjs[memObjIndex]);
}
}
// If all the mem objects used in this array are released, I can remove this call
if (newMemObjs.empty())
{
foundMem = *mem;
}
else
{
call->params.setValueParamAtIndex("num_mem_objects", ParamType::Tcl_uint,
newMemObjs.size(), 1);
setCLObjVectorMap(
newMemObjs.data(), newMemObjs.size(),
&call->params.getParam("mem_objects", ParamType::Tcl_memConstPointer, 2),
&angle::FrameCaptureShared::getIndex);
continue;
}
break;
}
default:
continue;
}
if (foundMem == *mem)
{
removeCLCall(calls, i);
--i;
}
}
}
void FrameCaptureShared::removeCLKernelOccurrences(const cl_kernel *kernel,
std::vector<CallCapture> *calls)
{
// This function gets called when it captures a clReleaseProgram prior to the starting frame
// that sets the program's reference count to 0. This ensures that the kernels in that program
// are/should be released as well, meaning that this cl_kernel object isn't necessary for
// the wanted frames. So, we can remove the calls that use it.
// We cannot remove cl_kernel occurrences at the time of clReleaseKernel because the kernel may
// be an input to clCloneKernel and clCreateKernelsInProgram.
for (size_t i = 0; i < calls->size(); ++i)
{
CallCapture *call = &calls->at(i);
cl_kernel foundKernel;
switch (call->entryPoint)
{
case EntryPoint::CLCreateKernel:
{
foundKernel = call->params.getReturnValue().value.cl_kernelVal;
break;
}
case EntryPoint::CLCloneKernel:
{
foundKernel = call->params.getReturnValue().value.cl_kernelVal;
if (foundKernel != *kernel)
{
foundKernel = call->params.getParam("source_kernel", ParamType::Tcl_kernel, 0)
.value.cl_kernelVal;
}
break;
}
case EntryPoint::CLRetainKernel:
case EntryPoint::CLReleaseKernel:
case EntryPoint::CLSetKernelArg:
case EntryPoint::CLSetKernelArgSVMPointer:
case EntryPoint::CLSetKernelExecInfo:
case EntryPoint::CLGetKernelInfo:
case EntryPoint::CLGetKernelArgInfo:
case EntryPoint::CLGetKernelWorkGroupInfo:
case EntryPoint::CLGetKernelSubGroupInfo:
{
foundKernel =
call->params.getParam("kernel", ParamType::Tcl_kernel, 0).value.cl_kernelVal;
break;
}
case EntryPoint::CLEnqueueNDRangeKernel:
case EntryPoint::CLEnqueueTask:
{
foundKernel =
call->params.getParam("kernel", ParamType::Tcl_kernel, 1).value.cl_kernelVal;
break;
}
default:
continue;
}
if (foundKernel == *kernel)
{
removeCLCall(calls, i);
--i;
}
}
}
void FrameCaptureShared::removeCLProgramOccurrences(const cl_program *program,
std::vector<CallCapture> *calls)
{
// This function gets called when it captures a clReleaseMemObj prior to the starting frame
// that sets the reference count to 0, and the program is not linked to any other program,
// meaning that this cl_mem object isn't necessary for the wanted frames. So, we can
// remove the calls that use it.
for (size_t i = 0; i < calls->size(); ++i)
{
CallCapture *call = &calls->at(i);
cl_program foundProgram;
switch (call->entryPoint)
{
case EntryPoint::CLCreateProgramWithSource:
case EntryPoint::CLCreateProgramWithBinary:
case EntryPoint::CLCreateProgramWithBuiltInKernels:
case EntryPoint::CLCreateProgramWithIL:
case EntryPoint::CLLinkProgram:
{
foundProgram = call->params.getReturnValue().value.cl_programVal;
break;
}
case EntryPoint::CLRetainProgram:
case EntryPoint::CLReleaseProgram:
case EntryPoint::CLBuildProgram:
case EntryPoint::CLGetProgramInfo:
case EntryPoint::CLGetProgramBuildInfo:
case EntryPoint::CLCreateKernel:
case EntryPoint::CLCreateKernelsInProgram:
case EntryPoint::CLUnloadPlatformCompiler:
case EntryPoint::CLCompileProgram:
{
uint8_t programIndex = call->entryPoint == EntryPoint::CLCompileProgram ? 1 : 0;
foundProgram =
call->params.getParam("program", ParamType::Tcl_program, programIndex)
.value.cl_programVal;
break;
}
default:
continue;
}
if (foundProgram == *program)
{
removeCLCall(calls, i);
--i;
}
}
if (mResourceTrackerCL.mCLProgramToKernels.find(*program) !=
mResourceTrackerCL.mCLProgramToKernels.end())
{
for (size_t i = 0; i < mResourceTrackerCL.mCLProgramToKernels[*program].size(); ++i)
{
removeCLKernelOccurrences(&mResourceTrackerCL.mCLProgramToKernels[*program].at(i),
calls);
}
mResourceTrackerCL.mCLProgramToKernels.erase(*program);
}
}
void FrameCaptureShared::removeCLCall(std::vector<CallCapture> *callVector, size_t &callIndex)
{
CallCapture *call = &callVector->at(callIndex);
const std::vector<ParamCapture> *params = &call->params.getParamCaptures();
cl_context context = nullptr;
// Checks if there is an event that is implicitly created during the deleted call.
// If there is, need to inject a clCreateUserEvent call and a clSetUserEventStatus call.
for (auto &param : *params)
{
if (param.type == ParamType::Tcl_context)
{
context = param.value.cl_contextVal;
}
else if (param.type == ParamType::Tcl_command_queue)
{
context =
param.value.cl_command_queueVal->cast<cl::CommandQueue>().getContext().getNative();
}
else if (param.type == ParamType::Tcl_eventPointer && param.value.cl_eventVal != nullptr &&
context)
{
// Capture the creation of a successful event if the CL call being removed created an
// event (ex: clEnqueueReadBuffer)
cl_event event = param.value.cl_eventVal;
callVector->insert(callVector->begin() + callIndex,
CaptureSetUserEventStatus(true, event, CL_COMPLETE, CL_SUCCESS));
callVector->insert(callVector->begin() + callIndex,
CaptureCreateUserEvent(true, context, nullptr, event));
callIndex += 2;
break;
}
}
callVector->erase(callVector->begin() + callIndex);
}
void FrameCaptureShared::maybeCapturePreCallUpdatesCL(CallCapture &call)
{
switch (call.entryPoint)
{
case EntryPoint::CLGetExtensionFunctionAddress:
case EntryPoint::CLGetExtensionFunctionAddressForPlatform:
{
uint32_t index = call.entryPoint == EntryPoint::CLGetExtensionFunctionAddress ? 0 : 1;
std::string funcName =
(const char *)call.params.getParam("func_name", ParamType::TcharConstPointer, index)
.value.charConstPointerPointerVal;
call.customFunctionName =
funcName + " = (" + funcName + "_fn)" + GetEntryPointName(call.entryPoint);
if (std::find(mExtFuncsAdded.begin(), mExtFuncsAdded.end(), funcName) ==
mExtFuncsAdded.end())
{
mExtFuncsAdded.push_back(funcName);
}
break;
}
case EntryPoint::CLCreateContext:
case EntryPoint::CLCreateContextFromType:
{
if (call.params.getParam("properties", ParamType::Tcl_context_propertiesConstPointer, 0)
.value.cl_context_propertiesConstPointerVal)
{
size_t propSize = 0;
size_t platformIDIndex = 0;
const cl_context_properties *propertiesData =
(cl_context_properties *)call.params
.getParam("properties", ParamType::Tcl_context_propertiesConstPointer, 0)
.data[0]
.data();
while (propertiesData[propSize] != 0)
{
if (propertiesData[propSize] == CL_CONTEXT_PLATFORM)
{
// "Each property name is immediately followed by the corresponding desired
// value"
platformIDIndex = propSize + 1;
}
++propSize;
}
++propSize;
if (platformIDIndex == 0)
{
ParamBuffer params;
params.addValueParam("propSize", ParamType::Tsize_t, propSize);
ParamCapture propertiesParam("propData",
ParamType::Tcl_context_propertiesConstPointer);
InitParamValue(ParamType::Tcl_context_propertiesConstPointer, propertiesData,
&propertiesParam.value);
CaptureMemory(propertiesData, propSize * sizeof(cl_context_properties),
&propertiesParam);
params.addParam(std::move(propertiesParam));
mFrameCalls.emplace_back(
CallCapture("UpdateCLContextPropertiesNoPlatform", std::move(params)));
call.params
.getParam("properties", ParamType::Tcl_context_propertiesConstPointer, 0)
.data.clear();
break;
}
// Create call to UpdateCLContextProperties
ParamBuffer params;
params.addValueParam("propSize", ParamType::Tsize_t, propSize);
ParamCapture propertiesParam("propData",
ParamType::Tcl_context_propertiesConstPointer);
InitParamValue(ParamType::Tcl_context_propertiesConstPointer, propertiesData,
&propertiesParam.value);
CaptureMemory(propertiesData, propSize * sizeof(cl_context_properties),
&propertiesParam);
params.addParam(std::move(propertiesParam));
params.addValueParam("platformIdxInProps", ParamType::Tsize_t, platformIDIndex);
params.addValueParam("platformIdxInMap", ParamType::Tsize_t,
getIndex((cl_platform_id *)&propertiesData[platformIDIndex]));
call.params.getParam("properties", ParamType::Tcl_context_propertiesConstPointer, 0)
.data.clear();
mFrameCalls.emplace_back(
CallCapture("UpdateCLContextPropertiesWithPlatform", std::move(params)));
}
break;
}
default:
break;
}
updateReadBufferSize(call.params.getReadBufferSize());
}
void FrameCaptureShared::addCLResetObj(const ParamCapture &param)
{
mResourceTrackerCL.mCLResetObjs.push_back(angle::ParamCapture("resetObj", param.type));
auto paramValue =
&mResourceTrackerCL.mCLResetObjs.at(mResourceTrackerCL.mCLResetObjs.size() - 1).value;
switch (param.type)
{
case ParamType::Tcl_device_id:
InitParamValue(param.type, param.value.cl_device_idVal, paramValue);
break;
case ParamType::Tcl_mem:
InitParamValue(param.type, param.value.cl_memVal, paramValue);
break;
case ParamType::Tcl_kernel:
InitParamValue(param.type, param.value.cl_kernelVal, paramValue);
break;
case ParamType::Tcl_program:
InitParamValue(param.type, param.value.cl_programVal, paramValue);
break;
case ParamType::Tcl_command_queue:
InitParamValue(param.type, param.value.cl_command_queueVal, paramValue);
break;
case ParamType::Tcl_context:
InitParamValue(param.type, param.value.cl_contextVal, paramValue);
break;
case ParamType::Tcl_sampler:
InitParamValue(param.type, param.value.cl_samplerVal, paramValue);
break;
case ParamType::Tcl_event:
InitParamValue(param.type, param.value.cl_eventVal, paramValue);
break;
default:
break;
}
}
void FrameCaptureShared::removeCLResetObj(const ParamCapture &param)
{
std::vector<ParamCapture> &mCLResetObjs = mResourceTrackerCL.mCLResetObjs;
for (size_t i = 0; i < mCLResetObjs.size(); ++i)
{
bool foundCLObj =
param.type == mCLResetObjs.at(i).type &&
((param.type == ParamType::Tcl_device_id &&
param.value.cl_device_idVal == mCLResetObjs.at(i).value.cl_device_idVal) ||
(param.type == ParamType::Tcl_mem &&
param.value.cl_memVal == mCLResetObjs.at(i).value.cl_memVal) ||
(param.type == ParamType::Tcl_kernel &&
param.value.cl_kernelVal == mCLResetObjs.at(i).value.cl_kernelVal) ||
(param.type == ParamType::Tcl_program &&
param.value.cl_programVal == mCLResetObjs.at(i).value.cl_programVal) ||
(param.type == ParamType::Tcl_command_queue &&
param.value.cl_command_queueVal == mCLResetObjs.at(i).value.cl_command_queueVal) ||
(param.type == ParamType::Tcl_context &&
param.value.cl_contextVal == mCLResetObjs.at(i).value.cl_contextVal) ||
(param.type == ParamType::Tcl_sampler &&
param.value.cl_samplerVal == mCLResetObjs.at(i).value.cl_samplerVal) ||
(param.type == ParamType::Tcl_event &&
param.value.cl_eventVal == mCLResetObjs.at(i).value.cl_eventVal));
if (foundCLObj)
{
mCLResetObjs.erase(mCLResetObjs.begin() + i);
break;
}
}
}
void FrameCaptureShared::printCLResetObjs(std::stringstream &stream)
{
std::vector<ParamCapture> &mCLResetObjs = mResourceTrackerCL.mCLResetObjs;
for (size_t i = 0; i < mCLResetObjs.size(); ++i)
{
stream << " ";
switch (mCLResetObjs.at(i).type)
{
case ParamType::Tcl_device_id:
stream << "clReleaseDevice(clDevicesMap["
<< std::to_string(getIndex(&mCLResetObjs.at(i).value.cl_device_idVal))
<< "]);";
break;
case ParamType::Tcl_mem:
stream << "clReleaseMemObject(clMemMap["
<< std::to_string(getIndex(&mCLResetObjs.at(i).value.cl_memVal)) << "]);";
break;
case ParamType::Tcl_kernel:
stream << "clReleaseKernel(clKernelsMap["
<< std::to_string(getIndex(&mCLResetObjs.at(i).value.cl_kernelVal)) << "]);";
break;
case ParamType::Tcl_program:
stream << "clReleaseProgram(clProgramsMap["
<< std::to_string(getIndex(&mCLResetObjs.at(i).value.cl_programVal))
<< "]);";
break;
case ParamType::Tcl_command_queue:
stream << "clReleaseCommandQueue(clCommandQueuesMap["
<< std::to_string(getIndex(&mCLResetObjs.at(i).value.cl_command_queueVal))
<< "]);";
break;
case ParamType::Tcl_context:
stream << "clReleaseContext(clContextsMap["
<< std::to_string(getIndex(&mCLResetObjs.at(i).value.cl_contextVal))
<< "]);";
break;
case ParamType::Tcl_sampler:
stream << "clReleaseSampler(clSamplersMap["
<< std::to_string(getIndex(&mCLResetObjs.at(i).value.cl_samplerVal))
<< "]);";
break;
case ParamType::Tcl_event:
stream << "clReleaseEvent(clEventsMap["
<< std::to_string(getIndex(&mCLResetObjs.at(i).value.cl_eventVal)) << "]);";
break;
default:
break;
}
stream << "\n";
}
}
void FrameCaptureShared::updateResourceCountsFromParamCaptureCL(const ParamCapture &param,
const CallCapture &call)
{
switch (param.type)
{
case ParamType::Tcl_platform_idPointer:
if (call.entryPoint == EntryPoint::CLIcdGetPlatformIDsKHR ||
call.entryPoint == EntryPoint::CLGetPlatformIDs)
{
mMaxCLParamsSize[param.type] +=
sizeof(cl_platform_id) * ((call.params.getParamCaptures()[0]).value.cl_uintVal);
}
break;
case ParamType::Tcl_device_idPointer:
if (call.entryPoint == EntryPoint::CLGetDeviceIDs)
{
mMaxCLParamsSize[param.type] +=
sizeof(cl_device_id) * ((call.params.getParamCaptures()[2]).value.cl_uintVal);
}
break;
case ParamType::Tcl_context:
if (call.entryPoint == EntryPoint::CLCreateContext ||
call.entryPoint == EntryPoint::CLCreateContextFromType)
{
if ((getIndex(&param.value.cl_contextVal) + 1) * sizeof(cl_context) >
mMaxCLParamsSize[param.type])
{
mMaxCLParamsSize[param.type] =
(uint32_t)((getIndex(&param.value.cl_contextVal) + 1) * sizeof(cl_context));
}
addCLResetObj(param);
}
break;
case ParamType::Tcl_command_queue:
if (call.entryPoint == EntryPoint::CLCreateCommandQueueWithProperties ||
call.entryPoint == EntryPoint::CLCreateCommandQueue)
{
if ((getIndex(&param.value.cl_command_queueVal) + 1) * sizeof(cl_command_queue) >
mMaxCLParamsSize[param.type])
{
mMaxCLParamsSize[param.type] =
(uint32_t)((getIndex(&param.value.cl_command_queueVal) + 1) *
sizeof(cl_command_queue));
}
addCLResetObj(param);
}
break;
case ParamType::Tcl_mem:
if (call.entryPoint == EntryPoint::CLCreateBufferWithProperties ||
call.entryPoint == EntryPoint::CLCreateBuffer ||
call.entryPoint == EntryPoint::CLCreateSubBuffer ||
call.entryPoint == EntryPoint::CLCreateImageWithProperties ||
call.entryPoint == EntryPoint::CLCreateImage ||
call.entryPoint == EntryPoint::CLCreateImage2D ||
call.entryPoint == EntryPoint::CLCreateImage3D)
{
if ((getIndex(&param.value.cl_memVal) + 1) * sizeof(cl_mem) >
mMaxCLParamsSize[param.type])
{
mMaxCLParamsSize[param.type] =
(uint32_t)((getIndex(&param.value.cl_memVal) + 1) * sizeof(cl_mem));
}
addCLResetObj(param);
}
break;
case ParamType::Tcl_eventPointer:
{
if (param.value.cl_eventVal)
{
if ((getIndex(&param.value.cl_eventVal) + 1) * sizeof(cl_event) >
mMaxCLParamsSize[param.type])
{
mMaxCLParamsSize[param.type] =
(uint32_t)((getIndex(&param.value.cl_eventVal) + 1) * sizeof(cl_event));
}
angle::ParamCapture eventParam("event", angle::ParamType::Tcl_event);
InitParamValue(angle::ParamType::Tcl_event, param.value.cl_eventVal,
&eventParam.value);
addCLResetObj(eventParam);
}
break;
}
case ParamType::Tcl_program:
if (call.entryPoint == EntryPoint::CLCreateProgramWithSource ||
call.entryPoint == EntryPoint::CLCreateProgramWithBinary ||
call.entryPoint == EntryPoint::CLCreateProgramWithBuiltInKernels ||
call.entryPoint == EntryPoint::CLLinkProgram ||
call.entryPoint == EntryPoint::CLCreateProgramWithIL)
{
if ((getIndex(&param.value.cl_programVal) + 1) * sizeof(cl_program) >
mMaxCLParamsSize[param.type])
{
mMaxCLParamsSize[param.type] =
(uint32_t)((getIndex(&param.value.cl_programVal) + 1) * sizeof(cl_program));
}
addCLResetObj(param);
}
break;
case ParamType::Tcl_kernel:
if (call.entryPoint == EntryPoint::CLCreateKernel ||
call.entryPoint == EntryPoint::CLCloneKernel)
{
if ((getIndex(&param.value.cl_kernelVal) + 1) * sizeof(cl_kernel) >
mMaxCLParamsSize[param.type])
{
mMaxCLParamsSize[param.type] =
(uint32_t)((getIndex(&param.value.cl_kernelVal) + 1) * sizeof(cl_kernel));
}
addCLResetObj(param);
}
break;
case ParamType::Tcl_sampler:
if (call.entryPoint == EntryPoint::CLCreateSampler ||
call.entryPoint == EntryPoint::CLCreateSamplerWithProperties)
{
if ((getIndex(&param.value.cl_samplerVal) + 1) * sizeof(cl_sampler) >
mMaxCLParamsSize[param.type])
{
mMaxCLParamsSize[param.type] =
(uint32_t)((getIndex(&param.value.cl_samplerVal) + 1) * sizeof(cl_sampler));
}
addCLResetObj(param);
}
break;
case ParamType::TvoidPointer:
if (call.entryPoint == EntryPoint::CLEnqueueMapImage ||
call.entryPoint == EntryPoint::CLEnqueueMapBuffer)
{
mMaxCLParamsSize[param.type] += sizeof(void *);
}
break;
default:
break;
}
}
void FrameCaptureShared::updateResourceCountsFromCallCaptureCL(const CallCapture &call)
{
for (const ParamCapture &param : call.params.getParamCaptures())
{
updateResourceCountsFromParamCaptureCL(param, call);
}
// Update resource IDs in the return value.
switch (call.entryPoint)
{
case EntryPoint::CLCreateContext:
case EntryPoint::CLCreateContextFromType:
setIndex(&call.params.getReturnValue().value.cl_contextVal);
updateResourceCountsFromParamCaptureCL(call.params.getReturnValue(), call);
break;
case EntryPoint::CLCreateBuffer:
case EntryPoint::CLCreateBufferWithProperties:
case EntryPoint::CLCreateSubBuffer:
case EntryPoint::CLCreateImageWithProperties:
case EntryPoint::CLCreateImage:
case EntryPoint::CLCreateImage2D:
case EntryPoint::CLCreateImage3D:
case EntryPoint::CLCreatePipe:
setIndex(&call.params.getReturnValue().value.cl_memVal);
updateResourceCountsFromParamCaptureCL(call.params.getReturnValue(), call);
break;
case EntryPoint::CLCreateSampler:
case EntryPoint::CLCreateSamplerWithProperties:
setIndex(&call.params.getReturnValue().value.cl_samplerVal);
updateResourceCountsFromParamCaptureCL(call.params.getReturnValue(), call);
break;
case EntryPoint::CLCreateCommandQueue:
case EntryPoint::CLCreateCommandQueueWithProperties:
setIndex(&call.params.getReturnValue().value.cl_command_queueVal);
updateResourceCountsFromParamCaptureCL(call.params.getReturnValue(), call);
break;
case EntryPoint::CLCreateProgramWithSource:
case EntryPoint::CLCreateProgramWithBinary:
case EntryPoint::CLCreateProgramWithBuiltInKernels:
case EntryPoint::CLLinkProgram:
case EntryPoint::CLCreateProgramWithIL:
setIndex(&call.params.getReturnValue().value.cl_programVal);
updateResourceCountsFromParamCaptureCL(call.params.getReturnValue(), call);
break;
case EntryPoint::CLCreateKernel:
case EntryPoint::CLCloneKernel:
setIndex(&call.params.getReturnValue().value.cl_kernelVal);
updateResourceCountsFromParamCaptureCL(call.params.getReturnValue(), call);
break;
case EntryPoint::CLEnqueueMapBuffer:
case EntryPoint::CLEnqueueMapImage:
case EntryPoint::CLSVMAlloc:
if (call.params.getReturnValue().value.voidPointerVal)
{
setCLVoidIndex(call.params.getReturnValue().value.voidPointerVal);
updateResourceCountsFromParamCaptureCL(call.params.getReturnValue(), call);
}
break;
case EntryPoint::CLCreateUserEvent:
setIndex(&call.params.getReturnValue().value.cl_eventVal);
break;
case EntryPoint::CLReleaseDevice:
case EntryPoint::CLReleaseCommandQueue:
case EntryPoint::CLReleaseContext:
case EntryPoint::CLReleaseEvent:
case EntryPoint::CLReleaseKernel:
case EntryPoint::CLReleaseMemObject:
case EntryPoint::CLReleaseProgram:
case EntryPoint::CLReleaseSampler:
removeCLResetObj(call.params.getParamCaptures()[0]);
break;
default:
break;
}
}
void FrameCaptureShared::captureCLCall(CallCapture &&inCall, bool isCallValid)
{
if (!mCallCaptured)
{
mReplayWriter.captureAPI = CaptureAPI::CL;
mBinaryData.clear();
mCallCaptured = true;
std::atexit(onCLProgramEnd);
}
if (mFrameIndex <= mCaptureEndFrame)
{
// Keep track of return values from OpenCL calls
updateResourceCountsFromCallCaptureCL(inCall);
// Set to true if the call signifies the end of a frame
// ex: clEnqueueNDRangeKernel
bool frameEnd = false;
// Covers pre call updates, like updating the read buffer size
maybeCapturePreCallUpdatesCL(inCall);
// If it's an unnecessary call for replay (ex: clGetDeviceInfo)
if (mCLOptionalCalls.find(inCall.entryPoint) == mCLOptionalCalls.end())
{
if (mCLEndFrameCalls.find(inCall.entryPoint) != mCLEndFrameCalls.end())
{
frameEnd = true;
}
mFrameCalls.emplace_back(std::move(inCall));
}
else
{
saveCLGetInfo(inCall);
return;
}
// For kernel argument memory snapshots
maybeCapturePostCallUpdatesCL();
if (mFrameIndex >= mCaptureStartFrame ||
(mFrameIndex + 1 == mCaptureStartFrame && frameEnd))
{
// Maybe add clEnqueueWrite* or memcpy for memory snapshots
captureUpdateCLObjs(&mFrameCalls);
}
if (frameEnd && mFrameIndex >= mCaptureStartFrame)
{
mActiveFrameIndices.push_back(mFrameIndex);
writeMainContextCppReplayCL();
if (mFrameIndex == mCaptureEndFrame)
{
writeCppReplayIndexFilesCL();
SaveBinaryData(mCompression, mOutDirectory, kNoContextId, mCaptureLabel,
mBinaryData);
}
reset();
}
if (frameEnd)
{
if (mFrameIndex == (mCaptureStartFrame == 0 ? 0 : mCaptureStartFrame - 1))
{
mCLSetupCalls = std::move(mFrameCalls);
}
++mFrameIndex;
}
}
}
void FrameCaptureShared::maybeCapturePostCallUpdatesCL()
{
CallCapture &lastCall = mFrameCalls.back();
switch (lastCall.entryPoint)
{
case EntryPoint::CLEnqueueMapBuffer:
{
// Recreate the map call to store in the mCLMapCall unordered_map
// so later upon the unmap call, the original map data will be available
cl_command_queue command_queue =
lastCall.params.getParam("command_queue", ParamType::Tcl_command_queue, 0)
.value.cl_command_queueVal;
cl_mem buffer =
lastCall.params.getParam("buffer", ParamType::Tcl_mem, 1).value.cl_memVal;
cl_bool blocking_map =
lastCall.params.getParam("blocking_map", ParamType::Tcl_bool, 2).value.cl_boolVal;
cl::MapFlags map_flags =
lastCall.params.getParam("map_flagsPacked", ParamType::TMapFlags, 3)
.value.MapFlagsVal;
size_t offset =
lastCall.params.getParam("offset", ParamType::Tsize_t, 4).value.size_tVal;
size_t size = lastCall.params.getParam("size", ParamType::Tsize_t, 5).value.size_tVal;
mResourceTrackerCL.mCLMapCall.emplace(
lastCall.params.getReturnValue().value.voidPointerVal,
CaptureEnqueueMapBuffer(true, command_queue, buffer, blocking_map, map_flags,
offset, size, 0, nullptr, nullptr, nullptr, nullptr));
break;
}
case EntryPoint::CLEnqueueMapImage:
{
// Recreate the map call to store in the mCLMapCall unordered_map
// so later upon the unmap call, the original map data will be available
cl_command_queue command_queue =
lastCall.params.getParam("command_queue", ParamType::Tcl_command_queue, 0)
.value.cl_command_queueVal;
cl_mem image = lastCall.params.getParam("image", ParamType::Tcl_mem, 1).value.cl_memVal;
cl_bool blocking_map =
lastCall.params.getParam("blocking_map", ParamType::Tcl_bool, 2).value.cl_boolVal;
cl::MapFlags map_flags =
lastCall.params.getParam("map_flagsPacked", ParamType::TMapFlags, 3)
.value.MapFlagsVal;
const size_t *origin =
lastCall.params.getParam("origin", ParamType::Tsize_tConstPointer, 4)
.value.size_tConstPointerVal;
const size_t *region =
lastCall.params.getParam("region", ParamType::Tsize_tConstPointer, 5)
.value.size_tConstPointerVal;
size_t *image_row_pitch =
lastCall.params.getParam("image_row_pitch", ParamType::Tsize_tPointer, 6)
.value.size_tPointerVal;
size_t *image_slice_pitch =
lastCall.params.getParam("image_slice_pitch", ParamType::Tsize_tPointer, 7)
.value.size_tPointerVal;
mResourceTrackerCL.mCLMapCall.emplace(
lastCall.params.getReturnValue().value.voidPointerVal,
CaptureEnqueueMapImage(true, command_queue, image, blocking_map, map_flags, origin,
region, image_row_pitch, image_slice_pitch, 0, nullptr,
nullptr, nullptr, nullptr));
mResourceTrackerCL.mCLMapCall.at(lastCall.params.getReturnValue().value.voidPointerVal)
.params.setValueParamAtIndex("image_row_pitch", ParamType::Tsize_t,
*image_row_pitch, 6);
mResourceTrackerCL.mCLMapCall.at(lastCall.params.getReturnValue().value.voidPointerVal)
.params.setValueParamAtIndex("image_slice_pitch", ParamType::Tsize_t,
image_slice_pitch == nullptr ? 0 : *image_slice_pitch,
7);
break;
}
case EntryPoint::CLEnqueueUnmapMemObject:
{
if (mFrameIndex >= mCaptureStartFrame)
{
// Mark as dirty
cl_mem *mem =
&lastCall.params.getParam("memobj", ParamType::Tcl_mem, 1).value.cl_memVal;
mResourceTrackerCL.mCLCurrentCommandQueue =
lastCall.params.getParam("command_queue", ParamType::Tcl_command_queue, 0)
.value.cl_command_queueVal;
CallCapture *mapCall = &mResourceTrackerCL.mCLMapCall.at(
lastCall.params.getParam("mapped_ptr", ParamType::TvoidPointer, 2)
.value.voidPointerVal);
auto it = std::find(mResourceTrackerCL.mCLDirtyMem.begin(),
mResourceTrackerCL.mCLDirtyMem.end(), *mem);
if (it == mResourceTrackerCL.mCLDirtyMem.end() &&
mapCall->params.getParam("map_flagsPacked", ParamType::TMapFlags, 3)
.value.MapFlagsVal.mask(CL_MAP_WRITE |
CL_MAP_WRITE_INVALIDATE_REGION) != 0u)
{
mResourceTrackerCL.mCLDirtyMem.push_back(*mem);
}
}
break;
}
case EntryPoint::CLEnqueueSVMUnmap:
{
// Mark as dirty
void *svm = &lastCall.params.getParam("svm_ptr", ParamType::TvoidPointer, 1)
.value.voidPointerVal;
mResourceTrackerCL.mCLCurrentCommandQueue =
lastCall.params.getParam("command_queue", ParamType::Tcl_command_queue, 0)
.value.cl_command_queueVal;
mResourceTrackerCL.mCLDirtySVM.push_back(svm);
break;
}
default:
break;
}
// OpenCL calls that come before the starting frame
if (mFrameIndex < mCaptureStartFrame)
{
std::unordered_map<cl_program, std::vector<cl_kernel>> &mCLProgramToKernels =
mResourceTrackerCL.mCLProgramToKernels;
switch (lastCall.entryPoint)
{
// There should be no unnecessary enqueue functions prior to the starting frame.
// captureUpdateCLObjs accounts for it by dynamically adding
// CLEnqueueWriteBuffer/CLEnqueueWriteImage to ensure the cl_mem objects
// have the needed info upon replay
case EntryPoint::CLEnqueueNDRangeKernel:
case EntryPoint::CLEnqueueNativeKernel:
case EntryPoint::CLEnqueueTask:
case EntryPoint::CLEnqueueReadBuffer:
case EntryPoint::CLEnqueueWriteBuffer:
case EntryPoint::CLEnqueueReadBufferRect:
case EntryPoint::CLEnqueueWriteBufferRect:
case EntryPoint::CLEnqueueReadImage:
case EntryPoint::CLEnqueueWriteImage:
case EntryPoint::CLEnqueueCopyBuffer:
case EntryPoint::CLEnqueueCopyBufferRect:
case EntryPoint::CLEnqueueCopyImage:
case EntryPoint::CLEnqueueCopyBufferToImage:
case EntryPoint::CLEnqueueCopyImageToBuffer:
case EntryPoint::CLEnqueueFillBuffer:
case EntryPoint::CLEnqueueFillImage:
case EntryPoint::CLEnqueueWaitForEvents:
case EntryPoint::CLEnqueueMarkerWithWaitList:
case EntryPoint::CLEnqueueBarrierWithWaitList:
case EntryPoint::CLEnqueueBarrier:
case EntryPoint::CLEnqueueMarker:
case EntryPoint::CLEnqueueMigrateMemObjects:
case EntryPoint::CLEnqueueSVMMemcpy:
case EntryPoint::CLEnqueueSVMMemFill:
case EntryPoint::CLEnqueueSVMMigrateMem:
{
size_t index = mFrameCalls.size() - 1;
removeCLCall(&mFrameCalls, index);
break;
}
case EntryPoint::CLCreateBuffer:
case EntryPoint::CLCreateBufferWithProperties:
case EntryPoint::CLCreateImage:
case EntryPoint::CLCreateImageWithProperties:
case EntryPoint::CLCreateImage2D:
case EntryPoint::CLCreateImage3D:
case EntryPoint::CLCreatePipe:
case EntryPoint::CLCreateSubBuffer:
{
const cl_mem *newBuff = &lastCall.params.getReturnValue().value.cl_memVal;
// Set the parent
if (lastCall.entryPoint == EntryPoint::CLCreateSubBuffer)
{
cl_mem parent =
lastCall.params.getParam("buffer", ParamType::Tcl_mem, 0).value.cl_memVal;
mResourceTrackerCL.mCLSubBufferToParent[*newBuff] = parent;
}
// Implicit retain
(*newBuff)->cast<cl::Memory>().retain();
// Add buffer as tracked
trackCLMemUpdate(newBuff, true);
break;
}
case EntryPoint::CLReleaseMemObject:
{
// Potentially remove buffer/image (and potentially parents) as tracked
trackCLMemUpdate(
&lastCall.params.getParam("memobj", ParamType::Tcl_mem, 0).value.cl_memVal,
false);
break;
}
case EntryPoint::CLCreateCommandQueue:
case EntryPoint::CLCreateCommandQueueWithProperties:
{
mResourceTrackerCL.mCLCurrentCommandQueue =
lastCall.params.getReturnValue().value.cl_command_queueVal;
break;
}
case EntryPoint::CLCreateProgramWithSource:
case EntryPoint::CLCreateProgramWithBinary:
case EntryPoint::CLCreateProgramWithBuiltInKernels:
case EntryPoint::CLCreateProgramWithIL:
{
mCLProgramToKernels[lastCall.params.getReturnValue().value.cl_programVal] =
std::vector<cl_kernel>();
trackCLProgramUpdate(&lastCall.params.getReturnValue().value.cl_programVal, true, 0,
nullptr);
break;
}
case EntryPoint::CLRetainProgram:
{
trackCLProgramUpdate(&lastCall.params.getParam("program", ParamType::Tcl_program, 0)
.value.cl_programVal,
true, 0, nullptr);
break;
}
case EntryPoint::CLCompileProgram:
{
const cl_program *program =
&lastCall.params.getParam("program", ParamType::Tcl_program, 0)
.value.cl_programVal;
trackCLProgramUpdate(
program, true,
lastCall.params.getParam("num_input_headers", ParamType::Tcl_uint, 4)
.value.cl_uintVal,
lastCall.params.getParam("input_headers", ParamType::Tcl_programConstPointer, 5)
.value.cl_programConstPointerVal);
break;
}
case EntryPoint::CLLinkProgram:
{
const cl_program *program = &lastCall.params.getReturnValue().value.cl_programVal;
mCLProgramToKernels[*program] = std::vector<cl_kernel>();
trackCLProgramUpdate(
program, true,
lastCall.params.getParam("num_input_programs", ParamType::Tcl_uint, 4)
.value.cl_uintVal,
lastCall.params
.getParam("input_programs", ParamType::Tcl_programConstPointer, 5)
.value.cl_programConstPointerVal);
break;
}
case EntryPoint::CLReleaseProgram:
{
trackCLProgramUpdate(&lastCall.params.getParam("program", ParamType::Tcl_program, 0)
.value.cl_programVal,
false, 0, nullptr);
break;
}
case EntryPoint::CLCreateKernel:
{
cl_program *program =
&lastCall.params.getParam("program", ParamType::Tcl_program, 0)
.value.cl_programVal;
const cl_kernel *kernel = &lastCall.params.getReturnValue().value.cl_kernelVal;
mCLProgramToKernels[*program].push_back(*kernel);
mResourceTrackerCL.mCLKernelToProgram[*kernel] = *program;
break;
}
case EntryPoint::CLCloneKernel:
{
cl_program *program =
&mResourceTrackerCL.mCLKernelToProgram[lastCall.params
.getParam("source_kernel",
ParamType::Tcl_kernel, 0)
.value.cl_kernelVal];
const cl_kernel *kernel = &lastCall.params.getReturnValue().value.cl_kernelVal;
mCLProgramToKernels[*program].push_back(*kernel);
mResourceTrackerCL.mCLKernelToProgram[*kernel] = *program;
break;
}
case EntryPoint::CLSVMAlloc:
{
void *svm = lastCall.params.getReturnValue().value.voidPointerVal;
// Potentially mark as dirty
auto it = std::find(mResourceTrackerCL.mCLDirtySVM.begin(),
mResourceTrackerCL.mCLDirtySVM.end(), svm);
if (it == mResourceTrackerCL.mCLDirtySVM.end())
{
mResourceTrackerCL.mCLDirtySVM.push_back(svm);
}
break;
}
case EntryPoint::CLSVMFree:
{
void *svm = lastCall.params.getParam("svm_pointer", ParamType::TvoidPointer, 1)
.value.voidPointerVal;
auto it = std::find(mResourceTrackerCL.mCLDirtySVM.begin(),
mResourceTrackerCL.mCLDirtySVM.end(), svm);
if (it != mResourceTrackerCL.mCLDirtySVM.end())
{
mResourceTrackerCL.mCLDirtySVM.erase(it);
}
break;
}
case EntryPoint::CLEnqueueSVMFree:
{
for (cl_uint svmIndex = 0;
svmIndex < lastCall.params.getParam("num_svm_pointers", ParamType::Tcl_uint, 1)
.value.cl_uintVal;
++svmIndex)
{
void *svm =
lastCall.params.getParam("svm_pointers", ParamType::TvoidPointerPointer, 1)
.value.voidPointerPointerVal[svmIndex];
auto it = std::find(mResourceTrackerCL.mCLDirtySVM.begin(),
mResourceTrackerCL.mCLDirtySVM.end(), svm);
if (it != mResourceTrackerCL.mCLDirtySVM.end())
{
mResourceTrackerCL.mCLDirtySVM.erase(it);
}
}
break;
}
default:
break;
}
}
}
void FrameCaptureShared::onCLProgramEnd()
{
if (cl::Platform::GetDefault()->getFrameCaptureShared()->onEndCLCapture())
{
delete cl::Platform::GetDefault()->getFrameCaptureShared();
}
}
bool FrameCaptureShared::onEndCLCapture()
{
if (mFrameIndex >= mCaptureStartFrame && mFrameIndex <= mCaptureEndFrame)
{
mActiveFrameIndices.push_back(mFrameIndex);
mCaptureEndFrame = mFrameIndex;
writeMainContextCppReplayCL();
writeCppReplayIndexFilesCL();
SaveBinaryData(mCompression, mOutDirectory, kNoContextId, mCaptureLabel, mBinaryData);
return true;
}
return false;
}
ResourceTrackerCL::ResourceTrackerCL() = default;
ResourceTrackerCL::~ResourceTrackerCL() = default;
void FrameCaptureShared::setCLPlatformIndices(cl_platform_id *platforms, size_t numPlatforms)
{
for (uint32_t i = 0; i < numPlatforms; ++i)
{
setIndex(&platforms[i]);
}
}
void FrameCaptureShared::setCLDeviceIndices(cl_device_id *devices, size_t numDevices)
{
for (uint32_t i = 0; i < numDevices; ++i)
{
setIndex(&devices[i]);
}
}
size_t FrameCaptureShared::getCLVoidIndex(const void *v)
{
if (mResourceTrackerCL.mCLVoidIndices.find(v) == mResourceTrackerCL.mCLVoidIndices.end())
{
return SIZE_MAX;
}
return mResourceTrackerCL.mCLVoidIndices[v];
}
void FrameCaptureShared::setCLVoidIndex(const void *v)
{
if (mResourceTrackerCL.mCLVoidIndices.find(v) == mResourceTrackerCL.mCLVoidIndices.end())
{
size_t tempSize = mResourceTrackerCL.mCLVoidIndices.size();
mResourceTrackerCL.mCLVoidIndices[v] = tempSize;
}
}
void FrameCaptureShared::setCLVoidVectorIndex(const void *pointers[],
size_t numPointers,
const angle::ParamCapture *paramCaptureKey)
{
mResourceTrackerCL.mCLParamIDToIndexVector[paramCaptureKey->uniqueID] = std::vector<size_t>();
for (size_t i = 0; i < numPointers; ++i)
{
mResourceTrackerCL.mCLParamIDToIndexVector[paramCaptureKey->uniqueID].push_back(
getCLVoidIndex(pointers[i]));
}
}
void FrameCaptureShared::setOffsetsVector(const void *args,
const void **argsLocations,
size_t numLocations,
const angle::ParamCapture *paramCaptureKey)
{
mResourceTrackerCL.mCLParamIDToIndexVector[paramCaptureKey->uniqueID] = std::vector<size_t>();
for (size_t i = 0; i < numLocations; ++i)
{
mResourceTrackerCL.mCLParamIDToIndexVector[paramCaptureKey->uniqueID].push_back(
(char *)argsLocations[i] - (char *)args);
}
}
std::vector<size_t> FrameCaptureShared::getCLObjVector(const angle::ParamCapture *paramCaptureKey)
{
if (mResourceTrackerCL.mCLParamIDToIndexVector.find(paramCaptureKey->uniqueID) !=
mResourceTrackerCL.mCLParamIDToIndexVector.end())
{
return mResourceTrackerCL.mCLParamIDToIndexVector[paramCaptureKey->uniqueID];
}
return std::vector<size_t>();
}
template <typename T>
std::unordered_map<T, size_t> &FrameCaptureShared::getMap()
{
ASSERT(false);
return std::unordered_map<T, size_t>();
}
template <>
std::unordered_map<cl_platform_id, size_t> &FrameCaptureShared::getMap<cl_platform_id>()
{
return mResourceTrackerCL.mCLPlatformIDIndices;
}
template <>
std::unordered_map<cl_device_id, size_t> &FrameCaptureShared::getMap<cl_device_id>()
{
return mResourceTrackerCL.mCLDeviceIDIndices;
}
template <>
std::unordered_map<cl_context, size_t> &FrameCaptureShared::getMap<cl_context>()
{
return mResourceTrackerCL.mCLContextIndices;
}
template <>
std::unordered_map<cl_event, size_t> &FrameCaptureShared::getMap<cl_event>()
{
return mResourceTrackerCL.mCLEventsIndices;
}
template <>
std::unordered_map<cl_command_queue, size_t> &FrameCaptureShared::getMap<cl_command_queue>()
{
return mResourceTrackerCL.mCLCommandQueueIndices;
}
template <>
std::unordered_map<cl_mem, size_t> &FrameCaptureShared::getMap<cl_mem>()
{
return mResourceTrackerCL.mCLMemIndices;
}
template <>
std::unordered_map<cl_sampler, size_t> &FrameCaptureShared::getMap<cl_sampler>()
{
return mResourceTrackerCL.mCLSamplerIndices;
}
template <>
std::unordered_map<cl_program, size_t> &FrameCaptureShared::getMap<cl_program>()
{
return mResourceTrackerCL.mCLProgramIndices;
}
template <>
std::unordered_map<cl_kernel, size_t> &FrameCaptureShared::getMap<cl_kernel>()
{
return mResourceTrackerCL.mCLKernelIndices;
}
void FrameCaptureShared::writeJSONCL()
{
JsonSerializer json;
json.startGroup("TraceMetadata");
json.addBool("IsBinaryDataCompressed", mCompression);
json.addScalar("CaptureRevision", GetANGLERevision());
json.addScalar("FrameStart", mCaptureStartFrame);
json.addScalar("FrameEnd", mFrameIndex);
json.addBool("IsOpenCL", true);
json.endGroup();
{
const std::vector<std::string> &traceFiles = mReplayWriter.getAndResetWrittenFiles();
json.addVectorOfStrings("TraceFiles", traceFiles);
}
{
std::stringstream jsonFileNameStream;
jsonFileNameStream << mOutDirectory << FmtCapturePrefix(kNoContextId, mCaptureLabel)
<< ".json";
std::string jsonFileName = jsonFileNameStream.str();
SaveFileHelper saveData(jsonFileName);
saveData.write(reinterpret_cast<const uint8_t *>(json.data()), json.length());
}
}
void FrameCaptureShared::saveCLGetInfo(const CallCapture &call)
{
std::string prevCall = "";
size_t size;
std::ostringstream objectStream;
std::string clObject;
JsonSerializer json;
json.startGroup(call.name());
// Below ONLY for clGetSupportedImageFormats
if (call.entryPoint == EntryPoint::CLGetSupportedImageFormats)
{
const cl_image_format *data =
call.params.getParam("image_formats", ParamType::Tcl_image_formatPointer, 4)
.value.cl_image_formatPointerVal;
if (!data)
{
return;
}
size_t *sizePointer =
call.params.getParam("num_image_formats", ParamType::Tcl_uintPointer, 5)
.value.size_tPointerVal;
if (!sizePointer)
{
size =
call.params.getParam("param_value_size", ParamType::Tcl_uint, 3).value.cl_uintVal;
}
else
{
size = *sizePointer;
}
cl_context context = call.params.getParamCaptures()[0].value.cl_contextVal;
objectStream << static_cast<void *>(context);
clObject = objectStream.str();
json.addString("context", clObject);
json.addScalar("flags", call.params.getParamCaptures()[1].value.MemFlagsVal.get());
cl::MemObjectType imageType =
call.params.getParam("image_typePacked", ParamType::TMemObjectType, 2)
.value.MemObjectTypeVal;
std::ostringstream oss;
oss << imageType;
std::string infoString = oss.str();
json.startGroup(infoString);
for (size_t j = 0; j < size; ++j)
{
std::ostringstream temp;
temp << (j + 1);
json.addScalar("image_channel_order" + temp.str(), data[j].image_channel_order);
json.addScalar("image_channel_data_type" + temp.str(), data[j].image_channel_data_type);
}
json.endGroup();
json.endGroup();
return;
}
// Get the param_value and size
bool offsetData = 0;
switch (call.entryPoint)
{
case EntryPoint::CLGetProgramBuildInfo:
case EntryPoint::CLGetKernelArgInfo:
case EntryPoint::CLGetKernelWorkGroupInfo:
{
offsetData = 1;
break;
}
default:
break;
}
const void *data = call.params.getParam("param_value", ParamType::TvoidPointer, 3 + offsetData)
.value.voidPointerVal;
if (!data)
{
return;
}
size_t *sizePointer =
call.params.getParam("param_value_size_ret", ParamType::Tsize_tPointer, 4 + offsetData)
.value.size_tPointerVal;
if (!sizePointer)
{
size = call.params.getParam("param_value_size", ParamType::Tsize_t, 2 + offsetData)
.value.size_tVal;
}
else
{
size = *sizePointer;
}
// Get string representation of OpenCL object specified
switch (call.entryPoint)
{
case EntryPoint::CLGetPlatformInfo:
{
cl_platform_id platform = call.params.getParamCaptures()[0].value.cl_platform_idVal;
objectStream << static_cast<void *>(platform);
break;
}
case EntryPoint::CLGetDeviceInfo:
{
cl_device_id device = call.params.getParamCaptures()[0].value.cl_device_idVal;
objectStream << static_cast<void *>(device);
break;
}
case EntryPoint::CLGetContextInfo:
{
cl_context context = call.params.getParamCaptures()[0].value.cl_contextVal;
objectStream << static_cast<void *>(context);
break;
}
case EntryPoint::CLGetCommandQueueInfo:
{
cl_command_queue commandQueue =
call.params.getParamCaptures()[0].value.cl_command_queueVal;
objectStream << static_cast<void *>(commandQueue);
break;
}
case EntryPoint::CLGetProgramInfo:
case EntryPoint::CLGetProgramBuildInfo:
{
cl_program program = call.params.getParamCaptures()[0].value.cl_programVal;
objectStream << static_cast<void *>(program);
break;
}
case EntryPoint::CLGetKernelInfo:
case EntryPoint::CLGetKernelArgInfo:
case EntryPoint::CLGetKernelWorkGroupInfo:
{
cl_kernel kernel = call.params.getParamCaptures()[0].value.cl_kernelVal;
objectStream << static_cast<void *>(kernel);
break;
}
case EntryPoint::CLGetEventInfo:
case EntryPoint::CLGetEventProfilingInfo:
{
cl_event event = call.params.getParamCaptures()[0].value.cl_eventVal;
objectStream << static_cast<void *>(event);
break;
}
case EntryPoint::CLGetMemObjectInfo:
case EntryPoint::CLGetImageInfo:
{
cl_mem mem = call.params.getParamCaptures()[0].value.cl_memVal;
objectStream << static_cast<void *>(mem);
break;
}
case EntryPoint::CLGetSamplerInfo:
{
cl_sampler sampler = call.params.getParamCaptures()[0].value.cl_samplerVal;
objectStream << static_cast<void *>(sampler);
break;
}
default:
break;
}
clObject = objectStream.str();
// Go through the param_name options
switch (call.entryPoint)
{
case EntryPoint::CLGetPlatformInfo:
{
cl::PlatformInfo info = call.params.getParamCaptures()[1].value.PlatformInfoVal;
std::ostringstream oss;
oss << info;
std::string infoString = oss.str();
json.addString("platform", clObject);
switch (ToCLenum(info))
{
case CL_PLATFORM_PROFILE:
case CL_PLATFORM_VERSION:
case CL_PLATFORM_NAME:
case CL_PLATFORM_VENDOR:
case CL_PLATFORM_EXTENSIONS:
case CL_PLATFORM_ICD_SUFFIX_KHR:
{
json.addCString(infoString, static_cast<const char *>(data));
break;
}
case CL_PLATFORM_EXTENSIONS_WITH_VERSION:
{
const cl_name_version *nameVersion = static_cast<const cl_name_version *>(data);
json.startGroup(infoString);
for (size_t j = 0; j < size / sizeof(cl_name_version); ++j)
{
json.addScalar(nameVersion[j].name, nameVersion[j].version);
}
json.endGroup();
break;
}
case CL_PLATFORM_NUMERIC_VERSION:
json.addScalar(infoString, *static_cast<const cl_uint *>(data));
break;
case CL_PLATFORM_HOST_TIMER_RESOLUTION:
case CL_PLATFORM_COMMAND_BUFFER_CAPABILITIES_KHR:
json.addScalar(infoString, *static_cast<const cl_ulong *>(data));
break;
case CL_PLATFORM_EXTERNAL_MEMORY_IMPORT_HANDLE_TYPES_KHR:
case CL_PLATFORM_SEMAPHORE_TYPES_KHR:
case CL_PLATFORM_SEMAPHORE_IMPORT_HANDLE_TYPES_KHR:
case CL_PLATFORM_SEMAPHORE_EXPORT_HANDLE_TYPES_KHR:
json.addVector(infoString,
std::vector<cl_uint>((cl_uint *)data,
(cl_uint *)data + size / sizeof(cl_uint)));
break;
default:
// Not supported or cannot add to JSON file
break;
}
break;
}
case EntryPoint::CLGetDeviceInfo:
{
cl::DeviceInfo info = call.params.getParamCaptures()[1].value.DeviceInfoVal;
std::ostringstream oss;
oss << info;
std::string infoString = oss.str();
json.addString("device", clObject);
switch (ToCLenum(info))
{
case CL_DEVICE_IL_VERSION:
case CL_DEVICE_LATEST_CONFORMANCE_VERSION_PASSED:
case CL_DEVICE_OPENCL_C_VERSION:
case CL_DEVICE_EXTENSIONS:
case CL_DEVICE_VERSION:
case CL_DEVICE_PROFILE:
case CL_DRIVER_VERSION:
case CL_DEVICE_VENDOR:
case CL_DEVICE_NAME:
json.addCString(infoString, static_cast<const char *>(data));
break;
case CL_DEVICE_TYPE:
case CL_DEVICE_MAX_MEM_ALLOC_SIZE:
case CL_DEVICE_LOCAL_MEM_SIZE:
case CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE:
case CL_DEVICE_GLOBAL_MEM_SIZE:
case CL_DEVICE_GLOBAL_MEM_CACHE_SIZE:
case CL_DEVICE_HALF_FP_CONFIG:
case CL_DEVICE_SINGLE_FP_CONFIG:
case CL_DEVICE_DEVICE_ENQUEUE_CAPABILITIES:
case CL_DEVICE_ATOMIC_FENCE_CAPABILITIES:
case CL_DEVICE_ATOMIC_MEMORY_CAPABILITIES:
case CL_DEVICE_SVM_CAPABILITIES:
case CL_DEVICE_QUEUE_ON_DEVICE_PROPERTIES:
case CL_DEVICE_PARTITION_AFFINITY_DOMAIN:
case CL_DEVICE_DOUBLE_FP_CONFIG:
case CL_DEVICE_QUEUE_ON_HOST_PROPERTIES:
case CL_DEVICE_EXECUTION_CAPABILITIES:
// cl_ulong and cl_bitfield
json.addScalar(infoString, *static_cast<const cl_ulong *>(data));
break;
case CL_DEVICE_VENDOR_ID:
case CL_DEVICE_MAX_COMPUTE_UNITS:
case CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS:
case CL_DEVICE_PREFERRED_VECTOR_WIDTH_CHAR:
case CL_DEVICE_PREFERRED_VECTOR_WIDTH_SHORT:
case CL_DEVICE_PREFERRED_VECTOR_WIDTH_INT:
case CL_DEVICE_PREFERRED_VECTOR_WIDTH_LONG:
case CL_DEVICE_PREFERRED_VECTOR_WIDTH_FLOAT:
case CL_DEVICE_PREFERRED_VECTOR_WIDTH_DOUBLE:
case CL_DEVICE_PREFERRED_VECTOR_WIDTH_HALF:
case CL_DEVICE_NATIVE_VECTOR_WIDTH_CHAR:
case CL_DEVICE_NATIVE_VECTOR_WIDTH_SHORT:
case CL_DEVICE_NATIVE_VECTOR_WIDTH_INT:
case CL_DEVICE_NATIVE_VECTOR_WIDTH_LONG:
case CL_DEVICE_NATIVE_VECTOR_WIDTH_FLOAT:
case CL_DEVICE_NATIVE_VECTOR_WIDTH_DOUBLE:
case CL_DEVICE_NATIVE_VECTOR_WIDTH_HALF:
case CL_DEVICE_MAX_CLOCK_FREQUENCY:
case CL_DEVICE_ADDRESS_BITS:
case CL_DEVICE_IMAGE_SUPPORT:
case CL_DEVICE_MAX_READ_IMAGE_ARGS:
case CL_DEVICE_MAX_WRITE_IMAGE_ARGS:
case CL_DEVICE_MAX_READ_WRITE_IMAGE_ARGS:
case CL_DEVICE_PIPE_SUPPORT:
case CL_DEVICE_GENERIC_ADDRESS_SPACE_SUPPORT:
case CL_DEVICE_WORK_GROUP_COLLECTIVE_FUNCTIONS_SUPPORT:
case CL_DEVICE_NON_UNIFORM_WORK_GROUP_SUPPORT:
case CL_DEVICE_NUMERIC_VERSION:
case CL_DEVICE_SUB_GROUP_INDEPENDENT_FORWARD_PROGRESS:
case CL_DEVICE_MAX_NUM_SUB_GROUPS:
case CL_DEVICE_PREFERRED_LOCAL_ATOMIC_ALIGNMENT:
case CL_DEVICE_PREFERRED_GLOBAL_ATOMIC_ALIGNMENT:
case CL_DEVICE_PREFERRED_PLATFORM_ATOMIC_ALIGNMENT:
case CL_DEVICE_PIPE_MAX_PACKET_SIZE:
case CL_DEVICE_PIPE_MAX_ACTIVE_RESERVATIONS:
case CL_DEVICE_MAX_PIPE_ARGS:
case CL_DEVICE_MAX_ON_DEVICE_EVENTS:
case CL_DEVICE_MAX_ON_DEVICE_QUEUES:
case CL_DEVICE_QUEUE_ON_DEVICE_MAX_SIZE:
case CL_DEVICE_QUEUE_ON_DEVICE_PREFERRED_SIZE:
case CL_DEVICE_IMAGE_BASE_ADDRESS_ALIGNMENT:
case CL_DEVICE_IMAGE_PITCH_ALIGNMENT:
case CL_DEVICE_PREFERRED_INTEROP_USER_SYNC:
case CL_DEVICE_REFERENCE_COUNT:
case CL_DEVICE_PARTITION_MAX_SUB_DEVICES:
case CL_DEVICE_LINKER_AVAILABLE:
case CL_DEVICE_HOST_UNIFIED_MEMORY:
case CL_DEVICE_COMPILER_AVAILABLE:
case CL_DEVICE_AVAILABLE:
case CL_DEVICE_ENDIAN_LITTLE:
case CL_DEVICE_ERROR_CORRECTION_SUPPORT:
case CL_DEVICE_LOCAL_MEM_TYPE:
case CL_DEVICE_MAX_CONSTANT_ARGS:
case CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE:
case CL_DEVICE_MIN_DATA_TYPE_ALIGN_SIZE:
case CL_DEVICE_MEM_BASE_ADDR_ALIGN:
case CL_DEVICE_MAX_SAMPLERS:
json.addScalar(infoString, *static_cast<const cl_uint *>(data));
break;
case CL_DEVICE_MAX_WORK_GROUP_SIZE:
case CL_DEVICE_IMAGE2D_MAX_WIDTH:
case CL_DEVICE_IMAGE2D_MAX_HEIGHT:
case CL_DEVICE_IMAGE3D_MAX_WIDTH:
case CL_DEVICE_IMAGE3D_MAX_HEIGHT:
case CL_DEVICE_IMAGE3D_MAX_DEPTH:
case CL_DEVICE_IMAGE_MAX_BUFFER_SIZE:
case CL_DEVICE_IMAGE_MAX_ARRAY_SIZE:
case CL_DEVICE_PREFERRED_WORK_GROUP_SIZE_MULTIPLE:
case CL_DEVICE_GLOBAL_VARIABLE_PREFERRED_TOTAL_SIZE:
case CL_DEVICE_MAX_GLOBAL_VARIABLE_SIZE:
case CL_DEVICE_PRINTF_BUFFER_SIZE:
case CL_DEVICE_PROFILING_TIMER_RESOLUTION:
case CL_DEVICE_MAX_PARAMETER_SIZE:
json.addScalar(infoString, *static_cast<const size_t *>(data));
break;
case CL_DEVICE_MAX_WORK_ITEM_SIZES:
json.addVector(infoString,
std::vector<size_t>((size_t *)data,
(size_t *)data + size / sizeof(size_t)));
break;
case CL_DEVICE_PARTITION_TYPE:
case CL_DEVICE_PARTITION_PROPERTIES:
json.addVector(infoString, std::vector<cl_ulong>(
(cl_ulong *)data,
(cl_ulong *)data + size / sizeof(cl_ulong)));
break;
case CL_DEVICE_PARENT_DEVICE:
{
std::ostringstream voidStream;
voidStream << static_cast<void *>(*(cl_device_id *)data);
std::string memLoc = voidStream.str();
json.addCString(infoString, memLoc.c_str());
break;
}
case CL_DEVICE_PLATFORM:
{
std::ostringstream voidStream;
voidStream << static_cast<void *>(*(cl_platform_id *)data);
std::string memLoc = voidStream.str();
json.addCString(infoString, memLoc.c_str());
break;
}
case CL_DEVICE_ILS_WITH_VERSION:
case CL_DEVICE_OPENCL_C_FEATURES:
case CL_DEVICE_OPENCL_C_ALL_VERSIONS:
case CL_DEVICE_BUILT_IN_KERNELS_WITH_VERSION:
case CL_DEVICE_EXTENSIONS_WITH_VERSION:
{
const cl_name_version *nameVersion = static_cast<const cl_name_version *>(data);
json.startGroup(infoString);
for (size_t j = 0; j < size / sizeof(cl_name_version); ++j)
{
json.addScalar(nameVersion[j].name, nameVersion[j].version);
}
json.endGroup();
break;
}
default:
// Not supported or cannot add to JSON file
break;
}
break;
}
case EntryPoint::CLGetContextInfo:
{
cl::ContextInfo info = call.params.getParamCaptures()[1].value.ContextInfoVal;
std::ostringstream oss;
oss << info;
std::string infoString = oss.str();
json.addString("context", clObject);
switch (ToCLenum(info))
{
case CL_PLATFORM_PROFILE:
case CL_PLATFORM_VERSION:
case CL_PLATFORM_NAME:
case CL_PLATFORM_VENDOR:
case CL_PLATFORM_EXTENSIONS:
case CL_PLATFORM_ICD_SUFFIX_KHR:
{
json.addCString(infoString, static_cast<const char *>(data));
break;
}
case CL_CONTEXT_REFERENCE_COUNT:
case CL_CONTEXT_NUM_DEVICES:
json.addScalar(infoString, *static_cast<const cl_uint *>(data));
break;
case CL_CONTEXT_PROPERTIES:
json.addVector(infoString, std::vector<cl_ulong>(
(cl_ulong *)data,
(cl_ulong *)data + size / sizeof(cl_ulong)));
break;
case CL_CONTEXT_DEVICES:
{
const cl_device_id *devices = static_cast<const cl_device_id *>(data);
std::vector<std::string> devicesStrings;
for (size_t j = 0; j < size / sizeof(cl_device_id); ++j)
{
std::ostringstream voidStream;
voidStream << static_cast<void *>(devices[j]);
devicesStrings.push_back(voidStream.str());
}
json.addVectorOfStrings(infoString, devicesStrings);
break;
}
default:
// Not supported or cannot add to JSON file
break;
}
break;
}
case EntryPoint::CLGetCommandQueueInfo:
{
cl::CommandQueueInfo info = call.params.getParamCaptures()[1].value.CommandQueueInfoVal;
std::ostringstream oss;
oss << info;
std::string infoString = oss.str();
json.addString("command_queue", clObject);
switch (ToCLenum(info))
{
case CL_QUEUE_DEVICE_DEFAULT:
{
std::ostringstream voidStream;
voidStream << static_cast<void *>(*(cl_command_queue *)data);
std::string memLoc = voidStream.str();
json.addCString(infoString, memLoc.c_str());
break;
}
case CL_QUEUE_CONTEXT:
{
std::ostringstream voidStream;
voidStream << static_cast<void *>(*(cl_context *)data);
std::string memLoc = voidStream.str();
json.addCString(infoString, memLoc.c_str());
break;
}
case CL_QUEUE_DEVICE:
{
std::ostringstream voidStream;
voidStream << static_cast<void *>(*(cl_device_id *)data);
std::string memLoc = voidStream.str();
json.addCString(infoString, memLoc.c_str());
break;
}
case CL_QUEUE_REFERENCE_COUNT:
case CL_QUEUE_SIZE:
{
json.addScalar(infoString, *static_cast<const cl_uint *>(data));
break;
}
case CL_QUEUE_PROPERTIES:
{
json.addScalar(infoString, *static_cast<const cl_ulong *>(data));
break;
}
case CL_QUEUE_PROPERTIES_ARRAY:
{
json.addVector(infoString, std::vector<cl_ulong>(
(cl_ulong *)data,
(cl_ulong *)data + size / sizeof(cl_ulong)));
break;
}
default:
// Not supported or cannot add to JSON file
break;
}
break;
}
case EntryPoint::CLGetProgramInfo:
{
cl::ProgramInfo info = call.params.getParamCaptures()[1].value.ProgramInfoVal;
std::ostringstream oss;
oss << info;
std::string infoString = oss.str();
json.addString("program", clObject);
switch (ToCLenum(info))
{
case CL_PROGRAM_SOURCE:
case CL_PROGRAM_IL:
case CL_PROGRAM_KERNEL_NAMES:
{
json.addCString(infoString, static_cast<const char *>(data));
break;
}
case CL_PROGRAM_CONTEXT:
{
std::ostringstream voidStream;
voidStream << static_cast<void *>(*(cl_context *)data);
std::string memLoc = voidStream.str();
json.addCString(infoString, memLoc.c_str());
break;
}
case CL_PROGRAM_REFERENCE_COUNT:
case CL_PROGRAM_NUM_DEVICES:
case CL_PROGRAM_SCOPE_GLOBAL_CTORS_PRESENT:
case CL_PROGRAM_SCOPE_GLOBAL_DTORS_PRESENT:
{
json.addScalar(infoString, *static_cast<const cl_uint *>(data));
break;
}
case CL_PROGRAM_DEVICES:
{
const cl_device_id *devices = static_cast<const cl_device_id *>(data);
std::vector<std::string> devicesStrings;
for (size_t j = 0; j < size / sizeof(cl_device_id); ++j)
{
std::ostringstream voidStream;
voidStream << static_cast<void *>(devices[j]);
devicesStrings.push_back(voidStream.str());
}
json.addVectorOfStrings(infoString, devicesStrings);
break;
}
case CL_PROGRAM_NUM_KERNELS:
{
json.addScalar(infoString, *static_cast<const size_t *>(data));
break;
}
case CL_PROGRAM_BINARY_SIZES:
json.addVector(infoString,
std::vector<size_t>((size_t *)data,
(size_t *)data + size / sizeof(size_t)));
break;
case CL_PROGRAM_BINARIES:
json.addVector(infoString,
std::vector<unsigned char>(
(unsigned char *)data,
(unsigned char *)data + size / sizeof(unsigned char)));
break;
default:
// Not supported or cannot add to JSON file
break;
}
break;
}
case EntryPoint::CLGetProgramBuildInfo:
{
cl::ProgramBuildInfo info = call.params.getParamCaptures()[2].value.ProgramBuildInfoVal;
std::ostringstream oss;
oss << info;
std::string infoString = oss.str();
json.addString("program", clObject);
cl_device_id device = call.params.getParamCaptures()[1].value.cl_device_idVal;
std::ostringstream objectStream2;
objectStream2 << static_cast<void *>(device);
clObject = objectStream2.str();
json.addString("device", clObject);
switch (ToCLenum(info))
{
case CL_PROGRAM_BUILD_OPTIONS:
case CL_PROGRAM_BUILD_LOG:
{
json.addCString(infoString, static_cast<const char *>(data));
break;
}
case CL_PROGRAM_BINARY_TYPE:
{
json.addScalar(infoString, *static_cast<const cl_uint *>(data));
break;
}
case CL_PROGRAM_BUILD_STATUS:
{
json.addScalar(infoString, *static_cast<const cl_int *>(data));
break;
}
case CL_PROGRAM_BUILD_GLOBAL_VARIABLE_TOTAL_SIZE:
{
json.addScalar(infoString, *static_cast<const size_t *>(data));
break;
}
default:
// Not supported or cannot add to JSON file
break;
}
break;
}
case EntryPoint::CLGetKernelInfo:
{
cl::KernelInfo info = call.params.getParamCaptures()[1].value.KernelInfoVal;
std::ostringstream oss;
oss << info;
std::string infoString = oss.str();
json.addString("kernel", clObject);
switch (ToCLenum(info))
{
case CL_KERNEL_FUNCTION_NAME:
case CL_KERNEL_ATTRIBUTES:
{
json.addCString(infoString, static_cast<const char *>(data));
break;
}
case CL_KERNEL_NUM_ARGS:
case CL_KERNEL_REFERENCE_COUNT:
{
json.addScalar(infoString, *static_cast<const cl_uint *>(data));
break;
}
case CL_KERNEL_CONTEXT:
{
std::ostringstream voidStream;
voidStream << static_cast<void *>(*(cl_context *)data);
std::string memLoc = voidStream.str();
json.addCString(infoString, memLoc.c_str());
break;
}
case CL_KERNEL_PROGRAM:
{
std::ostringstream voidStream;
voidStream << static_cast<void *>(*(cl_program *)data);
std::string memLoc = voidStream.str();
json.addCString(infoString, memLoc.c_str());
break;
}
default:
// Not supported or cannot add to JSON file
break;
}
break;
}
case EntryPoint::CLGetKernelArgInfo:
{
cl::KernelArgInfo info = call.params.getParamCaptures()[2].value.KernelArgInfoVal;
std::ostringstream oss;
oss << info;
std::string infoString = oss.str();
json.addString("kernel", clObject);
cl_uint index = call.params.getParamCaptures()[1].value.cl_uintVal;
json.addScalar("arg_index", index);
switch (ToCLenum(info))
{
case CL_KERNEL_ARG_TYPE_NAME:
case CL_KERNEL_ARG_NAME:
{
json.addCString(infoString, static_cast<const char *>(data));
break;
}
case CL_KERNEL_ARG_ADDRESS_QUALIFIER:
case CL_KERNEL_ARG_ACCESS_QUALIFIER:
{
json.addScalar(infoString, *static_cast<const cl_uint *>(data));
break;
}
case CL_KERNEL_ARG_TYPE_QUALIFIER:
{
json.addScalar(infoString, *static_cast<const cl_ulong *>(data));
break;
}
default:
// Not supported or cannot add to JSON file
break;
}
break;
}
case EntryPoint::CLGetKernelWorkGroupInfo:
{
cl::KernelWorkGroupInfo info =
call.params.getParamCaptures()[2].value.KernelWorkGroupInfoVal;
std::ostringstream oss;
oss << info;
std::string infoString = oss.str();
json.addString("kernel", clObject);
cl_device_id device = call.params.getParamCaptures()[1].value.cl_device_idVal;
std::ostringstream objectStream2;
objectStream2 << static_cast<void *>(device);
clObject = objectStream2.str();
json.addString("device", clObject);
switch (ToCLenum(info))
{
case CL_KERNEL_LOCAL_MEM_SIZE:
case CL_KERNEL_PRIVATE_MEM_SIZE:
{
json.addScalar(infoString, *static_cast<const cl_ulong *>(data));
break;
}
case CL_KERNEL_WORK_GROUP_SIZE:
case CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE:
{
json.addScalar(infoString, *static_cast<const size_t *>(data));
break;
}
case CL_KERNEL_GLOBAL_WORK_SIZE:
case CL_KERNEL_COMPILE_WORK_GROUP_SIZE:
{
json.addVector(infoString,
std::vector<size_t>((size_t *)data, (size_t *)data + 3));
break;
}
default:
// Not supported or cannot add to JSON file
break;
}
break;
}
case EntryPoint::CLGetEventInfo:
{
cl::EventInfo info = call.params.getParamCaptures()[1].value.EventInfoVal;
std::ostringstream oss;
oss << info;
std::string infoString = oss.str();
json.addString("event", clObject);
switch (ToCLenum(info))
{
case CL_EVENT_REFERENCE_COUNT:
case CL_EVENT_COMMAND_TYPE:
{
json.addScalar(infoString, *static_cast<const cl_uint *>(data));
break;
}
case CL_EVENT_COMMAND_EXECUTION_STATUS:
{
json.addScalar(infoString, *static_cast<const cl_int *>(data));
break;
}
case CL_EVENT_CONTEXT:
{
std::ostringstream voidStream;
voidStream << static_cast<void *>(*(cl_context *)data);
std::string memLoc = voidStream.str();
json.addCString(infoString, memLoc.c_str());
break;
}
case CL_EVENT_COMMAND_QUEUE:
{
std::ostringstream voidStream;
voidStream << static_cast<void *>(*(cl_command_queue *)data);
std::string memLoc = voidStream.str();
json.addCString(infoString, memLoc.c_str());
break;
}
default:
// Not supported or cannot add to JSON file
break;
}
break;
}
case EntryPoint::CLGetEventProfilingInfo:
{
cl::ProfilingInfo info = call.params.getParamCaptures()[1].value.ProfilingInfoVal;
std::ostringstream oss;
oss << info;
std::string infoString = oss.str();
json.addString("event", clObject);
switch (ToCLenum(info))
{
case CL_PROFILING_COMMAND_QUEUED:
case CL_PROFILING_COMMAND_SUBMIT:
case CL_PROFILING_COMMAND_START:
case CL_PROFILING_COMMAND_END:
case CL_PROFILING_COMMAND_COMPLETE:
{
json.addScalar(infoString, *static_cast<const cl_ulong *>(data));
break;
}
default:
// Not supported or cannot add to JSON file
break;
}
break;
}
case EntryPoint::CLGetMemObjectInfo:
{
cl::MemInfo info = call.params.getParamCaptures()[1].value.MemInfoVal;
std::ostringstream oss;
oss << info;
std::string infoString = oss.str();
json.addString("memObj", clObject);
switch (ToCLenum(info))
{
case CL_MEM_TYPE:
case CL_MEM_MAP_COUNT:
case CL_MEM_REFERENCE_COUNT:
case CL_MEM_USES_SVM_POINTER:
{
json.addScalar(infoString, *static_cast<const cl_uint *>(data));
break;
}
case CL_MEM_FLAGS:
{
json.addScalar(infoString, *static_cast<const cl_ulong *>(data));
break;
}
case CL_MEM_SIZE:
case CL_MEM_OFFSET:
{
json.addScalar(infoString, *static_cast<const size_t *>(data));
break;
}
case CL_MEM_HOST_PTR:
{
std::ostringstream voidStream;
voidStream << data;
std::string memLoc = voidStream.str();
json.addCString(infoString, memLoc.c_str());
break;
}
case CL_MEM_CONTEXT:
{
std::ostringstream voidStream;
voidStream << static_cast<void *>(*(cl_context *)data);
std::string memLoc = voidStream.str();
json.addCString(infoString, memLoc.c_str());
break;
}
case CL_MEM_ASSOCIATED_MEMOBJECT:
{
std::ostringstream voidStream;
voidStream << static_cast<void *>(*(cl_mem *)data);
std::string memLoc = voidStream.str();
json.addCString(infoString, memLoc.c_str());
break;
}
case CL_MEM_PROPERTIES:
{
json.addVector(infoString, std::vector<cl_mem_properties>(
(cl_mem_properties *)data,
(cl_mem_properties *)data +
size / sizeof(cl_mem_properties)));
break;
}
default:
// Not supported or cannot add to JSON file
break;
}
break;
}
case EntryPoint::CLGetImageInfo:
{
cl::ImageInfo info = call.params.getParamCaptures()[1].value.ImageInfoVal;
std::ostringstream oss;
oss << info;
std::string infoString = oss.str();
json.addString("image", clObject);
switch (ToCLenum(info))
{
case CL_IMAGE_FORMAT:
{
json.startGroup(infoString);
json.addScalar("image_channel_order",
static_cast<const cl_image_format *>(data)->image_channel_order);
json.addScalar(
"image_channel_data_type",
static_cast<const cl_image_format *>(data)->image_channel_data_type);
json.endGroup();
break;
}
case CL_IMAGE_NUM_MIP_LEVELS:
case CL_IMAGE_NUM_SAMPLES:
{
json.addScalar(infoString, *static_cast<const cl_uint *>(data));
break;
}
case CL_IMAGE_ELEMENT_SIZE:
case CL_IMAGE_ROW_PITCH:
case CL_IMAGE_SLICE_PITCH:
case CL_IMAGE_WIDTH:
case CL_IMAGE_HEIGHT:
case CL_IMAGE_DEPTH:
case CL_IMAGE_ARRAY_SIZE:
{
json.addScalar(infoString, *static_cast<const size_t *>(data));
break;
}
case CL_IMAGE_BUFFER:
{
std::ostringstream voidStream;
voidStream << static_cast<void *>(*(cl_mem *)data);
std::string memLoc = voidStream.str();
json.addCString(infoString, memLoc.c_str());
break;
}
default:
// Not supported or cannot add to JSON file
break;
}
break;
}
case EntryPoint::CLGetSamplerInfo:
{
cl::SamplerInfo info = call.params.getParamCaptures()[1].value.SamplerInfoVal;
std::ostringstream oss;
oss << info;
std::string infoString = oss.str();
json.addString("image", clObject);
switch (ToCLenum(info))
{
case CL_SAMPLER_REFERENCE_COUNT:
case CL_SAMPLER_NORMALIZED_COORDS:
case CL_SAMPLER_ADDRESSING_MODE:
case CL_SAMPLER_FILTER_MODE:
{
json.addScalar(infoString, *static_cast<const cl_uint *>(data));
break;
}
case CL_SAMPLER_PROPERTIES:
{
json.addVector(infoString, std::vector<cl_sampler_properties>(
(cl_sampler_properties *)data,
(cl_sampler_properties *)data +
size / sizeof(cl_sampler_properties)));
break;
}
case CL_SAMPLER_CONTEXT:
{
std::ostringstream voidStream;
voidStream << static_cast<void *>(*(cl_context *)data);
std::string memLoc = voidStream.str();
json.addCString(infoString, memLoc.c_str());
break;
}
default:
// Not supported or cannot add to JSON file
break;
}
break;
}
default:
break;
}
json.endGroup();
mCLInfoJson += std::string(json.data()) + ",\n";
}
void FrameCaptureShared::writeJSONCLGetInfo()
{
std::stringstream jsonFileNameStream;
jsonFileNameStream << mOutDirectory << FmtCapturePrefix(kNoContextId, mCaptureLabel)
<< "_OpenCL_info.json";
std::string jsonFileName = jsonFileNameStream.str();
SaveFileHelper saveData(jsonFileName);
saveData.write(reinterpret_cast<const uint8_t *>(mCLInfoJson.c_str()), mCLInfoJson.length());
}
void FrameCaptureShared::writeCppReplayIndexFilesCL()
{
// Ensure the last frame is written. This will no-op if the frame is already written.
mReplayWriter.saveFrame();
{
std::stringstream header;
header << "#pragma once\n";
header << "\n";
header << "#define CL_NO_EXTENSION_PROTOTYPES\n";
header << "#include <angle_cl.h>\n";
header << "#include <stdint.h>\n";
header << "#include \"trace_fixture_cl.h\"\n";
std::string includes = header.str();
mReplayWriter.setHeaderPrologue(includes);
}
{
std::stringstream source;
source << "#include \"" << FmtCapturePrefix(kNoContextId, mCaptureLabel) << ".h\"\n";
source << "#include \"trace_fixture_cl.h\"\n";
std::string sourcePrologue = source.str();
mReplayWriter.setSourcePrologue(sourcePrologue);
}
{
std::string proto = "void InitReplay(void)";
std::stringstream source;
source << proto << "\n";
source << "{\n";
WriteInitReplayCallCL(mCompression, source, mCaptureLabel, 0, mReadBufferSize,
mMaxCLParamsSize);
source << "}\n";
mReplayWriter.addPrivateFunction(proto, std::stringstream(), source);
}
{
std::string proto = "void ReplayFrame(uint32_t frameIndex)";
std::stringstream source;
source << proto << "\n";
source << "{\n";
source << " switch (frameIndex)\n";
source << " {\n";
for (uint32_t frameIndex : mActiveFrameIndices)
{
source << " case " << frameIndex << ":\n";
source << " " << FmtReplayFunction(kNoContextId, FuncUsage::Call, frameIndex)
<< ";\n";
source << " break;\n";
}
source << " default:\n";
source << " break;\n";
source << " }\n";
source << "}\n";
mReplayWriter.addPublicFunction(proto, std::stringstream(), source);
}
for (auto extFuncName : mExtFuncsAdded)
{
mReplayWriter.addStaticVariable(extFuncName + "_fn", extFuncName);
}
std::stringstream protoSetupStream;
protoSetupStream << "void SetupFirstFrame(void)";
std::string protoSetup = protoSetupStream.str();
std::stringstream headerStreamSetup;
std::stringstream bodyStreamSetup;
WriteCppReplayFunctionWithPartsCL(ReplayFunc::SetupFirstFrame, mReplayWriter,
mCaptureStartFrame, &mBinaryData, mCLSetupCalls,
headerStreamSetup, bodyStreamSetup);
mReplayWriter.addPublicFunction(protoSetup, headerStreamSetup, bodyStreamSetup);
{
std::string proto = "void ResetReplay(void)";
std::stringstream source;
source << proto << "\n" << "{\n";
printCLResetObjs(source);
source << "}\n";
mReplayWriter.addPublicFunction(proto, std::stringstream(), source);
}
{
std::stringstream fnameStream;
fnameStream << mOutDirectory << FmtCapturePrefix(kNoContextId, mCaptureLabel);
std::string fnamePattern = fnameStream.str();
mReplayWriter.setFilenamePattern(fnamePattern);
}
mReplayWriter.saveIndexFilesAndHeader();
writeJSONCL();
writeJSONCLGetInfo();
}
void FrameCaptureShared::writeMainContextCppReplayCL()
{
{
std::stringstream header;
header << "#include \"" << FmtCapturePrefix(kNoContextId, mCaptureLabel) << ".h\"\n";
header << "#include \"trace_fixture_cl.h\"\n";
std::string headerString = header.str();
mReplayWriter.setSourcePrologue(headerString);
}
uint32_t frameIndex = getReplayFrameIndex();
if (frameIndex == 1)
{
{
std::string proto = "void SetupReplay(void)";
std::stringstream out;
out << proto << "\n";
out << "{\n";
// Setup all of the shared objects.
out << " InitReplay();\n";
out << "}\n";
mReplayWriter.addPublicFunction(proto, std::stringstream(), out);
}
}
if (!mFrameCalls.empty())
{
std::stringstream protoStream;
protoStream << "void "
<< FmtReplayFunction(kNoContextId, FuncUsage::Prototype, mFrameIndex);
std::string proto = protoStream.str();
std::stringstream headerStream;
std::stringstream bodyStream;
WriteCppReplayFunctionWithPartsCL(ReplayFunc::Replay, mReplayWriter, mFrameIndex,
&mBinaryData, mFrameCalls, headerStream, bodyStream);
mReplayWriter.addPrivateFunction(proto, headerStream, bodyStream);
}
{
std::stringstream fnamePatternStream;
fnamePatternStream << mOutDirectory << FmtCapturePrefix(kNoContextId, mCaptureLabel);
std::string fnamePattern = fnamePatternStream.str();
mReplayWriter.setFilenamePattern(fnamePattern);
}
if (mFrameIndex == mCaptureEndFrame)
{
mReplayWriter.saveFrame();
}
}
} // namespace angle