DCU异构程序——矩阵转置

目录

一、概述

二、程序实现

三、编译运行

---

一、概述

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

HIP程序流程

​

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

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

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

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

二、程序实现

        下面是对矩阵转置的具体实现，MatrixTranspose.cpp：

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

#define WIDTH  1024
#define NUM (WIDTH * WIDTH)
#define THREADS_PER_BLOCK_X 4
#define THREADS_PER_BLOCK_Y 4
#define THREADS_PER_BLOCK_Z 1


/*核函数*/
__global__ void matrixTranspose(float *out, float *in, const int width)
{
    int x = blockDim.x * blockIdx.x + threadIdx.x;
    int y = blockDim.y * blockIdx.y + threadIdx.y;

    out[y * width + x] = in[x * width + y];
}

/*CPU端矩阵转置函数*/
void matrixTransposeCPUReference(float *output, float *input, const unsigned int width)
{
    for(unsigned int j = 0; j < width; j++)
    {
        for(unsigned int i = 0; i < width; i++)
        {
            output[i * width + j] = input[j * width + i];
        }
    }
}


int main(int argc, char *argv[])
{
    float *Matrix;
    float *TransposeMatrix;
    float *cpuTransposeMatrix;
    
    float *gpuMatrix;
    float *gpuTransposeMatrix;

    hipDeviceProp_t devProp;
    hipGetDeviceProperties(&devProp, 0);

    std::cout<< "Device name" <<devProp.name<<std::endl;

    int i;
    int errors;

    Matrix = (float*)malloc(NUM * sizeof(float));
    TransposeMatrix = (float*)malloc(NUM * sizeof(float));
    cpuTransposeMatrix = (float*)malloc(NUM * sizeof(float));

    for(i = 0; i < NUM; i++)
    {
        Matrix[i] = (float)i * 10.00f;
    }

    hipMalloc((void**)&gpuMatrix, NUM * sizeof(float));
    hipMalloc((void**)&gpuTransposeMatrix, NUM * sizeof(float));

    hipMemcpy(gpuMatrix, Matrix, NUM * sizeof(float), hipMemcpyHostToDevice);
    
    hipLaunchKernelGGL(matrixTranspose, dim3(WIDTH / THREADS_PER_BLOCK_X, WIDTH/THREADS_PER_BLOCK_Y), dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y), 0, 0, gpuTransposeMatrix, gpuMatrix, WIDTH);

    hipMemcpy(TransposeMatrix, gpuTransposeMatrix, NUM * sizeof(float), hipMemcpyDeviceToHost);

    matrixTransposeCPUReference(cpuTransposeMatrix, Matrix, WIDTH);

    errors = 0;
    double eps = 1.0e-6;
    
    for(i = 0; i < NUM; i++)
    {
        if(std::abs(TransposeMatrix[i] - cpuTransposeMatrix[i]) > eps)
        {
            errors++;
        }
    }

    if(errors != 0)
    {
        printf("FAILED: %d errors\n", errors);
    }
    else
    {
        printf("PASSED!\n");
    }

    hipFree(gpuMatrix);
    hipFree(gpuTransposeMatrix);

    free(Matrix);
    free(TransposeMatrix);
    free(cpuTransposeMatrix);

    return errors;
}

三、编译运行

        HIP程序采用hipcc编译

运行结果：