asc_vf_call

Applicability

Product	Supported
Atlas 350 Accelerator Card	√
Atlas A3 training product/Atlas A3 inference product	x
Atlas A2 training product/Atlas A2 inference product	x
Atlas 200I/500 A2 inference product	x
Atlas inference product AI Core	x
Atlas inference product Vector Core	x
Atlas training product	x

Function Usage

Starts a SIMT VF subtask to launch a specified number of threads for executing the specified SIMT kernel function by configuring parameters in SIMD and SIMT programming.

When using asc_vf_call to start SIMT VF subtasks, the subtask function must not be a class member function. You must use either an ordinary function or a class static function, and the entry function must be decorated with __simt_vf__.

When using asc_vf_call to start SIMT VF subtasks, only raw pointers and common basic data types can be passed as parameters. Structures and arrays cannot be passed.

Prototype

template <auto funcPtr, typename... Args>
__aicore__ inline void asc_vf_call(dim3 threadNums, Args &&...args)

Parameters

**Table 1** Template parameters
Parameter	Description
funcPtr	Specifies the SIMT entry kernel function.
Args	Defines variable parameters to pass parameters to the SIMT entry kernel function.

**Table 2** Parameters
Parameter	Input/Output	Description
threadNums	Input	A dim3 structure, defined as {dimx, dimy, dimz}, which specifies the number of threads in a SIMT thread block. The total number of threads is dimx * dimy * dimz, and the value must not exceed 2048.
args	Input	Defines variable parameters to pass parameters to the SIMT entry kernel function.

Returns

None

Restrictions

For SIMT programming, this API is not supported.

Header File to Be Included

To use this API, the simt_api/common_functions.h header file must be included.

#include "simt_api/common_functions.h"

Examples

For SIMD and SIMT programming:

Perform addition calculation on Global Memory data.

__simt_vf__ __launch_bounds__(2048) inline void SimtCompute(
    __gm__ float* dst, __gm__ float* src0, __gm__ float* src1, int count) const
{
    // SIMT code
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    dst[idx] = src0[idx] + src1[idx];
}

__global__ __vector__ void SimtComputeShell(__gm__ float* x, __gm__ float* y, __gm__ float* z, const int size)
{
    __gm__ float* dst = x;
    __gm__ float* src0 = y;
    __gm__ float* src1 = z;
    // Start the SIMT VF subtask by calling asc_vf_call.
    asc_vf_call<SimtCompute>(dim3{1024, 1, 1}, dst, src0, src1, size);
}

Parent topic: Kernel Function Definition