Programming Model

Host and Device Programming Model

Host: A host is an x86 or Arm server CPU that connects to one or more devices through a bus. The host utilizes the compute capabilities (for example, NN) provided by the device to implement services.

Device (NPU): A device refers to the hardware equipped with AI processors. It connects to the host through a bus (such as PCIe and HCCS), providing compute capabilities (for example, NN). HCCS stands for Huawei Cache Coherence System.

The following figure shows the relationship between the host and the device.

In summary:

The host and device each possess independent memory spaces. Runtime provides memory allocation APIs for both the host and device, along with memory copy APIs to transfer data between them. During programming, you must distinguish between host and device memory allocations and explicitly call the memory copy APIs to transfer data from the host to the device. This ensures that the device hardware accelerator achieves optimal performance when accessing its local memory.
The host and device use the asynchronous parallel execution mode.
After the host delivers a task to the device, the host does not wait for the task to complete on the device and immediately returns. The device then starts scheduling and executing the delivered task. The CPU on the host can work in parallel with the accelerator on the device. The API for delivering asynchronous tasks usually contains a stream parameter, indicating that the task is delivered to the corresponding stream for execution.

When the host needs to obtain the computation result from the device, it must explicitly call a synchronization API. The synchronization API blocks the host CPU until the task on the device is complete.

This asynchronous parallel execution mechanism effectively hides the host processing time or the data transfer latency between the host and the device, improving throughput and shortening the end-to-end execution time.
The asynchronous task is delivered to the stream for execution.
Tasks in a stream are executed in order, and tasks in different streams are executed in parallel. For example, in the following figure, the host starts the Kernel 1, Kernel 2, and Kernel 3 tasks in sequence. The execution sequence is as follows:
- Kernel1 and Kernel3 are in the same stream. Therefore, Kernel 3 is started only after Kernel 1 is completed.
- Kernel 2 is not in the same stream as Kernel 1 and Kernel 3. Therefore, Kernel 2 can be executed in parallel with Kernel 1 and Kernel 3.

Typical Execution Process

The following figure shows the typical execution process based on Runtime programming.

API Call	Handling
int32_t devId=0; aclrtSetDevice(devId);	Initialize device 0, including the following operations: 1. Create and initialize a device object. 2. Create a default context for the device. 3. Create a default stream for the default context. 4. Start the CPU executor process on the device.
aclrtStream stream1; aclrtCreateStream(&stream1);	Create a stream in the current context. 1. Call the driver to create a task queue. 2. Create a stream object on Runtime and associate the stream with the task queue. 3. Add the stream to the current context for management.
uint64_t size=1024; void hostPtr=nullptr; aclrtMallocHost*(&hostPtr, size);	Allocate the host memory.
void devPtr=nullptr; aclrtMalloc*(&devPtr, size);	Allocate the device memory.
aclrtMemcpy(devPtr, size, hostPtr, size, ACL_MEMCPY_HOST_TO_DEVICE)	Synchronous memory copy from the host to the device. The main operations are as follows: 1. Construct a Direct Memory Access (DMA) descriptor based on the source and destination addresses, and then deliver a DMA task. 2. Wait until the DMA task is complete and the API returns a result.
myKernel<<<numBlocks, nullptr, stream1>>>(devPtr);	Deliver the myKernel compute task to stream 1. numBlocks specifies the number of blocks in the compute task. The detailed operations are as follows: Register and load the binary code of myKernel to the device (only for the first time). Copy the execution parameters of myKernel to the device. Allocate task management resources and create a task descriptor. Deliver the task to the task queue corresponding to the stream. After the task is delivered, the API returns a result asynchronously. The task scheduler starts scheduling. The detailed operations are as follows: The scheduler uses the strict priority (SP) mode for scheduling based on the priority of the task queue. Parse the task descriptor delivered in the task queue. Check whether the corresponding acceleration unit is idle based on the task type. If there are idle cores and the number of cores meets the task requirements, schedule the task to the corresponding acceleration unit (for example, a CPU operator is scheduled to the CPU executor). The acceleration unit executes the task. The scheduler waits until the acceleration unit completes the execution and updates the task status and the busy/idle status of the acceleration unit. If an exception occurs during task execution, the scheduler handles the exception. Continue to schedule other tasks.
aclrtSynchronizeStream(stream1);	Stream synchronization. The detailed operations are as follows: 1. The API blocks the current CPU thread. 2. Synchronize the execution status of the task queue corresponding to the stream using the polling + interrupt mechanism. 3. After all tasks in the stream are executed, the API returns a result.
aclrtMemcpy(hostPtr, size, devPtr, size, ACL_MEMCPY_DEVICE_TO_HOST);	Synchronous memory copy from the device to the host. Copy the result back to the host. 1. Construct a DMA descriptor based on the source and destination addresses, and then deliver a DMA task. 2. Wait until the DMA task is complete and the API returns a result.
aclrtFree(devPtr); aclrtFreeHost(hostPtr);	Free the device memory. Free the host memory.
aclrtResetDeviceForce(devId);	Reset the device.

Main Programming Concepts of Runtime

A host is the abstraction of Runtime on the host.
A device represents the abstraction of the hardware device where the AI processor is installed. Typically, one host maps to N devices. The user application can call ACL APIs such as aclrtSetDevice to specify the compute device.
A context is the logical operating environment of a device. The relationship between contexts and devices is N:1, meaning that each context must belong to a unique device. A context manages the lifetime of its runtime resource objects (including streams, events, and Notifys, but excluding memory). Each unique context has isolated objects, and they cannot be synchronized with one another. Runtime errors are also isolated at the context level.
A stream is a logical task execution queue provided by a device. Tasks can be asynchronously added to a stream, and tasks in the same stream are strictly executed in FIFO mode. The relationship between streams and contexts is N:1. A stream must belong to a unique context.
A task is an execution task that can be added to a stream. Tasks can be classified into compute tasks, memory copy tasks, and event synchronization tasks. The relationship between tasks and streams is N:1. A task will be added to a unique stream.

The following figure shows the relationship between the device, context, and stream.

Thread-Context Association

Most Runtime APIs do not have the device ID parameter because the device on which these APIs operate is obtained from the context associated with the call thread. Therefore, when a host thread calls a Runtime API, the following requirements must be met:

A thread (CPU thread on the host) can correctly call Runtime APIs only after being associated with a context.
A thread can be associated with only one context at a time.
Applications can explicitly create contexts to meet the service requirements of running resource isolation. In this scenario, the context is visible to all threads in the same process, and the threads can switch the context by calling aclrtGetCurrentContext and aclrtSetCurrentContext.

The following sample code describes the association and switching process between a thread and a context when the thread calls context-related APIs. The sample code is for reference only and cannot be directly copied for compilation and running.

       
            // Initially, the thread is not associated with any context.
aclInit(nullptr);
aclrtSetDevice(0);  // aclrtSetDevice creates a default context and associates the thread with the default context.
// Context associated with the thread: default context

aclrtContext ctx1, ctx2, current_ctx;
aclrtCreateContext(&ctx1, 0); // Device 0 explicitly creates ctx1 and associates the thread with ctx1.
// Context associated with the thread: ctx1
aclrtGetCurrentContext(&current_ctx); // Obtain the context associated with the current thread. In this case, the returned current_ctx is ctx1.
// Context associated with the thread: ctx1

aclrtCreateContext(&ctx2, 0); // Device 0 explicitly creates ctx2 and associates the thread with ctx2.
// Context associated with the thread: ctx2

aclrtSetCurrentContext(current_ctx); // Switch the context. Because current_ctx=ctx1, the thread is associated with ctx1.
// Context associated with the thread: ctx1

aclrtSetCurrentContext(ctx2); // Switch the context.
// Context associated with the thread: ctx2
.....

aclrtDestroyContext(ctx2);  // The current thread is associated with ctx2. When ctx2 is destroyed, the thread is also disassociated from ctx2.
// Context associated with the thread: NA

// ctx2 has been destroyed. Switch to another ctx (for example, ctx1).
aclrtSetCurrentContext(ctx1);
// Context associated with the thread: ctx1
.....

aclrtDestroyContext(ctx1);  // The current thread is associated with ctx1. When ctx1 is destroyed, the thread is also disassociated from ctx1.
// Context associated with the thread: NA
aclrtResetDeviceForce(0);

Application Scenarios of Default Contexts and Default Streams

A context and a stream must exist before any requests are delivered to the device, and these can be created either explicitly or implicitly. The implicitly created context and stream are the default context and stream.
To pass the default stream to any API call, pass NULL directly.
If the default context is used, aclrtGetCurrentContext, aclrtSetCurrentContext, and aclrtDestroyContext are not available.
The default context and default stream is applicable to simple apps where only one compute device is needed. For a multithreaded application, you are advised to use the explicitly created context and stream.

The following code snippet shows key steps only, and is not ready to be built or run.

      
           // ......
uint32_t numBlocks = 32;
uint64_t size = 1024;
void *devPtr = nullptr;
aclInit(nullptr);
aclrtSetDevice(0); 
/* A default context has been created. In the default context, a default stream is created and is available in the current thread. */

......
aclrtMalloc(&devPtr, size );
myKernel<<<numBlocks, nullptr, nullptr>>>(devPtr);  // The third parameter nullptr in <<< >>> indicates that the operation is performed on the default stream.
aclrtSynchronizeStream(nullptr); 

/* Output the result as required when all compute tasks are complete. */
......
aclrtResetDeviceForce(0);  // Reset device 0. The lifetime of the corresponding default context and default stream ends.

Suggestions for Compiling High-Performance Applications

Comply with the following basic principles:

Perform asynchronous execution on the host and device. The host must be able to deliver sufficient tasks to the device in a timely manner to ensure that the acceleration hardware is always in the compute state.
Use the multi-stream mode to fully utilize different types of hardware accelerators on the device to implement concurrent execution. As shown in the following figure, CANN Runtime can schedule multiple types of hardware accelerators. The supported hardware accelerators vary with the AI processor generations. For details, see the actual hardware user manual.

The following is recommended:

Create and use multiple streams in a single thread. If the performance of a single thread is sufficient to deliver tasks to multiple streams to fully utilize the compute power of the device, the single-thread mode is recommended.
When the performance of a single thread is insufficient, the multi-thread mode can be used to improve the task delivery performance on the host. It is recommended that each thread create and use its own stream to deliver tasks. It is not recommended that multiple threads concurrently deliver tasks to the same stream, as this will introduce lock operations and the tasks delivered by multiple threads are out of order.
If a single task delivered by a stream does not fully utilize the AI Core, multiple streams can be used to deliver tasks that can be executed in parallel to fully utilize AI Core resources.
The AI processor contains multiple types of hardware accelerators, such as AI Core, AI CPU, Digital Vision Pre-Processing (DVPP), and Random (random number generator). These hardware accelerators correspond to different types of tasks. It is recommended that multiple streams be created based on the operator execution hardware.

Parent topic: Quick Start