Using Unified Buffer to Improve Memory Access Efficiency

This performance optimization suggestion applies to the following models:

• Atlas 350 Accelerator Card

[Priority] High

[Description] The granularity for SIMT to access the global memory is 128 bytes. When data in the global memory is randomly accessed, the memory access efficiency is low. When the amount of data to be accessed is far less than the maximum available Unified Buffer space (256 KB – 8 KB reserved for the system – 32 KB minimum Dcache), the SIMD transfer interface can be used to transfer data from the Global Memory to the Unified Buffer. This allows SIMT programming to directly read data from the Unified Buffer, thereby improving memory access efficiency and enhancing the overall performance of the operator.

[Example] The gather operator implemented in SIMD and SIMT hybrid programming mode is used as an example. This operator obtains 65536 pieces of data with specified indexes from a one-dimensional vector with a length of 8192. By pre-transferring the input data to the Unified Buffer, the efficiency of discrete memory access is improved, thereby enhancing the overall performance of the operator.

The SIMT thread hierarchy is as follows:

• Number of thread blocks: 64
• Number of threads in a single thread block: 1024

For a complete operator example, see using UB to improve memory access efficiency for hybrid programming of SIMD and SIMT.

[Negative Example]

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30

namespace {
    constexpr uint32_t THREAD_COUNT = 1024;
    constexpr uint32_t INPUT_SIZE = 8192;
    constexpr uint32_t INDEX_SIZE = 65536;
}

__simt_vf__ __launch_bounds__(THREAD_COUNT) inline void simt_gather(
    __gm__ float* input,
    __gm__ uint32_t* index,
    __gm__ float* output)
{
    int32_t idx = blockIdx.x * blockDim.x + threadIdx.x;

    if (idx >= INDEX_SIZE) {
        return;
    }

    uint32_t gather_idx = index[idx];
    if (gather_idx > INPUT_SIZE) {
        return;
    }

    output[idx] = input[gather_idx];
}


__global__ __vector__ void gather_kernel(__gm__ float* input, __gm__ uint32_t* index, __gm__ float* output)
{
    asc_vf_call<simt_gather>(dim3(THREAD_COUNT), input, index, output);
}

[Positive Example]

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39

namespace {
    constexpr uint32_t THREAD_COUNT = 1024;
    constexpr uint32_t INPUT_SIZE = 8192;
    constexpr uint32_t INDEX_SIZE = 65536;
}

__simt_vf__ __launch_bounds__(THREAD_COUNT) inline void simt_gather(
    __ubuf__ float* input,
    __gm__ uint32_t* index,
    __gm__ float* output)
{
    int32_t idx = blockIdx.x * blockDim.x + threadIdx.x;

    if (idx >= INDEX_SIZE) {
        return;
    }

    uint32_t gather_idx = index[idx];
    if (gather_idx >= INPUT_SIZE) {
        return;
    }

    output[idx] = input[gather_idx];
}


__global__ __vector__ void gather_kernel(__gm__ float* input, __gm__ uint32_t* index, __gm__ float* output)
{
    __ubuf__ float input_buf[INPUT_SIZE];
    uint32_t blk_length = INPUT_SIZE * sizeof(float);
    asc_copy_gm2ub_align(input_buf, input, 1, blk_length, 0, 0, false, 0, 0, 0);

    if ASC_IS_AIV {
        asc_sync_notify(PIPE_MTE2, PIPE_V, EVENT_ID0);
        asc_sync_wait(PIPE_MTE2, PIPE_V, EVENT_ID0);
    }

    asc_vf_call<simt_gather>(dim3(THREAD_COUNT), input_buf, index, output);
}

[Performance Comparison]

The following figure shows the pipeline of the negative example. There are two SIMT_LDG instructions in the thread, which correspond to loading data from Global Memory. The occupied range of the second instruction is unevenly distributed, and the instruction startup time varies significantly. In the same thread block, the input data is randomly accessed. 128 bytes of data are loaded during a single memory access. However, for single-precision floating-point numbers, the actual valid data is only 4 bytes, resulting in low memory access efficiency.

Figure 1 Instruction pipeline chart of the negative example (simulation data)

The following figure shows the memory load analysis of the negative example. The data transfer bandwidth from the L2 cache to the Dcache is 10.04 GB/s.

Figure 2 Memory workload analysis of the negative example (board data)

The following figure shows the pipeline of the positive example. There is only one SIMT_LDG instruction that occupies a large range, and the MOV_SRC_TO_DST_ALIGNv2 instruction is added to the MTE2 pipeline.

Figure 3 Instruction pipeline chart of the positive example (simulation data)

The following figure shows the memory load analysis of the positive example. The data transmission bandwidth from the L2 cache to the Dcache is reduced to 1.61 GB/s, and the data transmission bandwidth from the L2 cache to Unified Buffer is increased to 12.93 GB/s.

Figure 4 Positive example memory workload analysis (board data)

The operator running time is about 4.56 μs in the negative example and about 3.57 μs in the positive example. The overall performance is improved by about 21.71%.

Figure 5 Operator running time in the negative example (board data)

Figure 6 Positive example operator running time (board data)