Fusion Scope Marking Example

By calling the aclskScopeBegin and aclskScopeEnd APIs, if a valid scope name is passed, the operator can be fused into SuperKernel. If nullptr is passed, the operator will not be fused. This function takes effect only when SuperKernel optimization is enabled. The following code demonstrates how to use the SuperKernel technology to implement a simple addition operator on the NPU.

Example of key steps for fusing operators into the SuperKernel

The code only shows the core logic and cannot be directly compiled or run. It is for reference only.

        
         
           
           
             #include "acl/acl.h"
#include "super_kernel/super_kernel.h"
...
// Original kernel function (add_custom)
__global__ __vector__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z, AddCustomTilingData tiling)
{
    KernelAdd op;
    op.Init(x, y, z, tiling.totalLength, tiling.tileNum);
    op.Process();
}
// Define a parameter structure to encapsulate the parameters to be passed to SuperKernel.
struct ArgsStruct {
    GM_ADDR x;                     
    GM_ADDR y;                     
    GM_ADDR z;  
    AddCustomTilingData tiling;
};
// Define a SuperKernel subfunction with template parameters (add_custom_sk).
template<uint32_t splitNum>
__sk__ __vector__ void add_custom_sk(const ArgsStruct *args, sk::SkSystemArgs *sysArgs)
{
    // Obtain parameters from the structure.
    GM_ADDR x = args->x;
    GM_ADDR y = args->y;
    GM_ADDR z = args->z;
    AddCustomTilingData tiling = args->tiling;
    uint16_t blockNum = sysArgs->skNumBlocks;
    uint16_t blockIdx = sysArgs->skBlockIdx;
    // The logic is the same as that of the global function of the original kernel.
    KernelAdd op;
    op.Init(x, y, z, tiling.totalLength, tiling.tileNum);
    op.Process();
}

// Use SK_BIND to bind the add_custom kernel to multiple SK sub-functions, and instantiate four different symbols by specifying template parameters.
SK_BIND(add_custom, 4, add_custom_sk<0>, add_custom_sk<1>, add_custom_sk<2>, add_custom_sk<3>);

// Main function
int main()
{
    // Operator data preparation
    constexpr uint32_t totalLength = 8 * 2048;
    constexpr float valueX = 1.2f;
    constexpr float valueY = 2.3f;
    std::vector<float> x(totalLength, valueX);
    std::vector<float> y(totalLength, valueY);
    constexpr uint32_t numBlocks = 8;
    uint32_t totalLength = x.size();
    size_t totalByteSize = totalLength * sizeof(float);
    int32_t deviceId = 0;
    aclrtStream stream = nullptr;
    AddCustomTilingData tiling = {totalLength, 8};
    uint8_t *xHost = reinterpret_cast<uint8_t *>(x.data());
    // printf("0-host: x0:%d \n", xHost[0]);
    uint8_t *yHost = reinterpret_cast<uint8_t *>(y.data());
    uint8_t *zHost = nullptr;
    uint8_t *xDevice = nullptr;
    uint8_t *yDevice = nullptr;
    uint8_t *zDevice = nullptr;
    // Allocate memory.
    aclrtMallocHost((void **)(&zHost), totalByteSize);
    aclrtMalloc((void **)&xDevice, totalByteSize, ACL_MEM_MALLOC_HUGE_FIRST);
    aclrtMalloc((void **)&yDevice, totalByteSize, ACL_MEM_MALLOC_HUGE_FIRST);
    aclrtMalloc((void **)&zDevice, totalByteSize, ACL_MEM_MALLOC_HUGE_FIRST);
    // Transfer data.
    aclrtMemcpy(xDevice, totalByteSize, xHost, totalByteSize, ACL_MEMCPY_HOST_TO_DEVICE);
    aclrtMemcpy(yDevice, totalByteSize, yHost, totalByteSize, ACL_MEMCPY_HOST_TO_DEVICE);
    // Perform initialization.
    aclInit(nullptr);
    aclrtSetDevice(deviceId);
    aclrtCreateStream(&stream);
    // Build a graph.
    aclmdlRICaptureBegin(stream, ACL_MODEL_RI_CAPTURE_MODE_GLOBAL);
    aclmdlRI modelRI;
    // Mark the start position of SuperKernel fusion.
    aclskScopeBegin("sk1", stream);
    add_custom<<<numBlocks, nullptr, stream>>>(xDevice, yDevice, zDevice, tiling);
    // Mark the end position of SuperKernel fusion.
    aclskScopeEnd("sk1", stream);
    aclmdlRICaptureEnd(stream, &modelRI);
    // Enable SuperKernel optimization.
    aclskOptimize(modelRI, nullptr);
    // Execute the graph.
    aclmdlRIExecuteAsync(modelRI, stream);
    // Obtain the graph execution result.
    aclrtSynchronizeStream(stream);
    aclrtMemcpy(zHost, totalByteSize, zDevice, totalByteSize, ACL_MEMCPY_DEVICE_TO_HOST);
    std::vector<float> z((float *)zHost, (float *)(zHost + totalLength));
    // Release resources.
    aclrtFree(xDevice);
    aclrtFree(yDevice);
    aclrtFree(zDevice);
    aclrtFreeHost(zHost);
    aclrtDestroyStream(stream);
    aclrtResetDevice(deviceId);
    aclFinalize();
}

            

          

        
       

For details about the preceding ACL APIs, see Runtime APIs.

Example of key steps for not fusing operators into the SuperKernel

The following code only shows the core logic and cannot be directly compiled or run. It is for reference only.

        
         
           
           
             #include "acl/acl.h"
#include "super_kernel/super_kernel.h"
...
int main()
{
    // Operator data preparation
    ...
    // Perform initialization.
    aclInit(nullptr);
    aclrtSetDevice(deviceId);
    aclrtCreateStream(&stream);
    // Build a graph.
    aclmdlRICaptureBegin(stream, ACL_MODEL_RI_CAPTURE_MODE_GLOBAL);
    aclmdlRI modelRI;
    // Mark the start position of SuperKernel fusion. If nullptr is passed, fusion is not performed.
    aclskScopeBegin(nullptr, stream);
    add_custom<<<numBlocks, nullptr, stream>>>(xDevice, yDevice, zDevice, tiling);
    // Mark the end position of SuperKernel fusion. If nullptr is passed, fusion is not performed.
    aclskScopeEnd(nullptr, stream);
    aclmdlRICaptureEnd(stream, &modelRI);
    // Enable SuperKernel optimization.
    aclskOptimize(modelRI, nullptr);
    // Execute the graph.
    aclmdlRIExecuteAsync(modelRI, stream);
    // Obtain the graph execution result.
    ...
    // Release resources.
    aclrtFree(xDevice);
    aclrtFree(yDevice);
    aclrtFree(zDevice);
    aclrtFreeHost(zHost);
    aclrtDestroyStream(stream);
    aclrtResetDevice(deviceId);
    aclFinalize();
}

            

          

        
       

For details about the preceding ACL APIs, see Runtime APIs.

Parent topic: SuperKernel Fusion Scope Marking