Programming Sample

Currently, SIMT APIs are not supported in SIMT programming scenarios.

For details about the built-in keywords and APIs required for SIMT-based operator development, see SIMT Built-In Keywords and C APIs at the SIMT Language Extension Layer. Currently, SIMT programming does not support some syntax structures. For details about relevant restrictions, see C/C++ Syntax Restrictions.

Considering the following computation scenario: Obtain 12288 rows of data with a specified index from a 2D vector with the shape of 100000 × 128. The formula for computing data in the ith row in the operator output is as follows:

output[i] = input[index[i]]

Implement the computation logic for a row of data in the kernel function, and configure multiple threads to compute data in different rows. The implementation logic of the kernel function is as follows:

• Find the offset of the data to be computed by the current thread based on the unique index of each thread.

  1	int32_t out_row = blockIdx.x * blockDim.x + threadIdx.x;

  A thread completes a kernel function computation operation. The index offset is computed using blockIdx.x × blockDim.x + threadIdx.x within the kernel function. blockIdx indicates the index of the current thread block, blockDim indicates the number of threads enabled in a thread block, and threadIdx indicates the index of the current thread within the thread block. For more details, see SIMT Built-In Keywords.

• The input data at the offset position is copied to the output using the subscript offset, thereby obtaining the specified data.

  1 2 3 4 5 6 7

  uint32_t in_row = index[out_row]; for (int32_t col = 0; col < in_width; col++) { // Each thread completes the computation of a row of data. int input_idx = in_row * in_width + col; int output_idx = out_row * in_width + col; gather_output[output_idx] = input[input_idx]; }

Refer to the following code for kernel function implementation. For details about the complete sample, see pure_simt_gather operator sample.

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14

template <typename type_data, typename type_idx>
__global__ void gather_custom(type_data* input, type_idx* index, type_data* gather_output, uint32_t in_width, uint32_t index_total_length)
{
    // Compute the index offset.
    int32_t out_row = blockIdx.x * blockDim.x + threadIdx.x;
    // Obtain the row index to be processed from the index.
    uint32_t in_row = index[out_row];
    // Process all data in the row cyclically.
    for (int32_t col = 0; col < in_width; col++) {
        int input_idx = in_row * in_width + col;
        int output_idx = out_row * in_width + col;
        gather_output[output_idx] = input[input_idx]; // Copy the input data to the output.
    }
}

1
2
3
4
5
6

int main(int argc, char* argv[])
{
   ...    
    gather_custom<<<48, 256, 0, stream>>>(input_device, index_device, output_device, in_shape[1], index_total_length);
   ... 
}