Memory Hierarchy

Global Memory

The global memory on the device can be accessed by all threads in the entire grid. Its function is similar to that of the random access memory (RAM) in the CPU system. Kernel functions running on the device can directly access the global memory, which is the same as accessing the system memory by the code on the CPU.

The global memory is persistent. Data in the space allocated from the global memory will be retained until the space is released or the application is terminated. You can use runtime APIs to manage the global memory on the device. The host uses aclrtMalloc to allocate the global memory on the device, and uses aclrtMemcpy to copy data from the host to the global memory on the device or from the global memory on the device to the host. The global memory on the device allocated by aclrtMalloc must be freed using the aclrtFree API. For more information about runtime APIs, see Runtime API Reference.

During actual development, you need to allocate and initialize the global memory using runtime APIs before starting the kernel function. During kernel function execution, each thread in the SIMT can read data from and write data to the global memory. After the kernel function is executed, the result written to the global memory can be copied back to the host. The global memory can be accessed by all threads in a grid. Therefore, data races between threads must be strictly avoided.

The following is a simple code example of using the global memory: Arrays x, y, and z are stored in the global memory. The following kernel function is used to implement the global memory access and storage of each thread:

 1
 2
 3
 4
 5
 6
 7
 8
 9
10

__global__ void add_custom(float* x, float* y, float* z, uint64_t total_length)
{
    // Calculate global thread ID
    int32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
    // Maps to the row index of output tensor
    if (idx >= total_length) {
        return;
    }
    z[idx] = x[idx] + y[idx];
}

Shared Memory (Unified Buffer)

The shared memory can be accessed by all threads in the same thread block. It is located inside each Vector Core (AIV). Compared with the global memory, the shared memory has a smaller capacity but higher bandwidth and lower access latency. It can be considered as a high-speed cache resource managed by the user during kernel execution. The shared memory can be accessed by all threads in a thread block. Therefore, data races between threads in the same thread block must be avoided. The asc_syncthreads API can be used to synchronize threads in the same thread block. This function blocks all threads in the thread block until all threads reach the position where the API is called.

You can allocate the shared memory dynamically or statically.

1. Static allocation: Allocate a segment of memory space with a specified size. The space size is determined during compilation and cannot be dynamically modified. You can allocate and use the memory space through arrays. This mode will be supported in later versions.

   1	__ubuf__ half staticBuf[1024];

2. Dynamic allocation: You need to specify the size of the dynamic memory space using the dynUBufSize parameter in <<<>>>. The space size is determined during runtime. In SIMT programming, you can allocate and use the dynamic memory in the following ways. This mode will be supported in later versions.

   1	extern __ubuf__ char dynamicBuf[];

The Unified Buffer not only serves as shared memory, but also reserves some memory space for internal use. Therefore, when allocating shared memory, ensure that the shared memory is not exhausted. As shown in the following figure, the size of the Unified Buffer is 256 KB, which is divided into four main areas based on functions. From the low address to high address, these areas are the static memory, dynamic memory, reserved space, and data cache.

The structure is as follows:

1. The static memory and dynamic memory correspond to the memory allocated in static and dynamic allocation modes, respectively.
2. Reserved space: space reserved for the compiler and Ascend C. The size is fixed at 8 KB.
3. Data cache: The data cache is dedicated for SIMT, which is used to cache data when SIMT threads access the global memory. The size of the data cache space can be configured to range from 32 KB to 128 KB. The actual memory size is affected by the sizes of the static and dynamic memory configured by the user. The simple calculation formula is as follows: Data cache size = UB size (256 KB) – Static memory – Dynamic memory – Reserved space (8 KB). You need to properly configure the sizes of the static and dynamic memory to ensure that the data cache is greater than or equal to 32 KB.

The static memory allocation and the method of allocating the dynamic memory through dynamic arrays are under development and will be supported in later versions.

• If the data cache is less than 32 KB, a verification error will occur.
• In the SIMT scenario, an operator cannot use all Unified Buffer space. In addition to the 8 KB reserved space, at least 32 KB of data cache space must be reserved for SIMT.

Register

As described in the preceding table, the more threads in each thread block, the fewer registers available for each thread. If the number of threads set by the user is too large and the computation complexity of each thread is high, the compiler may temporarily store data in the stack space due to insufficient registers for storing local variables. As a result, register overflow (stack spill) may occur, affecting operator performance. Therefore, you need to properly configure blockDim based on the actual operator complexity.