asc_set_l0c_copy_prequant

产品支持情况

• Ascend 950PR/Ascend 950DT：支持

• Atlas A3 训练系列产品/Atlas A3 推理系列产品：不支持

• Atlas A2 训练系列产品/Atlas A2 推理系列产品：不支持

• Atlas 200I/500 A2 推理产品：不支持

• Atlas 推理系列产品AI Core：不支持

• Atlas 推理系列产品Vector Core：不支持

• Atlas 训练系列产品：不支持

功能说明

头文件路径为："c_api/cube_datamove/cube_datamove.h"。

本接口用于在L0C Buffer搬出过程中进行随路Scalar量化时，配置量化计算所需的缩放系数、偏移量和量化结果符号。配置完成后，由后续调用的asc_copy_l0c2gm、asc_copy_l0c2l1或asc_copy_l0c2ub接口使用，并且需要与搬出接口的quant_pre_mode搭配使用。

quant_pre_mode量化模式参数支持的Scalar量化枚举值如下：

• DEQF16：int32_t量化成half， scalar量化。
• REQ4：int32_t量化成int4b_t，scalar量化。
• REQ8：int32_t量化成int8_t/uint8_t，scalar量化。
• QS322BF16_PRE：int32_t量化成bfloat16_t，scalar量化。
• QF322F16_PRE：float量化成half，scalar量化。
• QF322BF16_PRE：float量化成bfloat16_t，scalar量化。
• QF322S4_PRE：float量化成int4b_t，scalar量化。
• QF322B8_PRE：float量化成int8_t/uint8_t，scalar量化。
• QF322FP8_PRE：float量化成fp8_e4m3fn_t，scalar量化。
• QF322HIF8_PRE：float量化成hifloat8_t(Half to Away Round)，scalar量化。
• QF322HIF8_PRE_HYBRID：float量化成hifloat8_t(Hybrid Round)，scalar量化。
• QF322F32_PRE：float量化成float，scalar量化，精度可以达到双千分之一，无法达到双万分之一。

本接口为矩阵搬出相关配置接口，仅在AIC上生效。

函数原型

C++

__aicore__ inline void asc_set_l0c_copy_prequant(float scale, uint16_t offset, bool is_signed);

__aicore__ inline void asc_set_l0c_copy_prequant(uint64_t config);

参数说明

表1 参数说明

参数名	输入/输出	描述
scale	输入	随路scalar量化所使用的缩放系数。取值不能为inf或nan。
offset	输入	随路量化计算所需的偏移量，取值范围为[0, 511]。不使用偏移量时，配置为0。在REQ8/QF322B8_PRE/REQ4/QF322S4_PRE模式下生效。
is_signed	输入	量化结果是否为有符号数据。取值为true时，量化结果为int8_t；取值为false时，量化结果为uint8_t。在REQ8/QF322B8_PRE模式下生效。
config	输入	Scalar量化参数，由scale，offset，is_signed拼装而成，config各比特位含义参见表2。

表2 config比特位含义

比特位数	变量名	描述
0~12	-	无效比特位。
13~31	scale	数据类型视为float，硬件以(1, 8, 10)格式进行计算。其中第31位为符号位，第23~30位为指数位，第13~22位为尾数位。取值不能为inf或nan。
32~36	-	无效比特位。
37~45	offset	9位整型数据。源数据乘以量化系数后的结果与Offset相加。不使用偏移量时，配置为0。在REQ8/QF322B8_PRE/REQ4/QF322S4_PRE模式下生效。
46	is_signed	取值为1时，量化结果为int8_t；取值为0时，量化结果为uint8_t。在REQ8/QF322B8_PRE模式下生效。
47~63	-	无效比特位。

当前提供了REQ8量化算法Python示例代码，请参考(V)REQ8量化算法。

返回值说明

无

流水类型

PIPE_S

约束说明

• 本接口非AIC调用直接返回。
• 本接口需在对应的L0C Buffer搬出接口执行前调用，并在搬出接口中配置对应的Scalar量化模式。

调用示例

将代码保存为example.asc后，可通过bisheng命令编译运行，其中--npu-arch参数需根据实际产品型号指定对应的NPU架构，具体产品与NPU架构的映射关系请参考__NPU_ARCH__。

以Ascend 950PR/Ascend 950DT产品（对应NPU架构为dav-3510）为例，编译运行命令如下：

Bash

bisheng example.asc -o main --npu-arch=dav-3510 && ./main

以下调用示例代码仅Ascend 950PR/Ascend 950DT产品支持。

C++

#include <cstdint>
#include <iostream>
#include <vector>
#include "c_api/asc_simd.h"
#include "acl/acl.h"

namespace {
constexpr uint32_t DIM = 16;
constexpr uint32_t ELEMENTS = DIM * DIM;
constexpr uint16_t HALF_ONE = 0x3c00;
constexpr float SCALE = 2.0f;
constexpr uint16_t OFFSET = 3;

__global__ __cube__ void asc_set_l0c_copy_prequant_kernel(
    __gm__ uint16_t* a, __gm__ uint16_t* b, __gm__ int8_t* output)
{
    asc_init();
    __cbuf__ half a_l1[ELEMENTS], b_l1[ELEMENTS];
    __ca__ half a_l0[ELEMENTS];
    __cb__ half b_l0[ELEMENTS];
    __cc__ float c_l0[ELEMENTS];

    asc_set_gm2l1_nz_para(1, 1, 16, 0);
    asc_copy_gm2l1_nd2nz(a_l1, reinterpret_cast<__gm__ half*>(a), DIM * sizeof(half), 0, DIM, DIM, 0, false);
    asc_set_gm2l1_nz_para(1, 1, 16, 0);
    asc_copy_gm2l1_nd2nz(b_l1, reinterpret_cast<__gm__ half*>(b), DIM * sizeof(half), 0, DIM, DIM, 0, false);
    asc_sync_notify(PIPE_MTE2, PIPE_MTE1, EVENT_ID0);
    asc_sync_wait(PIPE_MTE2, PIPE_MTE1, EVENT_ID0);
    asc_copy_l12l0a(a_l0, a_l1, 0, 0, 1, 1, 1, 1);
    asc_copy_l12l0b_transpose(b_l0, b_l1, 0, 0, 1, 1, 1, 1);
    asc_sync_notify(PIPE_MTE1, PIPE_M, EVENT_ID0);
    asc_sync_wait(PIPE_MTE1, PIPE_M, EVENT_ID0);
    asc_mmad(c_l0, a_l0, b_l0, DIM, DIM, DIM, 0, false, false, true);
    asc_sync_notify(PIPE_M, PIPE_FIX, EVENT_ID0);
    asc_sync_wait(PIPE_M, PIPE_FIX, EVENT_ID0);

    asc_set_l0c_copy_prequant(SCALE, OFFSET, true);
    asc_set_l0c_copy_nz_para(1, 0, 0);
    asc_copy_l0c2gm(output, c_l0, DIM, DIM, DIM, DIM, 0, 0, 0,
        static_cast<uint64_t>(QuantMode_t::QF322B8_PRE), 0, false, true,
        static_cast<uint64_t>(QuantMode_post::NoConv), 0, false, 0, false, false, false, false);
    asc_sync_pipe(PIPE_ALL);
}

void print_row(const char* label, const std::vector<int8_t>& data)
{
    std::cout << label << ':';
    for (uint32_t i = 0; i < 8; ++i) std::cout << ' ' << static_cast<int32_t>(data[i]);
    std::cout << " ..." << std::endl;
}
} // namespace

int main()
{
    std::vector<uint16_t> a(ELEMENTS), b(ELEMENTS);
    std::vector<int8_t> output(ELEMENTS), golden(ELEMENTS, static_cast<int8_t>(OFFSET));
    for (uint32_t i = 0; i < DIM; ++i) {
        a[i * DIM + i] = HALF_ONE;
        b[i * DIM + i] = HALF_ONE;
        golden[i * DIM + i] = static_cast<int8_t>(SCALE + OFFSET);
    }

    aclInit(nullptr);
    aclrtSetDevice(0);
    uint16_t *a_device = nullptr, *b_device = nullptr;
    int8_t* output_device = nullptr;
    aclrtMalloc(reinterpret_cast<void**>(&a_device), a.size() * sizeof(uint16_t), ACL_MEM_MALLOC_HUGE_FIRST);
    aclrtMalloc(reinterpret_cast<void**>(&b_device), b.size() * sizeof(uint16_t), ACL_MEM_MALLOC_HUGE_FIRST);
    aclrtMalloc(reinterpret_cast<void**>(&output_device), output.size() * sizeof(int8_t), ACL_MEM_MALLOC_HUGE_FIRST);
    aclrtMemcpy(a_device, a.size() * sizeof(uint16_t), a.data(), a.size() * sizeof(uint16_t), ACL_MEMCPY_HOST_TO_DEVICE);
    aclrtMemcpy(b_device, b.size() * sizeof(uint16_t), b.data(), b.size() * sizeof(uint16_t), ACL_MEMCPY_HOST_TO_DEVICE);
    asc_set_l0c_copy_prequant_kernel<<<1, 0>>>(a_device, b_device, output_device);
    aclrtSynchronizeDevice();
    aclrtMemcpy(output.data(), output.size() * sizeof(int8_t), output_device, output.size() * sizeof(int8_t),
        ACL_MEMCPY_DEVICE_TO_HOST);
    print_row("Quantized output row 0", output);
    print_row("Golden row 0", golden);
    const bool passed = output == golden;
    std::cout << (passed ? "[Success] asc_set_l0c_copy_prequant scalar quantization passed."
                         : "[Failed] asc_set_l0c_copy_prequant result mismatch.") << std::endl;
    aclrtFree(a_device); aclrtFree(b_device); aclrtFree(output_device);
    aclrtResetDevice(0);
    aclFinalize();
    return passed ? 0 : 1;
}