ICPU_RUN_KF

Product Support

Product	Supported
Atlas A3 training products/Atlas A3 inference products	√
Atlas A2 training products/Atlas A2 inference products	√
Atlas 200I/500 A2 inference products	√
Atlas inference product's AI Core	√
Atlas inference product's Vector Core	x
Atlas training products	√

Function

Functions as the CPU debugging entry and completes calls to CPU operator programs during verification of the CPU-side operation of the kernel function.

Prototype

1	#define ICPU_RUN_KF(func, blkdim, ...)

Parameters

Parameter	Input/Output	Description
func	Input	Pointer to the operator kernel function.
blkdim	Input	Number of operator cores (corenum).
...	Input	All input and output parameters entered in sequence. The maximum number of parameters is 32. If the number of parameters exceeds 32, a compilation error occurs.

Returns

None

Restrictions

Except func and blkdim, other variables must be pointers to the shared memory allocated by GmAlloc. The number and sequence of input parameters must be the same as those of the kernel function.

Example

The following code uses the add_custom operator as an example to describe how to compile the application called by the operator when the operator kernel function is verified on the CPU. When implementing your own applications, pay attention to the modifications caused by different operator kernel functions, including different operator kernel function names and input and output parameters. Properly allocate memory, copy memory, and read/write files. You can directly reuse the calling methods of related APIs.

Include header files as required. The ASCENDC_CPU_DEBUG macro is used to distinguish the header files that need to be included on the CPU side from those on the NPU side.

#include "data_utils.h"
#ifndef ASCENDC_CPU_DEBUG
#include "acl/acl.h"
#else
#include "tikicpulib.h"
extern "C" __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z); // Kernel function declaration.
#endif

Perform runtime verification on the CPU. To implement the runtime verification of the operator kernel function on the CPU, perform the following steps.

Figure 1 Runtime verification on the CPU

int32_t main(int32_t argc, char* argv[])
{
    uint32_t blockDim = 8;
    size_t inputByteSize = 8 * 2048 * sizeof(uint16_t);
    size_t outputByteSize = 8 * 2048 * sizeof(uint16_t);

    // Allocate shared memory and initialize data using GmAlloc.
    uint8_t* x = (uint8_t*)AscendC::GmAlloc(inputByteSize);
    uint8_t* y = (uint8_t*)AscendC::GmAlloc(inputByteSize);
    uint8_t* z = (uint8_t*)AscendC::GmAlloc(outputByteSize);

    ReadFile("./input/input_x.bin", inputByteSize, x, inputByteSize);
    ReadFile("./input/input_y.bin", inputByteSize, y, inputByteSize);
    // Set the kernel mode to AIV for vector operators.
    AscendC::SetKernelMode(KernelMode::AIV_MODE);
    // Call the ICPU_RUN_KF debugging macro to call the kernel function on the CPU.
    ICPU_RUN_KF(add_custom, blockDim, x, y, z);
    // Write the output data.
    WriteFile("./output/output_z.bin", z, outputByteSize);
    // Release the allocated resources using GmFree.
    AscendC::GmFree((void *)x);
    AscendC::GmFree((void *)y);
    AscendC::GmFree((void *)z);
    return 0;
}

Parent topic: CPU Twin Debugging