SIMT Built-In Keywords

Function Execution Space Qualifier

The function execution space qualifier indicates whether the function is executed on the host or device, and the space scope that can be called.

A function modified by __global__ is the entry of a kernel function and has the following restrictions:

• The return type of the function must be void. It cannot be a member function of a class, struct, or union.
• Recursive calls are not supported.
• The __global__ function is called asynchronously. After the function is called, the host thread is returned.
• It can be called only by functions on the host and executed on the device.

Functions decorated with __aicore__ can be executed only on the device and can be called only by __global__ functions or other __aicore__ functions.

Functions decorated with __host__ can be called and executed only on the host.

Memory Space Qualifier

The memory space qualifier __ubuf__ is used to indicate dynamic and static memory. The size of the static memory is determined during compilation, and the size of the dynamic memory is determined during kernel function execution.

The current version does not support dynamic and static memory.

• The static memory is allocated through arrays.

  1	__ubuf__ half staticBuf[1024];

• The dynamic memory is allocated using the following method:

  1	extern __ubuf__ half dynamicBuf[];

  The actual size of the dynamic memory needs to be configured when the kernel function is started. For details, see Kernel Function Configuration.

Built-in Variables

• dim3

  Specifies and obtains the built-in structure of the thread grid and thread block in the x, y, and z dimensions.

  dim3 consists of three unsigned integers. The structure is defined as {dimx, dimy, dimz}, which is used to specify the sizes of three different dimensions. The total number of three dimensions is dimx × dimy × dimz. You can create a dim3 structure as follows:

  1 2 3

  dim3(x); // Create a one-dimensional structure. The default value 1 is used for dimy and dimz. dim3(x, y); // Create a two-dimensional structure. The default value 1 is used for dimz. dim3(x, y, z); // Create a three-dimensional structure.

Currently, the following built-in variables of the dim3 structure are provided, which are available only on the device:

• gridDim

  Built-in global variable, which can be used only in kernel functions. It indicates the number of thread blocks in each dimension of the entire compute task.

  • gridDim.x indicates the number of thread blocks in the x dimension.
  • gridDim.y indicates the number of thread blocks in the y dimension. Currently, only 1 can be returned.
  • gridDim.z indicates the number of thread blocks in the z dimension. Currently, only 1 can be returned.
• blockDim

  Built-in global variable, which can be directly used in the kernel function to obtain the three-dimensional hierarchy of threads configured in the thread block, that is, the value of the dim3 structure instance configured when the VF is started. blockDim.x, blockDim.y, and blockDim.z indicate the numbers of threads in the three dimensions of a thread block, respectively.

• blockIdx

  Built-in global variable, which can be used only in the kernel function to obtain the block index. It indicates the coordinates of the thread block where the current thread is located in the entire grid.

  • The value of blockIdx.x ranges from 0 to gridDim.x – 1.
  • The value of blockIdx.y ranges from 0 to gridDim.y – 1. Currently, only 0 can be returned.
  • The value of blockIdx.z ranges from 0 to gridDim.z – 1. Currently, only 0 can be returned.
• threadIdx

  Built-in global variable, which can be directly used in the kernel function to obtain the index of the current thread in the thread block. threadIdx.x, threadIdx.y, and threadIdx.z indicate the indexes of the current thread in three dimensions, respectively. The value range is [0, blockDim.x) for threadIdx.x, [0, blockDim.y) for threadIdx.y, and [0, blockDim.z) for threadIdx.z. The mapping between the thread index and thread ID in a thread block is as follows:

  • For a one-dimensional thread block, the thread ID is threadIdx.x.
  • For a two-dimensional thread block, the thread ID is (threadIdx.x + threadIdx.y × blockDim.x).
  • For a three-dimensional thread block, the thread ID is (threadIdx.x + threadIdx.y × blockDim.x + threadIdx.z × blockDim.x × blockDim.y).

Currently, the following built-in variables of the int type are provided, which are available only on the device:

• warpSize

  Runtime variable, indicating the number of threads in a warp. The value is fixed at 32.

Built-in Data Types

Operator

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17
18
19
20
21
22
23

// Addition operation
res[idx] = x[idx] + y[idx]; 

// Inversion operation
x[idx] = (-x[idx]);

// Comparison operation
if (x[idx] > y[idx]) {
    res[idx] = x[idx];
} else {
    res[idx] = y[idx];
}

// Bitwise AND operation
res[idx] = x[idx] & y[idx];

// Logical OR operation
if (x[idx] || y[idx]) {
    res[idx] = 1;
}

// Condition operation
res[idx] = x[idx] > y[idx] ? x[idx] : y[idx];

Kernel Function Configuration

The execution configuration must be specified when a function modified by the __global__ qualifier is called. The execution configuration is inserted between the parameter list enclosed in parentheses following the function name, for example:

1

<<<grid_dim, block_dim, dynamic_mem_size, stream>>>

Where:

• grid_dim: int or dim3 type, which is used to specify the dimension and size of a grid. The result of grid_dim.x × grid_dim.y × grid_dim.z is equal to the total number of started thread blocks.
• block_dim: int or dim3 type, which is used to specify the dimension and size of each thread block. The result of block_dim.x × block_dim.y × block_dim.z is equal to the number of threads in each thread block.
• dynamic_mem_size: size_t type. This parameter specifies the number of bytes of shared memory dynamically allocated to each thread block in this call, except the statically allocated memory.
• stream: aclrtStream type pointer, which specifies the associated stream and is used to maintain the execution sequence of asynchronous operations.

1
2
3
4

// Declaration
__global__ void add_custom(float* x, float* y, float* z, uint64_t total_length);
// Call
add_custom<<<block_num, thread_num_per_block, dyn_ubuf_size, stream>>>(x, y, z, 1024);

Before executing the function, the preceding configuration parameters are validated. If grid_dim or block_dim exceeds the maximum size allowed by the device, or dynamic_smem_bytes exceeds the remaining available shared memory after static memory allocation, the function will fail to execute.

When multiple threads are executed concurrently, using fewer registers per thread allows more threads and thread blocks to reside in the Ascend AI Processor, thereby improving performance. Therefore, the compiler uses heuristic algorithms to minimize register spilling and the number of instructions, while also reducing the number of used registers. Applications can use the __launch_bounds__() qualifier in the __global__ function definition to restrict the launch bounds, providing additional information to assist the compiler in optimizing this process. This is an optional configuration.

• __launch_bounds__(N) is a function marking macro that can be optionally configured on the SIMT VF entry function. It is used to specify the maximum number of threads that can be started by the SIMT VF during compilation. If __launch_bounds__ is not set, the maximum number of threads is 1024 by default. The parameter N must meet the following requirements:
  • N ≥ dimx × dimy × dimz; dimx, dimy, and dimz are the dim3 structure that represents threads.
  • The value of N ranges from 1 to 2048.

    The maximum number of threads determines the number of registers that can be allocated to each thread. The following table lists the mapping between the maximum number of threads and the number of available registers for each thread. Registers are used to store local variables in threads. If the number of local variables exceeds the number of registers, problems such as stack overflow may occur. It is recommended that the maximum number of threads be the same as the number of dim3 threads for starting the VF task.

    Table 5 Mapping between the number of threads in __launch_bounds__ and the number of available registers for each thread

    Number of Threads	Number of Available Registers for Each Thread
    1025–2048	16
    513–1024	32
    257–512	64
    1–256	127

    The following is an example of setting the maximum number of threads for the SIMT function to 512:

    1	__simt_vf__ __launch_bounds__(512) inline void add(__gm__ uint8_t* x, __gm__ uint8_t* y, __gm__ uint8_t* z)