Data Replication

Copy Kinds

Common copy kinds are as follows.

The synchronous API aclrtMemcpy returns after the copy is complete. The asynchronous API aclrtMemcpyAsync dispatches the copy task to a specified stream. If the API returns, it only indicates that the task has been dispatched, not that the copy task is complete. Before reading the destination memory, it is necessary to call aclrtSynchronizeStream or aclrtSynchronizeDevice, or use the event or other mechanisms, to ensure that the task is complete.

Host-Device Copy

Data copy between the host and device is the most common input and output movement mode. The typical process is as follows: Allocate host memory and device memory, copy the input from the host to the device, launch the operator on the stream, and finally copy the output from the device back to the host.

aclrtStream stream = nullptr;
aclrtCreateStream(&stream);

void *hostIn = nullptr;
void *hostOut = nullptr;
void *devIn = nullptr;
void *devOut = nullptr;
aclrtMallocHost(&hostIn, size);
aclrtMallocHost(&hostOut, size);
aclrtMalloc(&devIn, size, ACL_MEM_MALLOC_HUGE_FIRST);
aclrtMalloc(&devOut, size, ACL_MEM_MALLOC_HUGE_FIRST);

// Copy the input data from the host to the device.
aclrtMemcpyAsync(devIn, size, hostIn, size, ACL_MEMCPY_HOST_TO_DEVICE, stream);

// Dispatch the compute task, which is executed in order with the preceding copy task in the same stream.
myKernel<<<numBlocks, nullptr, stream>>>(devIn, devOut);

// Copy the output data from the device back to the host.
aclrtMemcpyAsync(hostOut, size, devOut, size, ACL_MEMCPY_DEVICE_TO_HOST, stream);

// Read hostOut after the copy and compute tasks in the same stream are complete.
aclrtSynchronizeStream(stream);

If a synchronous copy API is used, the API returns after the copy is complete, but does not perform implicit stream synchronization. If there are other asynchronous tasks occurring before or after the copy, stream ordering, events, or explicit synchronization must be used to ensure data dependencies.

Asynchronous Copy of Non-page-locked Host Memory

The host memory allocated by traditional APIs such as malloc and mmap is non-page-locked memory. Asynchronous APIs such as aclrtMemcpyAsync, aclrtMemcpy2dAsync, and aclrtMemcpyBatchAsync support non-page-locked host memory. Note that:

• When the host memory involved in the copy is page-locked, the asynchronous copy API only dispatches the copy task and returns immediately. The host thread can continue executing other tasks, while the copy operation is completed in order within the stream.
• If the host memory involved in the copy is non-page-locked memory, the asynchronous copy API returns only after the memory copy task is complete. Although the API name contains Async, the host thread cannot be executed in parallel with the copy operation, and data transfer cannot fully overlap with compute.
• To reliably achieve overlap between copy and compute, it is advised to use aclrtMallocHost to allocate page-locked memory or register existing host memory as page-locked memory using aclrtHostRegisterV2 on supported systems.
• When ACL Graph captures an asynchronous memory copy task, the host memory, if involved in the copy, must be the page-locked memory allocated by a Runtime API. Otherwise, an error will be returned during the capture.

Therefore, non-page-locked host memory is more suitable for low-frequency, temporary, and concurrency-insensitive data movement, while page-locked host memory should be used for performance-sensitive input and output channels.

Intra-Device Copy

When aclrtMemcpy or aclrtMemcpyAsync is used for device-to-device copy, set the copy type to ACL_MEMCPY_DEVICE_TO_DEVICE. This enumeration indicates intra-device or inter-device data copy. Runtime determines the actual copy path based on the source and destination addresses. You do not need to explicitly specify ACL_MEMCPY_INNER_DEVICE_TO_DEVICE or ACL_MEMCPY_INTER_DEVICE_TO_DEVICE.

ACL_MEMCPY_INNER_DEVICE_TO_DEVICE indicates only intra-device copy and needs to be explicitly specified in descriptor-based copy scenarios, for example, when using aclrtMemcpyAsyncWithDesc API. The descriptor mode is suitable for scenarios where copy parameters can be reused. You need to obtain the descriptor size and set the source address, destination address, and length, and then use the descriptor in the stream to dispatch a copy task.

Cross-device Copy

Cross-device copy is used for data movement between different devices in the same process. Before calling the copy API, check whether data exchange is supported between two devices and enable, based on the access direction, data exchange between the two devices.

int32_t canAccess = 0;
aclrtDeviceCanAccessPeer(&canAccess, srcDeviceId, dstDeviceId);
if (canAccess != 0) {
    aclrtSetDevice(dstDeviceId);
    aclrtDeviceEnablePeerAccess(srcDeviceId, 0);
}

For details about cross-device memory copy, see Cross-Device Data Exchange.

2D and Batch Copy

The 2D copy APIs aclrtMemcpy2d and aclrtMemcpy2dAsync are applicable to matrices, images, or pitched 2D data. When calling the APIs, you need to specify the pitch for both source and destination memory, alongside the width and height of the region to be copied. For asynchronous 2D copy involving host memory, the rule remains: Page-locked host memory is required for truly asynchronous operation.

The batch copy APIs aclrtMemcpyBatch and aclrtMemcpyBatchAsync, as well as their V2 versions, are applicable to the scenario where multiple segments are copied from the host to the device or from the device to the host at a time. When performing batch copy operations, note that:

• Copies within a batch are not guaranteed to execute in array order.
• The destination address, source address, and copy size arrays must match in length for each copy operation.
• The copy kind within a batch can only be either host-to-device or device-to-host.
• If an API is for trial use, the product description and restrictions mentioned in the API reference prevail.

Initializing Memory

aclrtMemset and aclrtMemsetAsync are used to initialize device memory. The synchronous API returns after the initialization is complete. The asynchronous API dispatches the memory initialization task to a specified stream. If the API returns, it only indicates that the task has been dispatched, not that the copy task is complete.

It is recommended that the device memory used as a synchronization flag or compute input be explicitly initialized before use to avoid reading undefined data.