Skip to content

Commit 3fa8fc4

Browse files
marandjemarandje
authored
fix(clr): Optimize D2D 3D batch copies (#12206)
## Motivation - Optimize D2D batch copies to make use of batch HSA API ## Technical Details - Enqueue D2D as batch copies instead of enqueuing them sequentially one by one. ## Issue Tracking ISSUE ID ROCM-31101 ## Test Plan - Add new 3D batch tests to cover optimized path. - Run performance tests. ## Test Result - All newly added tests are passing. - Batch copies are now showing up to 15x improvements over sequential copies. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/rocm-systems/blob/develop/CONTRIBUTING.md.
1 parent 08c6e8f commit 3fa8fc4

28 files changed

Lines changed: 1693 additions & 472 deletions

‎projects/clr/hipamd/src/hip_graph.cpp‎

Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -185,6 +185,11 @@ hipError_t ihipGraphAddMemcpyNode(hip::GraphNode** pGraphNode, hip::Graph* graph
185185
if (status != hipSuccess) {
186186
return status;
187187
}
188+
const HIP_MEMCPY3D desc = hip::getDrvMemcpy3DDesc(*pCopyParams);
189+
status = ihipDrvMemcpy3D_validate(&desc);
190+
if (status != hipSuccess) {
191+
return status;
192+
}
188193
*pGraphNode = new hip::GraphMemcpyNode(pCopyParams);
189194
status = ihipGraphAddNode(*pGraphNode, graph, pDependencies, numDependencies, capture);
190195
return status;

‎projects/clr/hipamd/src/hip_memory.cpp‎

Lines changed: 149 additions & 34 deletions
Original file line numberDiff line numberDiff line change
@@ -2763,13 +2763,6 @@ hipError_t ihipMemcpy3D_validate(const hipMemcpy3DParms* p) {
27632763
// not have the same element size.
27642764
if ((p->srcArray != nullptr) && (p->dstArray != nullptr)) {
27652765
if (hip::getElementSize(p->srcArray) != hip::getElementSize(p->dstArray)) {
2766-
return hipErrorInvalidValue;
2767-
}
2768-
2769-
// If both src and dst are arrays, verify they have the same extents
2770-
if (p->srcArray->width != p->dstArray->width ||
2771-
p->srcArray->height != p->dstArray->height ||
2772-
p->srcArray->depth != p->dstArray->depth) {
27732766
return hipErrorInvalidValue;
27742767
}
27752768
}
@@ -3015,6 +3008,26 @@ static hipError_t EnqueueBatchCommands(std::vector<std::vector<Operation>>& oper
30153008
return hipSuccess;
30163009
}
30173010

3011+
// ================================================================================================
3012+
// Releases the batch commands that EnqueueBatchCommands placed on other GPUs' null streams.
3013+
// When status is hipSuccess, the caller's stream first receives a marker that waits for them.
3014+
static hipError_t EnqueueBatchMarker(hipError_t status, hip::Stream& stream,
3015+
amd::Command::EventWaitList& marker_wait_list) {
3016+
if (status == hipSuccess && !marker_wait_list.empty()) {
3017+
amd::Command* dependent_marker = new amd::Marker(stream, true, marker_wait_list);
3018+
if (dependent_marker == nullptr) {
3019+
status = hipErrorOutOfMemory;
3020+
} else {
3021+
dependent_marker->enqueue();
3022+
dependent_marker->release();
3023+
}
3024+
}
3025+
for (amd::Event* cmd : marker_wait_list) {
3026+
cmd->release();
3027+
}
3028+
return status;
3029+
}
3030+
30183031
// ================================================================================================
30193032
// Returns the attribute flags that apply to copy `copyIdx`.
30203033
static inline unsigned int getBatchCopyFlags(hipMemcpyAttributes* attrs, size_t* attrsIdxs,
@@ -3242,29 +3255,7 @@ hipError_t ihipMemcpyBatch(void** dsts, void** srcs, size_t* sizes, size_t count
32423255
stream_wait_cmd->release();
32433256
}
32443257

3245-
if (status != hipSuccess) {
3246-
for (auto* cmd : marker_wait_list) {
3247-
cmd->release();
3248-
}
3249-
return status;
3250-
}
3251-
3252-
if (!marker_wait_list.empty()) {
3253-
amd::Command* dependent_marker = new amd::Marker(stream, true, marker_wait_list);
3254-
if (dependent_marker == nullptr) {
3255-
for (auto* cmd : marker_wait_list) {
3256-
cmd->release();
3257-
}
3258-
return hipErrorOutOfMemory;
3259-
}
3260-
dependent_marker->enqueue();
3261-
dependent_marker->release();
3262-
for (auto* cmd : marker_wait_list) {
3263-
cmd->release();
3264-
}
3265-
}
3266-
3267-
return hipSuccess;
3258+
return EnqueueBatchMarker(status, stream, marker_wait_list);
32683259
}
32693260

32703261
// ================================================================================================
@@ -3339,22 +3330,146 @@ hipError_t hipMemcpyBatchAsync(void** dsts, void** srcs, size_t* sizes, size_t c
33393330
hipError_t hipMemcpy3DBatchAsync(size_t numOps, struct hipMemcpy3DBatchOp* opList, size_t* failIdx,
33403331
unsigned long long flags, hipStream_t stream) {
33413332
HIP_INIT_API(hipMemcpy3DBatchAsync, numOps, opList, failIdx, flags, stream);
3333+
size_t ignored_fail_idx = SIZE_MAX;
3334+
if (failIdx == nullptr) {
3335+
failIdx = &ignored_fail_idx;
3336+
}
33423337
if (flags != 0 || opList == nullptr || numOps == 0) {
3338+
*failIdx = SIZE_MAX;
33433339
HIP_RETURN(hipErrorInvalidValue);
33443340
}
33453341
if (!hip::isValid(stream)) {
33463342
HIP_RETURN(hipErrorInvalidResourceHandle);
33473343
}
33483344
CHECK_STREAM_DETACHED_API(stream);
33493345

3350-
hipError_t status = hipSuccess;
3351-
33523346
*failIdx = SIZE_MAX;
3353-
for (int i = 0; i < numOps; ++i) {
3347+
hip::Stream* hip_stream = hip::getStream(stream);
3348+
3349+
std::vector<std::vector<amd::BatchCopyRectOp>> copy_ops_by_device(g_devices.size());
3350+
std::vector<std::pair<size_t, HIP_MEMCPY3D>> per_entry_copies;
3351+
std::vector<HIP_MEMCPY3D> host_copies;
3352+
for (size_t i = 0; i < numOps; ++i) {
3353+
const hipMemcpySrcAccessOrder access_order = opList[i].srcAccessOrder;
3354+
if (access_order < hipMemcpySrcAccessOrderStream || access_order > hipMemcpySrcAccessOrderAny) {
3355+
*failIdx = i;
3356+
HIP_RETURN(hipErrorInvalidValue);
3357+
}
33543358
hipMemcpy3DParms parms = getMemcpy3DParms(opList[i]);
3355-
status = ihipMemcpy3D(&parms, stream, true);
3359+
hipError_t status = ihipMemcpy3D_validate(&parms);
3360+
if (status != hipSuccess) {
3361+
*failIdx = i;
3362+
HIP_RETURN(status);
3363+
}
3364+
HIP_MEMCPY3D desc = getDrvMemcpy3DDesc(parms);
3365+
if (desc.WidthInBytes == 0 || desc.Height == 0 || desc.Depth == 0) {
3366+
*failIdx = i;
3367+
HIP_RETURN(hipErrorInvalidValue);
3368+
}
3369+
3370+
hipMemoryType src_memory_type;
3371+
hipMemoryType dst_memory_type;
3372+
ihipCopyMemParamSet(&desc, src_memory_type, dst_memory_type);
3373+
3374+
amd::Coord3D src_origin = {desc.srcXInBytes, desc.srcY, desc.srcZ};
3375+
amd::Coord3D dst_origin = {desc.dstXInBytes, desc.dstY, desc.dstZ};
3376+
amd::Coord3D copy_region = {desc.WidthInBytes, desc.Height, desc.Depth};
3377+
amd::BufferRect src_rect;
3378+
amd::BufferRect dst_rect;
3379+
amd::Image* src_image = nullptr;
3380+
amd::Image* dst_image = nullptr;
3381+
status = ihipDrvMemcpy3D_validate(&desc, src_origin, dst_origin, copy_region, &src_rect,
3382+
&dst_rect, &src_image, &dst_image);
33563383
if (status != hipSuccess) {
33573384
*failIdx = i;
3385+
HIP_RETURN(status);
3386+
}
3387+
if (src_memory_type == hipMemoryTypeArray || dst_memory_type == hipMemoryTypeArray) {
3388+
per_entry_copies.emplace_back(i, desc);
3389+
continue;
3390+
}
3391+
3392+
amd::Memory* src_memory = nullptr;
3393+
amd::Memory* dst_memory = nullptr;
3394+
hip::MemcpyType type;
3395+
if (src_memory_type == hipMemoryTypeHost) {
3396+
type = (dst_memory_type == hipMemoryTypeHost) ? hipHostToHost : hipWriteBuffer;
3397+
} else if (dst_memory_type == hipMemoryTypeHost) {
3398+
type = hipReadBuffer;
3399+
} else {
3400+
size_t src_offset = 0;
3401+
size_t dst_offset = 0;
3402+
getMemoryObjectPairs(hip::getCurrentDevice(), desc.srcDevice, desc.dstDevice, src_memory,
3403+
dst_memory, src_offset, dst_offset);
3404+
const bool src_on_host = getMemoryType(src_memory) == hipMemoryTypeHost &&
3405+
!IsManagedMemory(src_memory->getMemFlags());
3406+
const bool dst_on_host = getMemoryType(dst_memory) == hipMemoryTypeHost &&
3407+
!IsManagedMemory(dst_memory->getMemFlags());
3408+
type = (src_on_host && dst_on_host)
3409+
? hipHostToHost
3410+
: ihipGetMemcpyType(src_memory, dst_memory, hipMemcpyDefault);
3411+
}
3412+
3413+
switch (type) {
3414+
case hipCopyBuffer:
3415+
case hipCopyBufferSDMA: {
3416+
const int device_id = (getMemoryType(src_memory) == hipMemoryTypeDevice &&
3417+
getMemoryType(dst_memory) == hipMemoryTypeHost)
3418+
? src_memory->getUserData().deviceId
3419+
: dst_memory->getUserData().deviceId;
3420+
copy_ops_by_device[device_id].emplace_back(src_memory, dst_memory, src_rect, dst_rect,
3421+
copy_region);
3422+
break;
3423+
}
3424+
case hipCopyBufferP2P: {
3425+
const int device_id = hip_stream->DeviceId();
3426+
copy_ops_by_device[device_id].emplace_back(src_memory, dst_memory, src_rect, dst_rect,
3427+
copy_region);
3428+
break;
3429+
}
3430+
case hipHostToHost:
3431+
host_copies.push_back(desc);
3432+
break;
3433+
case hipWriteBuffer:
3434+
case hipReadBuffer:
3435+
per_entry_copies.emplace_back(i, desc);
3436+
break;
3437+
}
3438+
}
3439+
3440+
if (!host_copies.empty()) {
3441+
hip_stream->finish();
3442+
for (const HIP_MEMCPY3D& desc : host_copies) {
3443+
for (size_t z = 0; z < desc.Depth; ++z) {
3444+
for (size_t y = 0; y < desc.Height; ++y) {
3445+
const size_t src_offset =
3446+
((desc.srcZ + z) * desc.srcHeight + desc.srcY + y) * desc.srcPitch + desc.srcXInBytes;
3447+
const size_t dst_offset =
3448+
((desc.dstZ + z) * desc.dstHeight + desc.dstY + y) * desc.dstPitch + desc.dstXInBytes;
3449+
std::memcpy(static_cast<char*>(desc.dstHost) + dst_offset,
3450+
static_cast<const char*>(desc.srcHost) + src_offset, desc.WidthInBytes);
3451+
}
3452+
}
3453+
}
3454+
}
3455+
3456+
amd::Command::EventWaitList marker_wait_list;
3457+
amd::Command* stream_wait_cmd = hip_stream->getLastQueuedCommand(true);
3458+
hipError_t status = EnqueueBatchCommands<amd::BatchCopyMemoryRectCommand>(
3459+
copy_ops_by_device, ROCCLR_COMMAND_BATCH_COPY_BUFFER_RECT, *hip_stream, true, stream_wait_cmd,
3460+
marker_wait_list);
3461+
if (stream_wait_cmd != nullptr) {
3462+
stream_wait_cmd->release();
3463+
}
3464+
status = EnqueueBatchMarker(status, *hip_stream, marker_wait_list);
3465+
if (status != hipSuccess) {
3466+
HIP_RETURN(status);
3467+
}
3468+
3469+
for (auto& [index, desc] : per_entry_copies) {
3470+
status = ihipMemcpyParam3D(&desc, stream, true);
3471+
if (status != hipSuccess) {
3472+
*failIdx = index;
33583473
break;
33593474
}
33603475
}

‎projects/clr/hipamd/src/profiler/hip_clr_profiler.cpp‎

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -298,6 +298,8 @@ static HipCopyKindExt ToCopyKindExt(uint32_t cl_kind) {
298298
case CL_COMMAND_COPY_IMAGE_TO_BUFFER: return HIP_COPY_KIND_IMAGE_TO_BUFFER_EXT;
299299
case CL_COMMAND_FILL_BUFFER: return HIP_COPY_KIND_FILL_EXT;
300300
case ROCCLR_COMMAND_BATCH_COPY_BUFFER: return HIP_COPY_KIND_BATCH_EXT;
301+
case ROCCLR_COMMAND_BATCH_COPY_BUFFER_RECT:
302+
return HIP_COPY_KIND_BATCH_EXT;
301303
case ROCCLR_COMMAND_BATCH_WRITE_BUFFER:
302304
return HIP_COPY_KIND_BATCH_EXT;
303305
case ROCCLR_COMMAND_BATCH_READ_BUFFER:

‎projects/clr/rocclr/device/blit.cpp‎

Lines changed: 11 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -361,6 +361,17 @@ bool HostBlitManager::copyBufferBatch(const std::vector<amd::BatchCopyOp>& copyO
361361
return true;
362362
}
363363

364+
bool HostBlitManager::CopyBufferRectBatch(const std::vector<amd::BatchCopyRectOp>& copy_ops) const {
365+
for (const amd::BatchCopyRectOp& op : copy_ops) {
366+
device::Memory* src_dev_mem = op.src_memory->getDeviceMemory(dev_);
367+
device::Memory* dst_dev_mem = op.dst_memory->getDeviceMemory(dev_);
368+
if (!copyBufferRect(*src_dev_mem, *dst_dev_mem, op.src_rect, op.dst_rect, op.size)) {
369+
return false;
370+
}
371+
}
372+
return true;
373+
}
374+
364375
bool HostBlitManager::WriteBufferBatch(
365376
const std::vector<amd::BatchWriteMemoryOp>& write_ops) const {
366377
for (const amd::BatchWriteMemoryOp& op : write_ops) {

‎projects/clr/rocclr/device/blit.hpp‎

Lines changed: 10 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -150,6 +150,11 @@ class BlitManager {
150150
const std::vector<amd::BatchCopyOp>& copyOps //!< Batch of copy operations
151151
) const = 0;
152152

153+
//! Copies multiple pitched boxes between buffer objects in a batch
154+
virtual bool CopyBufferRectBatch(
155+
const std::vector<amd::BatchCopyRectOp>& copy_ops //!< Batch of box copy operations
156+
) const = 0;
157+
153158
//! Copies pageable host-to-device operations in a batch
154159
virtual bool WriteBufferBatch(
155160
const std::vector<amd::BatchWriteMemoryOp>& write_ops //!< Batch of write operations
@@ -365,6 +370,11 @@ class HostBlitManager : public device::BlitManager {
365370
const std::vector<amd::BatchCopyOp>& copyOps //!< Batch of copy operations
366371
) const;
367372

373+
//! Copies multiple pitched boxes between buffer objects in a batch
374+
bool CopyBufferRectBatch(
375+
const std::vector<amd::BatchCopyRectOp>& copy_ops //!< Batch of box copy operations
376+
) const override;
377+
368378
//! Copies pageable host-to-device operations in a batch
369379
virtual bool WriteBufferBatch(
370380
const std::vector<amd::BatchWriteMemoryOp>& write_ops //!< Batch of write operations

‎projects/clr/rocclr/device/device.hpp‎

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -52,6 +52,7 @@ class FillMemoryCommand;
5252
class CopyMemoryCommand;
5353
class CopyMemoryP2PCommand;
5454
class BatchCopyMemoryCommand;
55+
class BatchCopyMemoryRectCommand;
5556
class BatchWriteMemoryCommand;
5657
class BatchReadMemoryCommand;
5758
class MapMemoryCommand;
@@ -1333,6 +1334,7 @@ class VirtualDevice : public amd::ReferenceCountedObject {
13331334
virtual void submitCopyMemory(amd::CopyMemoryCommand& cmd) = 0;
13341335
virtual void submitCopyMemoryP2P(amd::CopyMemoryP2PCommand& cmd) = 0;
13351336
virtual void submitBatchCopyMemory(amd::BatchCopyMemoryCommand& cmd) = 0;
1337+
virtual void SubmitBatchCopyMemoryRect(amd::BatchCopyMemoryRectCommand& cmd) = 0;
13361338
virtual void SubmitBatchWriteMemory(amd::BatchWriteMemoryCommand& cmd) = 0;
13371339
virtual void SubmitBatchReadMemory(amd::BatchReadMemoryCommand& cmd) = 0;
13381340
virtual void submitMapMemory(amd::MapMemoryCommand& cmd) = 0;

‎projects/clr/rocclr/device/pal/palvirtual.cpp‎

Lines changed: 37 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -2126,6 +2126,43 @@ void VirtualGPU::submitBatchCopyMemory(amd::BatchCopyMemoryCommand& cmd) {
21262126
profilingEnd(cmd);
21272127
}
21282128

2129+
void VirtualGPU::SubmitBatchCopyMemoryRect(amd::BatchCopyMemoryRectCommand& cmd) {
2130+
// Make sure VirtualGPU has an exclusive access to the resources
2131+
std::scoped_lock lock(execution());
2132+
2133+
profilingBegin(cmd);
2134+
2135+
const std::vector<amd::BatchCopyRectOp>& copy_ops = cmd.CopyOps();
2136+
2137+
device::Memory::SyncFlags sync_flags;
2138+
sync_flags.skipEntire_ = false;
2139+
for (const amd::BatchCopyRectOp& op : copy_ops) {
2140+
Memory* src_dev_mem = dev().getGpuMemory(op.src_memory);
2141+
Memory* dst_dev_mem = dev().getGpuMemory(op.dst_memory);
2142+
if (src_dev_mem == nullptr || dst_dev_mem == nullptr) {
2143+
LogError("SubmitBatchCopyMemoryRect: Invalid memory objects!");
2144+
cmd.setStatus(CL_INVALID_MEM_OBJECT);
2145+
profilingEnd(cmd);
2146+
return;
2147+
}
2148+
dst_dev_mem->syncCacheFromHost(*this, sync_flags);
2149+
src_dev_mem->syncCacheFromHost(*this);
2150+
}
2151+
2152+
if (!blitMgr().CopyBufferRectBatch(copy_ops)) {
2153+
LogError("SubmitBatchCopyMemoryRect: Batch rect copy failed!");
2154+
cmd.setStatus(CL_OUT_OF_RESOURCES);
2155+
profilingEnd(cmd);
2156+
return;
2157+
}
2158+
2159+
for (const amd::BatchCopyRectOp& op : copy_ops) {
2160+
op.dst_memory->signalWrite(&dev());
2161+
}
2162+
2163+
profilingEnd(cmd);
2164+
}
2165+
21292166
void VirtualGPU::SubmitBatchWriteMemory(amd::BatchWriteMemoryCommand& cmd) {
21302167
// Make sure VirtualGPU has an exclusive access to the resources
21312168
std::scoped_lock lock(execution());

‎projects/clr/rocclr/device/pal/palvirtual.hpp‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -317,6 +317,7 @@ class VirtualGPU : public device::VirtualDevice {
317317
void submitCopyMemory(amd::CopyMemoryCommand& vcmd);
318318
void submitCopyMemoryP2P(amd::CopyMemoryP2PCommand& vcmd);
319319
void submitBatchCopyMemory(amd::BatchCopyMemoryCommand& vcmd);
320+
void SubmitBatchCopyMemoryRect(amd::BatchCopyMemoryRectCommand& cmd);
320321
void SubmitBatchWriteMemory(amd::BatchWriteMemoryCommand& cmd);
321322
void SubmitBatchReadMemory(amd::BatchReadMemoryCommand& cmd);
322323
void submitMapMemory(amd::MapMemoryCommand& vcmd);

0 commit comments

Comments
 (0)