已关闭
[Requirement|需求建议]: acl调用代码相比Cuda C更复杂,多约 70%代码量,能否做优化 #580
gao_dafa创建于  4月28日关闭于  5月18日
gao_dafa成员
4月28日 创建

Thanks for sending an requirement! Please fill in the following template to help quickly solve your problem.

Backgroud(背景信息)

CUDA 最简流程

// 无需显式 init,第一次调用 CUDA API 时自动初始化
cudaSetDevice(0); // 1

float* d; cudaMalloc(&d, size); // 2
cudaMemcpy(d, host_data, size, cudaMemcpyHostToDevice); // 3
kernel<<<grid, block>>>(d); // 4
cudaDeviceSynchronize(); // 5
cudaMemcpy(host_out, d, size, cudaMemcpyDeviceToHost); // 6
cudaFree(d); // 7

// 无需显式 finalize,进程退出自动清理

7 步,无需 init/finalize,stream 可选。

AscendC 最简流程

aclInit(nullptr); // 1 ← 多的
aclrtSetDevice(0); // 2
aclrtStream stream = nullptr;
aclrtCreateStream(&stream); // 3 ←
多的(必须显式创建)

float* d; aclrtMalloc((void**)&d, size, ACL_MEM_MALLOC_HUGE_FIRST); //
4(多一个策略参数)
aclrtMemcpy(d, size, host, size, ACL_MEMCPY_HOST_TO_DEVICE); //
5(多一个 dstSize)
kernel<<<grid, block, 0, stream>>>(d); // 6
aclrtSynchronizeStream(stream); // 7
aclrtMemcpy(host_out, size, d, size, ACL_MEMCPY_DEVICE_TO_HOST); // 8
aclrtFree(d); // 9

aclrtDestroyStream(stream); // 10 ← 多的
aclrtResetDevice(0); // 11 ← 多的
aclFinalize(); // 12 ← 多的

12 步,多了 5 个必须的调用。

逐项对比

操作 CUDA AscendC 差异
全局初始化 隐式(自动) aclInit(nullptr) 必须 AscendC 多一步
选设备 cudaSetDevice(0) aclrtSetDevice(0) 相同
创建 stream 可选(默认 stream 自动可用) aclrtCreateStream 必须 AscendC 多一步
分配内存 cudaMalloc(&p, size) 2参数 aclrtMalloc(&p, size, strategy) 3参数 AscendC 多一个策略参数
拷贝内存 cudaMemcpy(d,s,n,kind) 4参数 aclrtMemcpy(d,dn,s,sn,kind) 5参数 AscendC 多一个 dstSize
启动 kernel kernel<<<g,b>>>() kernel<<<g,b,0,stream>>>() AscendC stream 必传
同步 cudaDeviceSynchronize() aclrtSynchronizeStream(stream) CUDA 可全局同步
释放内存 cudaFree(d) aclrtFree(d) 相同
销毁 stream 可选 aclrtDestroyStream 必须 AscendC 多一步
重置设备 可选 cudaDeviceReset() aclrtResetDevice 必须 AscendC 多一步
全局清理 不需要(进程退出自动清) aclFinalize() 必须 AscendC 多一步
总步数 7 步 12 步 AscendC 多约 70%

Origin(信息来源)

用户@brian66237 在issue #490中提出,为方便issue管理和后续查询,单独提issue跟踪

Benefit / Necessity (价值/作用)

简化资源创建销毁流程,减少代码量

Design(设计方案)

参考描述中cuda的实现,简化多余步骤

likedislike
Ggao_dafa成员
4月28日 添加了label:requirement
wk0911成员
4月28日 评论:

可以AscendC 极简形式可以:

#include "debug/asc_printf.h"
#include "acl/acl.h"

__global__ __vector__ void hello()
{
    printf("hello\n");
}

int main(int argc, char*argv[])
{
    aclrtSetDevice(0);
    hello<<<4>>>();
    aclrtSynchronizeDevice();
    return 0;
}

后续会将这种极简形式的代码合入asc的样例代码仓敬请关注

likedislike
Wwk0911成员
5月12日 关联了pull request:Add hello_world samples
CANN-robotCANN-robot成员
5月18日 关闭了 issue
CANN-robotCANN-robot成员
5月18日 添加了label:resolved