使用HIP编写GPU 算子向量加法

HIP (Heterogeneous-compute Interface for Portability)来编写一个 GPU 算子operator。HIP 是 AMD 推出的 GPU 编程接口类似 CUDA但可在 AMD 和 NVIDIA GPU 上运行。下面我给你一个完整示例演示如何写一个简单算子比如向量加法Vector Add并用 HIP 编译运行。实践1️⃣ 基本 HIP 算子示例向量加法#includehip/hip_runtime.h#includeiostream#defineN1024// 向量长度#defineTHREADS_PER_BLOCK256// HIP kernel: 向量加法__global__voidvector_add(constfloat*A,constfloat*B,float*C,intn){intidxblockIdx.x*blockDim.xthreadIdx.x;if(idxn){C[idx]A[idx]B[idx];}}intmain(){// 主机向量float*h_Anewfloat[N];float*h_Bnewfloat[N];float*h_Cnewfloat[N];// 初始化向量for(inti0;iN;i){h_A[i]i*1.0f;h_B[i](N-i)*1.0f;}// 设备向量float*d_A,*d_B,*d_C;hipMalloc(d_A,N*sizeof(float));hipMalloc(d_B,N*sizeof(float));hipMalloc(d_C,N*sizeof(float));// 拷贝数据到设备hipMemcpy(d_A,h_A,N*sizeof(float),hipMemcpyHostToDevice);hipMemcpy(d_B,h_B,N*sizeof(float),hipMemcpyHostToDevice);// 计算线程块和网格大小intblocks(NTHREADS_PER_BLOCK-1)/THREADS_PER_BLOCK;// 启动 kernelhipLaunchKernelGGL(vector_add,dim3(blocks),dim3(THREADS_PER_BLOCK),0,0,d_A,d_B,d_C,N);// 拷贝结果回主机hipMemcpy(h_C,d_C,N*sizeof(float),hipMemcpyDeviceToHost);// 验证结果for(inti0;i10;i){// 打印前 10 个元素std::couth_A[i] h_B[i] h_C[i]std::endl;}// 释放内存delete[]h_A;delete[]h_B;delete[]h_C;hipFree(d_A);hipFree(d_B);hipFree(d_C);return0;}2️⃣ 编译 运行# 编译hipcc vector_add.cpp-ovector_add# 运行./vector_add输出前 10 个元素类似0 1024 1024 1 1023 1024 2 1022 1024 ...3️⃣ 说明__global__修饰函数表示这是一个 GPU kernel。hipMalloc/hipFree对应 CUDA 的cudaMalloc/cudaFree。hipMemcpy负责主机和设备间数据拷贝。hipLaunchKernelGGL用来启动 kernelhipLaunchKernelGGL(kernel_name,dim3(blocks),dim3(threads_per_block),shared_mem_bytes,stream,args...)HIP 的算子逻辑和 CUDA 几乎一致只要把cudaXXX改成hipXXX就能跑在 AMD GPU。运行时报错/workspace/models/liuysh# hipcc vector_add.cpp -o vector_add vector_add.cpp:29:5: warning: ignoring return value of function declared with nodiscard attribute [-Wunused-result] 29 | hipMalloc(d_A, N * sizeof(float)); | ^~~~~~~~~ ~~~~~~~~~~~~~~~~~~~~~~~ vector_add.cpp:30:5: warning: ignoring return value of function declared with nodiscard attribute [-Wunused-result] 30 | hipMalloc(d_B, N * sizeof(float)); | ^~~~~~~~~ ~~~~~~~~~~~~~~~~~~~~~~~ vector_add.cpp:31:5: warning: ignoring return value of function declared with nodiscard attribute [-Wunused-result] 31 | hipMalloc(d_C, N * sizeof(float)); | ^~~~~~~~~ ~~~~~~~~~~~~~~~~~~~~~~~ vector_add.cpp:34:5: warning: ignoring return value of function declared with nodiscard attribute [-Wunused-result] 34 | hipMemcpy(d_A, h_A, N * sizeof(float), hipMemcpyHostToDevice); | ^~~~~~~~~ ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ vector_add.cpp:35:5: warning: ignoring return value of function declared with nodiscard attribute [-Wunused-result] 35 | hipMemcpy(d_B, h_B, N * sizeof(float), hipMemcpyHostToDevice); | ^~~~~~~~~ ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ vector_add.cpp:44:5: warning: ignoring return value of function declared with nodiscard attribute [-Wunused-result] 44 | hipMemcpy(h_C, d_C, N * sizeof(float), hipMemcpyDeviceToHost); | ^~~~~~~~~ ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ vector_add.cpp:53:5: warning: ignoring return value of function declared with nodiscard attribute [-Wunused-result] 53 | hipFree(d_A); hipFree(d_B); hipFree(d_C); | ^~~~~~~ ~~~ vector_add.cpp:53:19: warning: ignoring return value of function declared with nodiscard attribute [-Wunused-result] 53 | hipFree(d_A); hipFree(d_B); hipFree(d_C); | ^~~~~~~ ~~~ vector_add.cpp:53:33: warning: ignoring return value of function declared with nodiscard attribute [-Wunused-result] 53 | hipFree(d_A); hipFree(d_B); hipFree(d_C); | ^~~~~~~ ~~~ 9 warnings generated when compiling for gfx906. vector_add.cpp:29:5: warning: ignoring return value of function declared with nodiscard attribute [-Wunused-result] 29 | hipMalloc(d_A, N * sizeof(float)); | ^~~~~~~~~ ~~~~~~~~~~~~~~~~~~~~~~~ vector_add.cpp:30:5: warning: ignoring return value of function declared with nodiscard attribute [-Wunused-result] 30 | hipMalloc(d_B, N * sizeof(float)); | ^~~~~~~~~ ~~~~~~~~~~~~~~~~~~~~~~~ vector_add.cpp:31:5: warning: ignoring return value of function declared with nodiscard attribute [-Wunused-result] 31 | hipMalloc(d_C, N * sizeof(float)); | ^~~~~~~~~ ~~~~~~~~~~~~~~~~~~~~~~~ vector_add.cpp:34:5: warning: ignoring return value of function declared with nodiscard attribute [-Wunused-result] 34 | hipMemcpy(d_A, h_A, N * sizeof(float), hipMemcpyHostToDevice); | ^~~~~~~~~ ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ vector_add.cpp:35:5: warning: ignoring return value of function declared with nodiscard attribute [-Wunused-result] 35 | hipMemcpy(d_B, h_B, N * sizeof(float), hipMemcpyHostToDevice); | ^~~~~~~~~ ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ vector_add.cpp:44:5: warning: ignoring return value of function declared with nodiscard attribute [-Wunused-result] 44 | hipMemcpy(h_C, d_C, N * sizeof(float), hipMemcpyDeviceToHost); | ^~~~~~~~~ ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ vector_add.cpp:53:5: warning: ignoring return value of function declared with nodiscard attribute [-Wunused-result] 53 | hipFree(d_A); hipFree(d_B); hipFree(d_C); | ^~~~~~~ ~~~ vector_add.cpp:53:19: warning: ignoring return value of function declared with nodiscard attribute [-Wunused-result] 53 | hipFree(d_A); hipFree(d_B); hipFree(d_C); | ^~~~~~~ ~~~ vector_add.cpp:53:33: warning: ignoring return value of function declared with nodiscard attribute [-Wunused-result] 53 | hipFree(d_A); hipFree(d_B); hipFree(d_C); | ^~~~~~~ ~~~ 9 warnings generated when compiling for gfx926. vector_add.cpp:29:5: warning: ignoring return value of function declared with nodiscard attribute [-Wunused-result] 29 | hipMalloc(d_A, N * sizeof(float)); | ^~~~~~~~~ ~~~~~~~~~~~~~~~~~~~~~~~ vector_add.cpp:30:5: warning: ignoring return value of function declared with nodiscard attribute [-Wunused-result] 30 | hipMalloc(d_B, N * sizeof(float)); | ^~~~~~~~~ ~~~~~~~~~~~~~~~~~~~~~~~ vector_add.cpp:31:5: warning: ignoring return value of function declared with nodiscard attribute [-Wunused-result] 31 | hipMalloc(d_C, N * sizeof(float)); | ^~~~~~~~~ ~~~~~~~~~~~~~~~~~~~~~~~ vector_add.cpp:34:5: warning: ignoring return value of function declared with nodiscard attribute [-Wunused-result] 34 | hipMemcpy(d_A, h_A, N * sizeof(float), hipMemcpyHostToDevice); | ^~~~~~~~~ ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ vector_add.cpp:35:5: warning: ignoring return value of function declared with nodiscard attribute [-Wunused-result] 35 | hipMemcpy(d_B, h_B, N * sizeof(float), hipMemcpyHostToDevice); | ^~~~~~~~~ ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ vector_add.cpp:44:5: warning: ignoring return value of function declared with nodiscard attribute [-Wunused-result] 44 | hipMemcpy(h_C, d_C, N * sizeof(float), hipMemcpyDeviceToHost); | ^~~~~~~~~ ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ vector_add.cpp:53:5: warning: ignoring return value of function declared with nodiscard attribute [-Wunused-result] 53 | hipFree(d_A); hipFree(d_B); hipFree(d_C); | ^~~~~~~ ~~~ vector_add.cpp:53:19: warning: ignoring return value of function declared with nodiscard attribute [-Wunused-result] 53 | hipFree(d_A); hipFree(d_B); hipFree(d_C); | ^~~~~~~ ~~~ vector_add.cpp:53:33: warning: ignoring return value of function declared with nodiscard attribute [-Wunused-result] 53 | hipFree(d_A); hipFree(d_B); hipFree(d_C); | ^~~~~~~ ~~~ 9 warnings generated when compiling for gfx928. vector_add.cpp:29:5: warning: ignoring return value of function declared with nodiscard attribute [-Wunused-result] 29 | hipMalloc(d_A, N * sizeof(float)); | ^~~~~~~~~ ~~~~~~~~~~~~~~~~~~~~~~~ vector_add.cpp:30:5: warning: ignoring return value of function declared with nodiscard attribute [-Wunused-result] 30 | hipMalloc(d_B, N * sizeof(float)); | ^~~~~~~~~ ~~~~~~~~~~~~~~~~~~~~~~~ vector_add.cpp:31:5: warning: ignoring return value of function declared with nodiscard attribute [-Wunused-result] 31 | hipMalloc(d_C, N * sizeof(float)); | ^~~~~~~~~ ~~~~~~~~~~~~~~~~~~~~~~~ vector_add.cpp:34:5: warning: ignoring return value of function declared with nodiscard attribute [-Wunused-result] 34 | hipMemcpy(d_A, h_A, N * sizeof(float), hipMemcpyHostToDevice); | ^~~~~~~~~ ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ vector_add.cpp:35:5: warning: ignoring return value of function declared with nodiscard attribute [-Wunused-result] 35 | hipMemcpy(d_B, h_B, N * sizeof(float), hipMemcpyHostToDevice); | ^~~~~~~~~ ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ vector_add.cpp:44:5: warning: ignoring return value of function declared with nodiscard attribute [-Wunused-result] 44 | hipMemcpy(h_C, d_C, N * sizeof(float), hipMemcpyDeviceToHost); | ^~~~~~~~~ ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ vector_add.cpp:53:5: warning: ignoring return value of function declared with nodiscard attribute [-Wunused-result] 53 | hipFree(d_A); hipFree(d_B); hipFree(d_C); | ^~~~~~~ ~~~ vector_add.cpp:53:19: warning: ignoring return value of function declared with nodiscard attribute [-Wunused-result] 53 | hipFree(d_A); hipFree(d_B); hipFree(d_C); | ^~~~~~~ ~~~ vector_add.cpp:53:33: warning: ignoring return value of function declared with nodiscard attribute [-Wunused-result] 53 | hipFree(d_A); hipFree(d_B); hipFree(d_C); | ^~~~~~~ ~~~ 9 warnings generated when compiling for gfx936. vector_add.cpp:29:5: warning: ignoring return value of function declared with nodiscard attribute [-Wunused-result] 29 | hipMalloc(d_A, N * sizeof(float)); | ^~~~~~~~~ ~~~~~~~~~~~~~~~~~~~~~~~ vector_add.cpp:30:5: warning: ignoring return value of function declared with nodiscard attribute [-Wunused-result] 30 | hipMalloc(d_B, N * sizeof(float)); | ^~~~~~~~~ ~~~~~~~~~~~~~~~~~~~~~~~ vector_add.cpp:31:5: warning: ignoring return value of function declared with nodiscard attribute [-Wunused-result] 31 | hipMalloc(d_C, N * sizeof(float)); | ^~~~~~~~~ ~~~~~~~~~~~~~~~~~~~~~~~ vector_add.cpp:34:5: warning: ignoring return value of function declared with nodiscard attribute [-Wunused-result] 34 | hipMemcpy(d_A, h_A, N * sizeof(float), hipMemcpyHostToDevice); | ^~~~~~~~~ ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ vector_add.cpp:35:5: warning: ignoring return value of function declared with nodiscard attribute [-Wunused-result] 35 | hipMemcpy(d_B, h_B, N * sizeof(float), hipMemcpyHostToDevice); | ^~~~~~~~~ ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ vector_add.cpp:44:5: warning: ignoring return value of function declared with nodiscard attribute [-Wunused-result] 44 | hipMemcpy(h_C, d_C, N * sizeof(float), hipMemcpyDeviceToHost); | ^~~~~~~~~ ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ vector_add.cpp:53:5: warning: ignoring return value of function declared with nodiscard attribute [-Wunused-result] 53 | hipFree(d_A); hipFree(d_B); hipFree(d_C); | ^~~~~~~ ~~~ vector_add.cpp:53:19: warning: ignoring return value of function declared with nodiscard attribute [-Wunused-result] 53 | hipFree(d_A); hipFree(d_B); hipFree(d_C); | ^~~~~~~ ~~~ vector_add.cpp:53:33: warning: ignoring return value of function declared with nodiscard attribute [-Wunused-result] 53 | hipFree(d_A); hipFree(d_B); hipFree(d_C); | ^~~~~~~ ~~~ 9 warnings generated when compiling for gfx938. vector_add.cpp:29:5: warning: ignoring return value of function declared with nodiscard attribute [-Wunused-result] 29 | hipMalloc(d_A, N * sizeof(float)); | ^~~~~~~~~ ~~~~~~~~~~~~~~~~~~~~~~~ vector_add.cpp:30:5: warning: ignoring return value of function declared with nodiscard attribute [-Wunused-result] 30 | hipMalloc(d_B, N * sizeof(float)); | ^~~~~~~~~ ~~~~~~~~~~~~~~~~~~~~~~~ vector_add.cpp:31:5: warning: ignoring return value of function declared with nodiscard attribute [-Wunused-result] 31 | hipMalloc(d_C, N * sizeof(float)); | ^~~~~~~~~ ~~~~~~~~~~~~~~~~~~~~~~~ vector_add.cpp:34:5: warning: ignoring return value of function declared with nodiscard attribute [-Wunused-result] 34 | hipMemcpy(d_A, h_A, N * sizeof(float), hipMemcpyHostToDevice); | ^~~~~~~~~ ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ vector_add.cpp:35:5: warning: ignoring return value of function declared with nodiscard attribute [-Wunused-result] 35 | hipMemcpy(d_B, h_B, N * sizeof(float), hipMemcpyHostToDevice); | ^~~~~~~~~ ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ vector_add.cpp:44:5: warning: ignoring return value of function declared with nodiscard attribute [-Wunused-result] 44 | hipMemcpy(h_C, d_C, N * sizeof(float), hipMemcpyDeviceToHost); | ^~~~~~~~~ ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ vector_add.cpp:53:5: warning: ignoring return value of function declared with nodiscard attribute [-Wunused-result] 53 | hipFree(d_A); hipFree(d_B); hipFree(d_C); | ^~~~~~~ ~~~ vector_add.cpp:53:19: warning: ignoring return value of function declared with nodiscard attribute [-Wunused-result] 53 | hipFree(d_A); hipFree(d_B); hipFree(d_C); | ^~~~~~~ ~~~ vector_add.cpp:53:33: warning: ignoring return value of function declared with nodiscard attribute [-Wunused-result] 53 | hipFree(d_A); hipFree(d_B); hipFree(d_C); | ^~~~~~~ ~~~ 9 warnings generated when compiling for host. rootf10r1n06:/workspace/models/liuysh# ./vector_add 0 1024 1024 1 1023 1024 2 1022 1024 3 1021 1024 4 1020 1024 5 1019 1024 6 1018 1024 7 1017 1024 8 1016 1024 9 1015 1024 Warning: Resource leak detected by SharedSignalPool, 68 Signals leaked.根据编译和运行日志来看遇到的“报错”主要是编译时的警告和运行时的资源泄漏提示1️⃣ 编译时警告日志里反复出现warning: ignoring return value of function declared with nodiscard attribute [-Wunused-result]出现在hipMalloc、hipMemcpy、hipFree这几行。原因从 HIP 5.x 版本开始这些函数都加了[[nodiscard]]属性意思是返回值不能忽略因为返回值中可能包含错误码。你现在的代码写法直接调用了函数没有检查返回值所以会产生警告。示例hipError_t errhipMalloc(d_A,N*sizeof(float));if(err!hipSuccess){std::cerrhipMalloc failed: hipGetErrorString(err)std::endl;}✅ 建议编译通过不会阻止程序运行但在生产环境中最好检查返回值。2️⃣ 运行时提示Warning: Resource leak detected by SharedSignalPool, 68 Signals leaked.这是 HIP / ROCm 在运行 kernel 后检测到 GPU 资源没有完全释放。可能原因kernel 启动后没有同步 (hipDeviceSynchronize())。某些 GPU 信号或流被泄漏在你的简单示例里可能是内部 runtime 的管理问题。解决方法在 kernel 调用后加同步确保所有操作完成hipLaunchKernelGGL(vector_add,dim3(blocks),dim3(THREADS_PER_BLOCK),0,0,d_A,d_B,d_C,N);hipDeviceSynchronize();// 等待 kernel 执行完成对于资源释放问题保证hipFree被调用即可。你原来的代码已经调用了hipFree(d_A)、hipFree(d_B)、hipFree(d_C)所以主要是 runtime 内部 signal 还未完全回收。加同步通常可以消除警告。✅ 总结编译警告nodiscard警告可以通过检查返回值消除。运行时警告SharedSignalPool 资源泄漏一般加hipDeviceSynchronize()后消失。程序逻辑你的向量加法结果是正确的所以算子逻辑没有问题。