DCU 异构程序——数组相加

目录

一、概述

二、实现

三、编译运行

---

        上一篇文章我们介绍了常见的异构程序设计语言（↓传送门↓），感兴趣的小伙伴可以跳转。本节我们基于国产DCU加速卡使用HIP进行编程示范。

常见异构程序设计语言-优快云博客https://blog.youkuaiyun.com/L_peanut/article/details/144058178

一、概述

        HIP属于显式编程模型，需要在程序中明确写出并行控制语句，包括数据传输、核函数启动等。核函数是运行在DCU上的函数，在CPU端运行的部分称为主机端（主要是执行管理和启动），DCU端运行的部分称为设备端（用于执行计算）。大概的流程如下图：

HIP程序流程

        ①主机端将需要并行计算的数据通过hipMemcpy()传递给DCU（将CPU存储的内容传递给DCU的显存）；

        ②调用核函数启动函数hipLaunchKernelGGL()启动DCU，开始执行计算；

        ③设备端将计算好的结果数据通过hipMemcpy()从DCU复制回CPU。

        hipMemcpy()是阻塞式的，数据复制完成后才可以执行后续的程序；hipLanuchKernelGGL()是非阻塞式的，执行完后程序继续向后执行，但是在Kernel没有计算完成之前，最后一个hipMemcpy()是不会开始的，这是由于HIP的Stream机制。

二、实现

        下面是对数组相加的具体实现，vectoradd.cpp：

#include <iostream>
#include "hip/hip_runtime.h"

#define N 10000

/*核函数*/
__global__ void add(float *d_A, float *d_B, float *d_C)
{
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    if(tid < N)
    {
        d_C[tid] = d_A[tid] + d_B[tid];
    }
}


int main(int argc, char *argv[])
{
    /*申请数组空间*/
    float *A = (float *)malloc(N * sizeof(float));
    float *B = (float *)malloc(N * sizeof(float));
    float *C = (float *)malloc(N * sizeof(float));
    float *d_A = NULL;
    float *d_B = NULL;
    float *d_C = NULL;
    hipMalloc((void **) &d_A, N * sizeof(float));
    hipMalloc((void **) &d_B, N * sizeof(float));
    hipMalloc((void **) &d_C, N * sizeof(float));

    /*数据初始化*/
    for(int i = 0; i < N; i++)
    {
        A[i] = 1;
        B[i] = 1;
        C[i] = 0;
    }

    /*数据复制*/
    hipMemcpy(d_A, A, sizeof(float) * N, hipMemcpyHostToDevice);
    hipMemcpy(d_B, B, sizeof(float) * N, hipMemcpyHostToDevice);
    hipMemcpy(d_C, C, sizeof(float) * N, hipMemcpyHostToDevice);
    dim3 blocksize(256, 1);
    dim3 gridsize(N / 256 + 1; 1);

    /*数组相加*/
    add<<<gridsize, blocksize>>>(d_A, d_B, d_C);

    /*结果验证*/
    hipMemcpy(C, d_C, sizeof(float) * N, hipMemcpyDeviceToHost);
    for(int i = 0; i < N; i++)
    {
        std::cout<<C[i]<<std::endl
    }

    /*释放内存空间*/
    free(A);
    free(B);
    free(C);
    hipFree(d_A);
    hipFree(d_B);
    hipFree(d_C);

    return 0;
}

三、编译运行

        HIP程序采用hipcc编译。

运行结果：