SIMD and SIMT Programming

Abstract Hardware Architecture

Hybrid programming of SIMD and SIMT on the AI Core combines the parallel compute capability of SIMD with the discrete memory access advantage of SIMT, achieving efficient collaboration between vector-level parallelism and thread-level parallelism. Currently, this technology is supported only by Atlas 350 Accelerator Card.

The entire execution process uses the Vector Function (VF) as the basic scheduling unit. A VF is a basic function block. SIMD and SIMT programming supports flexible switching between SIMD and SIMT execution modes within the same operator. Two different types of VFs can be quickly switched. Each VF represents an independent compute task segment, which usually corresponds to a segment of logic that can be processed in parallel within the operator. This achieves a better balance between performance, energy efficiency, and development efficiency. In SIMD and SIMT programming:

• A kernel function can contain multiple VFs.
• Each VF can be programmed in either SIMD or SIMT mode.
• Different types of VFs can be quickly switched, with the switching granularity being a single VF.
• At any given moment, an AIV core can execute only SIMT or SIMD tasks.

In SIMD and SIMT programming, SIMT simplifies the development of complex operators and irregular control flows. SIMD, on the other hand, implements efficient data parallel processing based on vector registers and instructions, allowing a single instruction to process multiple pieces of data and improving the throughput per cycle. SIMD and SIMT programming allows you to perform refined mapping based on operator characteristics. Regular element-wise operations achieve high bandwidth and compute power utilization through SIMD, while irregular or branch-containing computations are handled by SIMT to mitigate divergence and control complexity. At the system level, this helps improve hardware utilization and energy efficiency. It also facilitates optimizations such as operator fusion and data reuse. An operator can contain both continuous regular computations that are well-suited to SIMD and discrete access tasks that are better handled by SIMT, thereby leveraging the advantages of both SIMD and SIMT within the same operator.

As shown in Figure 1, the internal execution processes of SIMD and SIMT are as follows:

• The Scalar unit emits VFs to the vector function queue.
• The SIMD and SIMT programming mode is switched at the granularity of VF. The data cache (DCache) data in SIMT mode is retained when the VF is switched.
• VFs of SIMD and SIMT are executed in sequence. At any given moment, an AIV core can execute only SIMD or SIMT tasks.
• After the VF execution is complete, the result data is written back to the Unified Buffer or Global Memory.

Figure 1 SIMD and SIMT programming hardware architecture

The following table describes the differences between SIMD and SIMT programming.

Although SIMD and SIMT differ significantly in their programming models and execution mechanisms, they share the following key resources at the hardware level:

• SIMT VFs and SIMD VFs share the instruction cache (ICache) to improve instruction prefetch efficiency.
• SIMT and SIMD share the vector ALU unit, and the functions executed based on this unit and the performance are basically the same.
• A part of the Unified Buffer is shared by SIMT and SIMD, and the other part is used as the data cache of SIMT.

Memory Hierarchy

In the SIMD and SIMT programming scenario, multiple types of memory spaces can be accessed. The following table summarizes the scopes and lifecycles of common memory types.

• The Global Memory has a large storage space and low bandwidth. The lifecycle of the global memory is the same as that of the entire application. It is usually used to store input and output data.
• The shared memory is an independent Unified Buffer (UB) owned by each AIV core. The lifecycle of the shared memory is the same as that of the AIV core function. It is usually used as the cache of the global memory.
• Registers are closest to the vector unit. The SIMT and SIMD modes each have a private register memory.

The following figure shows the overall memory architecture.

Figure 2 SIMD and SIMT programming memory model diagram

In SIMT mode, the working process of each memory is as follows:

• Each thread has an independent register and stack, which are used to store local variables. The number of available registers depends on the number of threads in a thread block. For details, see Table 5.
• All threads in a thread block share the Unified Buffer. The Unified Buffer is accessed by all threads in the thread block, and its lifetime is the same as that of the thread block.
• All threads can access the Global Memory (GM) through the data cache (DCache). The data cache is a cache space divided from the UB. The memory size can be configured in the range from 32 KB to 128 KB. For details about the configuration method, see UB Allocation.

In SIMD mode, the working process of each memory is as follows:

• The register file (RF) of the SIMD contains multiple types of vector compute registers, which are used to store compute data in the SIMD VF function. For details about the register types, see Register Data Types.
• All VF registers in a single core share the UB.
• In SIMD mode, data cannot be directly loaded from the global memory to the vector compute register. Instead, data must be transferred from the global memory to the UB, and then explicitly loaded from the UB to the vector compute register using the Load/Store instructions for computation.

UB Allocation

The total memory space of the UB is 256 KB. As shown in Figure 3, the UB is divided into four main areas based on functions: static memory, dynamic memory, reserved space, and data cache, in ascending order of addresses. The structure is as follows:

1. Static memory: The static memory is a segment of memory space (with a specified size) allocated from the start address of the memory. The size of the static memory is determined during compilation and cannot be dynamically modified.

   1 2	// The static memory is allocated through arrays. Example: __ubuf__ char staticBuf[1024];

2. Dynamic memory (will be supported in later versions): The dynamic memory is located after the static memory. The dynamic memory size is specified by the dynUBufSize parameter in <<<>>>. You can allocate the dynamic memory in the following ways:
   • Allocate the memory by calling relevant TPipe APIs.
   • Allocate the memory by calling the Alloc API of LocalMemAllocator.
   • Allocate the memory using dynamic arrays.

     // The dynamic memory is allocated through dynamic arrays. Example:
     extern __ubuf__ char dynamicBuf[];

   All the preceding three methods allocate memory starting from the end position of the static memory. Address space overlapping will occur if these methods are used at the same time, causing undefined behavior. Therefore, you can use only one of the methods to allocate the dynamic memory.

3. Reserved space: space reserved for the compiler and Ascend C. The size is fixed at 8 KB.
4. Data cache: The data cache is dedicated for SIMT. Its size must be greater than or equal to 32 KB.

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

• Data cache = Total UB size (256 KB) – Static memory – Dynamic memory – Reserved space (8 KB)
• If the data cache is less than 32 KB, a verification error will occur.
• In the SIMD and SIMT programming 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.

Figure 3 UB allocation diagram

Kernel Function Definition

• Kernel function definition methods
  • SIMT VF function definition:

    When defining a SIMT VF kernel function, __launch_bounds__(thread_num) is optional and is used to specify the maximum number of threads to be started by the kernel function during compilation. If thread_num is not configured, the default value 1024 is used.

    In SIMD and SIMT programming, the __simt_vf__ and __gm__ modifiers defined in the SIMT VF kernel function need to be identified separately. For details about the restrictions on SIMT VF function programming, see Appendix.

    1	__simt__vf__ __launch_bounds__(thread_num) inline void simt_vector_function(__ubuf__ float* input, ...)

  • SIMD VF function definition:
    The SIMD VF kernel function is identified by the __simd_vf__ modifier.

    1	__simd_vf__ inline void my_kernel(__gm__ uint8_t* x, __gm__ uint8_t* y, __gm__ uint8_t* z);

    

    The input parameters of SIMD_VF and SIMT_VF support only the Plain Old Data (PoD) data type.

    • PoD data types include basic data types (such as int32_t and float) and arrays and structures composed of these basic data types, excluding classes or structures of constructors, destructors, copy constructors, copy assignment operators, non-static member functions, and virtual functions.
  • Definition of SIMD and SIMT programming kernel functions:
    1. The kernel function is identified by the __global__ and __aicore__ modifier.
    2. The input parameters of the kernel function are used in the same way as those of the SIMD function.
    3. Call the SIMT VF and SIMD VF functions in the SIMD and SIMT programming kernel function.

    1	__global__ __aicore__ void my_kernel(__gm__ float*,...)

• SIMD and SIMT kernel function call mode:
  1. For details about how to call the kernel function, see Kernel Function. The execution configuration is determined by the following parameters:
     • numBlocks: number of cores enabled by the kernel function, which is passed using <<<...>>>.
     • dynUBufSize: size of the dynamic memory. For details about how to allocate dynamic memory, see dynamic memory in UB allocation.
     • stream: The value is of the aclrtStream type, which is used to preserve the order of a stack of asynchronous operations being executed on the device.
  2. Ensure that the size of the dynamic memory used in the kernel function does not exceed the value of dynUBufSize. If the size exceeds the value, out-of-bounds access to the reserved space or data cache occurs, causing undefined behavior.
  3. Maximum configurable dynamic memory size = 256 KB – Reserved space (8 KB) – 32 KB (minimum DCache) – Static memory.

  1	kernel_name<<<numBlocks, dynUBufSize, stream>>>(args...)

Call Hierarchy

• Kernel functions are identified by __global__ __aicore__, which are entry functions on the device and can be called by the host using the <<<...>>> syntax.
• __aicore__ functions are executed on the device. Kernel functions can call __aicore__ functions.
• The SIMD VF function is identified by __simd_vf__ and can be called by kernel functions through asc_vf_call. Only the __simd_callee__ and constexpr functions can be called in the SIMD VF function.
• The SIMT VF function is identified by __simt_vf__ and can be called by kernel functions through asc_vf_call. Only the __simt_callee__ and constexpr functions can be called in the SIMT VF function.
• __simd_callee__ subfunctions can be called in the SIMD VF function. The subfunctions may return values or pass parameters by reference. Such subfunctions are identified by __simd_callee__. Only the __simd_callee__ and constexpr functions can be called in the __simd_callee__ function.
• __simt_callee__ subfunctions can be called in the SIMT VF function. The subfunctions may return values or pass parameters by reference. Such subfunctions are identified by __simt_callee__. Only the __simt_callee__ and constexpr functions can be called in the __simt_callee__ function.

The following figure shows the supported call relationship.

Figure 4 Function call relationship

Programming Sample

 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
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77

__simt_vf__ __launch_bounds__(THREAD_COUNT) inline void simt_gather(
    __gm__ float* input,
    __gm__ uint32_t* index,
    __ubuf__ float* gather_output,
    uint32_t input_total_length,
    uint32_t index_total_length,
    uint32_t output_total_length)
{
    if (threadIdx.x >= output_total_length) {
        return;
    }
    // blockIdx will be supported later.
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx >= index_total_length) {
        return;
    }

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

    gather_output[threadIdx.x] = input[gather_idx];
}

__simd_vf__ inline void simd_adds(__ubuf__ float* output, __ubuf__ float* input,
    uint32_t count, uint32_t one_repeat_size, uint16_t repeat_times)
{
    AscendC::Reg::RegTensor<float> src_reg0;
    AscendC::Reg::RegTensor<float> dst_reg0;
    AscendC::Reg::MaskReg mask_reg; 

    for (uint16_t i = 0; i < repeat_times; i++) {
        mask_reg = AscendC::Reg::UpdateMask<float>(count);
        AscendC::Reg::LoadAlign(src_reg0, input + i * one_repeat_size);
        AscendC::Reg::Adds(dst_reg0, src_reg0, ADDS_ADDEND, mask_reg);
        AscendC::Reg::StoreAlign(output + i * one_repeat_size, dst_reg0, mask_reg);
    }
}

__global__ __aicore__ void gather_and_adds_kernel(__gm__ float* input, __gm__ uint32_t* index, __gm__ float* output,
    uint32_t input_total_length, uint32_t index_total_length)
{
    KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY);
    AscendC::LocalMemAllocator<AscendC::Hardware::UB> ub_allocator;
    // 1. gather numbers from input.
    uint32_t index_total_length_per_block = index_total_length / AscendC::GetBlockNum();
    AscendC::LocalTensor<float> gather_output = ub_allocator.Alloc<float>(index_total_length_per_block);
    asc_vf_call<simt_gather>(dim3(THREAD_COUNT),input, index,
                             (__ubuf__ float *)gather_output.GetPhyAddr(),
                             input_total_length,
                             index_total_length,
                             index_total_length_per_block);

    // 2. use reg compute api to do addition.
    AscendC::LocalTensor<float> adds_output = ub_allocator.Alloc<float>(index_total_length_per_block);
    constexpr uint32_t one_repeat_size = AscendC::GetVecLen() / sizeof(float);
    uint16_t repeat_times = (index_total_length_per_block + one_repeat_size - 1) / one_repeat_size;
    asc_vf_call<simd_adds>((__ubuf__ float *)adds_output.GetPhyAddr(),
        (__ubuf__ float *)gather_output.GetPhyAddr(), index_total_length_per_block, one_repeat_size, repeat_times);

    AscendC::SetFlag<AscendC::HardEvent::V_MTE3>(0);
    AscendC::WaitFlag<AscendC::HardEvent::V_MTE3>(0);

    // 3. copy data to global memory.
    AscendC::GlobalTensor<float> output_global_tensor;
    output_global_tensor.SetGlobalBuffer(output + index_total_length_per_block * AscendC::GetBlockIdx());
    AscendC::DataCopy(output_global_tensor, adds_output, index_total_length_per_block);
}

int main(int argc, char *argv[])
{
     ...
     //numBlocks only supports one dimension currently.
     gather_and_adds_kernel<<<numBlocks, dynUBufSize, stream>>>(input_device, index_device, output_device, input_total_length, index_total_length);
     ...
}