CANN/asc-devkit函数语法限制

函数

【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言，原生支持C和C++标准规范，主要由类库和语言扩展层构成，提供多层级API，满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit

核函数参数列表限制

核函数的参数列表（argument list）仅支持以下类型：

• 基础数据类型，如int32_t、float等。
• 基础数据类型组成的结构体，支持嵌套结构体形式，但必须为POD（Plain Old Data）类型。
• 基础数据类型的指针类型，如int32_t*、float*等，实际上这些指针指向的是Global Memory内存。

SIMT VF函数限制

__simt_vf__标记的SIMT VF函数需要遵循以下约束和限制：

• 标量参数不允许使用指针和引用。

• 不允许将栈数组作为参数传入SIMT VF函数；SIMT VF函数的指针形参必须显式使用__gm__或__ubuf__地址空间限定符。以下错误示例会在编译期报错。

  __simt_vf__ __launch_bounds__(1024) inline void foo(__gm__ int* a, __gm__ int* b, __gm__ int* c, int* array) { int idx = blockIdx.x * blockDim.x + threadIdx.x; a[idx] = b[idx] + c[idx] + array[0]; } __global__ __aicore__ void foo(__gm__ int* a, __gm__ int* b, __gm__ int* c) { int array[5] = {0,1,2,3,4}; asc_vf_call<foo>(dim3{256}, a, b, c, array); // error: cannot initialize a parameter of type '__gm__ int *' with an lvalue of type 'int[5]' }

• 如果传参中出现多级指针，不允许使用内层栈地址指针访问，以下为错误示例。

  __simt_vf__ __launch_bounds__(1024) inline void foo(__gm__ int* a, __gm__ int* b, __gm__ int* c, __ubuf__ uint64_t* s) { int idx = blockIdx.x * blockDim.x + threadIdx.x; int* stack = (int*)(s[0]); // error: *stack表示从多级指针中读取，不允许使用 a[idx] = b[idx] + c[idx] + *stack; } __global__ __aicore__ void foo(__gm__ int* a, __gm__ int* b, __gm__ int* c) { int stack = 0; __ubuf__ uint64_t* s = ...; s[0] = &stack; asc_vf_call<foo>(dim3{256}, a, b, c, s); }

• 不支持通过函数指针进行间接调用，被调用的__simt_vf__函数需要在编译期确定。

• 函数的inline行为由编译器决定，添加的always_inline或noinline将被忽略。

• 不允许使用结构体作为参数，以下为错误示例。

  __simt_vf__ __launch_bounds__(1024) inline void foo(__gm__ int* a, __gm__ int* b, __gm__ int* c, struct S s) { // error: s表示结构体参数，不允许使用 }

SIMD与SIMT混合编程场景结构体使用限制

SIMD与SIMT混合编程场景下，当使用位于Global Memory或Unified Buffer中的结构体或类拷贝构造新结构体或类时，需要开发者为结构体或类手动实现拷贝构造函数，编译器不会默认隐式生成。

struct TestS { int a; float b; // 手动实现Global Memory内存空间的拷贝构造函数 __callee__ TestS(const __gm__ TestS& other) { a = other.a; b = other.b; } // 手动实现Unified Buffer内存空间的拷贝构造函数 __callee__ TestS(const __ubuf__ TestS& other) { a = other.a; b = other.b; } }; __simt_vf__ void simt_kernel(__gm__ TestS *data1, __ubuf__ TestS *data2) { uint32_t idx = threadIdx.x; TestS t1 = data1[idx]; // 调用Global Memory拷贝构造函数 TestS t2 = data2[idx]; // 调用Unified Buffer拷贝构造函数 ... }

在异构编译场景中，Host和Device可共用结构体定义代码，但Host侧无法识别__callee__、__gm__等Device侧关键字，因此无法根据这些关键字实现函数重载。当需要定义多个地址空间重载函数时，需通过宏__NPU_ARCH__将相关Device代码隔离，以避免函数定义冲突。__NPU_ARCH__是由编译预定义的宏，可用于区分Device侧代码，具体可参考《毕昇编译器》中的“基本编程指导 > AI Core编程指导 > 预定义宏和内建变量”章节。

此外，SIMD与SIMT混合编程场景下不支持直接向Global Memory或Unified Buffer地址空间写入结构体。

struct TestS { int a; float b; //device 侧代码隔离 #ifdef __NPU_ARCH__ __callee__ TestS(const __gm__ TestS& other) { a = other.a; b = other.b; } __callee__ TestS(const __ubuf__ TestS& other) { a = other.a; b = other.b; } #endif // host 侧代码 TestS(const TestS& other) { a = other.a; b = other.b; } }; __simt_vf__ void simt_kernel(__gm__ TestS *data1, __ubuf__ TestS *data2) { uint32_t idx = threadIdx.x; TestS t1 = data1[idx]; // 调用Global Memory拷贝构造函数 data2[0] = t; // 不支持，编译报错 ... } __global__ __vector__ void kernel_func(__gm__ TestS *data1, __gm__ TestS *data2) { TestS t = data1[0]; // 调用Global Memory拷贝构造函数 data2[0] = t; // 不支持，编译报错 ... }

【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言，原生支持C和C++标准规范，主要由类库和语言扩展层构成，提供多层级API，满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit