CANN/asc-devkit函数语法限制 函数【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit核函数参数列表限制核函数的参数列表argument list仅支持以下类型基础数据类型如int32_t、float等。基础数据类型组成的结构体支持嵌套结构体形式但必须为PODPlain 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_callfoo(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_callfoo(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创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考