异构计算——OPENCL计算矩阵幂

OPENCL计算矩阵幂

实验内容：
用OpenCL编程模型实现矩阵的次幂。
要求实现暴力算法和高效算法，同时对比分析一下相同OpenCL程序分别运行在纯多核CPU环境下以及异构GPU环境下的性能。
对于一个 的方阵 ，计算的次幂。首先，生成一个的方阵，保证每行每列元素之和满足(0,1])。
·暴力算法
N个矩阵相乘
·高效算法
利用矩阵乘法的结合律

暴力算法：

#include <CL/cl.h>   //MACOS 是OPENCL.h
#include <iostream>
#include <fstream>
#include <sstream>
#include <unistd.h>
#include<time.h>
#include<stdio.h>
#include<stdlib.h>
#include <time.h>
#include <boost/algorithm/string.hpp>


using namespace std;


const int N = 256; //矩阵大小
const int M = 20000; //幂次数


//一、 选择OpenCL平台并创建一个上下文
cl_context CreateContext()
{
    cl_int errNum;
    cl_uint numPlatforms;
    cl_platform_id firstPlatformId;
    cl_context context = NULL;

    //选择可用的平台中的第一个
    errNum = clGetPlatformIDs(1, &firstPlatformId, &numPlatforms);
    if (errNum != CL_SUCCESS || numPlatforms <= 0)
    {
        std::cerr << "Failed to find any OpenCL platforms." << std::endl;
        return NULL;
    }

    //创建一个OpenCL上下文环境
    cl_context_properties contextProperties[] =
    {
        CL_CONTEXT_PLATFORM,
        (cl_context_properties)firstPlatformId,
        0
    };
    context = clCreateContextFromType(contextProperties, CL_DEVICE_TYPE_GPU,
                                      NULL, NULL, &errNum);

    return context;
}


//二、 创建设备并创建命令队列
cl_command_queue CreateCommandQueue(cl_context context, cl_device_id *device)
{
    cl_int errNum;
    cl_device_id *devices;
    cl_command_queue commandQueue = NULL;
    size_t deviceBufferSize = -1;

    // 获取设备缓冲区大小
    errNum = clGetContextInfo(context, CL_CONTEXT_DEVICES, 0, NULL, &deviceBufferSize);

    if (deviceBufferSize <= 0)
    {
        std::cerr << "No devices available.";
        return NULL;
    }

    // 为设备分配缓存空间
    devices = new cl_device_id[deviceBufferSize / sizeof(cl_device_id)];
    errNum = clGetContextInfo(context, CL_CONTEXT_DEVICES, deviceBufferSize, devices, NULL);

    //选取可用设备中的第一个
    commandQueue = clCreateCommandQueue(context, devices[0], 0, NULL);

    *device = devices[0];
    delete[] devices;
    return commandQueue;
}


// 三、创建和构建程序对象
cl_program CreateProgram(cl_context context, cl_device_id device, const char* fileName)
{
    cl_int errNum;
    cl_program program;

    std::ifstream kernelFile(fileName, std::ios::in);
    if (!kernelFile.is_open())
    {
        std::cerr << "Failed to open file for reading: " << fileName << std::endl;
        return NULL;
    }

    std::ostringstream oss;
    oss << kernelFile.rdbuf();

    std::string srcStdStr = oss.str();
    const char *srcStr = srcStdStr.c_str();
    program = clCreateProgramWithSource(context, 1,
                                        (const char**)&srcStr,
                                        NULL, NULL);

    errNum = clBuildProgram(program, 0, NULL, NULL, NULL, NULL);

    return program;
}

//创建和构建程序对象
bool CreateMemObjects(cl_context context, cl_mem memObjects[3],
                      int *a, int *b)
{
    memObjects[0] = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
                                   sizeof(int) * N*N, a, NULL);
    memObjects[1] = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
                                   sizeof(int) * N*N, b, NULL);
    memObjects[2] = clCreateBuffer(context, CL_MEM_READ_WRITE,
                                   sizeof(int) * N*N, NULL, NULL);
    return true;
}


// 释放OpenCL资源
void Cleanup(cl_context context, cl_command_queue commandQueue,
             cl_program program, cl_kernel kernel, cl_mem memObjects[3])
{
    for (int i = 0; i < 3; i++)
    {
        if (memObjects[i] != 0)
            clReleaseMemObject(memObjects[i]);
    }
    if (commandQueue != 0)
        clReleaseCommandQueue(commandQueue);

    if (kernel != 0)
        clReleaseKernel(kernel);

    if (program != 0)
        clReleaseProgram(program);

    if (context != 0)
        clReleaseContext(context);
}


int main(int argc, char** argv)
{



    cl_context context = 0;
    cl_command_queue commandQueue = 0;
    cl_program program = 0;
    cl_device_id device = 0;
    cl_kernel kernel = 0;
    cl_mem memObjects[3] = { 0, 0, 0 };
    cl_int errNum;
   // uint64_t t1,t2,t3;
    clock_t t1,t2,t3;


    const char* filename = "/home/stu3019244205/data/stupid.cl";
    // 一、选择OpenCL平台并创建一个上下文
    context = CreateContext();

    // 二、 创建设备并创建命令队列
    commandQueue = CreateCommandQueue(context, &device);

    //创建和构建程序对象
    program = CreateProgram(context, device, filename);

    // 四、 创建OpenCL内核并分配内存空间
    kernel = clCreateKernel(program, "hello_kernel", NULL);

    //创建要处理的数据
    int result[N*N]{0};
    int CPU_result[N*N]{0};
    int a[N*N];
    int b[N*N];
    int temp1[N*N];
    for (int i = 0; i < N; i++)
    {
        for(int j = 0;j < N;j++)
        {
            a[i*N+j]=2;
            b[i*N+j]=2;
            temp1[i*N+j]=2;
        }

    }
    t1 = clock();  
    printf("t1 = %.8f\n",(double)t1);
    //cpu串行处理代码
    int n = M-1;
    while(n--){

    for(int i=0;i<N;i++){
        for(int j=0;j<N;j++){
            for(int k=0;k<N;k++){
            CPU_result[i*N+j] += temp1[i*N+k]*a[k*N+j]; 
        }
      }
     }
     
    for(int i=0;i<N;i++)
    {
        for(int j=0;j<N;j++)
        {
            temp1[i*N+j] = CPU_result[i*N+j];
        }
    }
 }

    t2 = clock(); //mach_absolute_time();
    printf("t2 = %.8f\n",(double)t2);

    //创建内存对象
    if (!CreateMemObjects(context, memObjects, a, b))
    {
        Cleanup(context, commandQueue, program, kernel, memObjects);
        return 1;
    }

    // 五、 设置内核数据并执行内核
    errNum = clSetKernelArg(kernel, 0, sizeof(cl_mem), &memObjects[0]);
    errNum |= clSetKernelArg(kernel, 1, sizeof(cl_mem), &memObjects[1]);
    errNum |= clSetKernelArg(kernel, 2, sizeof(cl_mem), &memObjects[2]);
    errNum |= clSetKernelArg(kernel, 3, sizeof(int), &N);
    errNum |= clSetKernelArg(kernel, 4, sizeof(int), &M);
    
    size_t globalWorkSize[2] = {4,4};
    size_t localWorkSize[2] = {4,4};

    errNum = clEnqueueNDRangeKernel(commandQueue, kernel, 1, NULL, //参数1代表globalWorkSize等是一（零）维数。
                                    globalWorkSize, localWorkSize,
                                    0, NULL, NULL);

    // 六、 读取执行结果并释放OpenCL资源
    errNum = clEnqueueReadBuffer(commandQueue, memObjects[2], CL_TRUE,
                                 0, N*N * sizeof(int), result,
                                 0, NULL, NULL);

    t3 = clock();  




    printf("CPU串行 t = %.8f\n",(float)(t2-t1)/CLOCKS_PER_SEC);
    printf("异构GPU t = %.8f \n",(double)(t3-t2)/CLOCKS_PER_SEC);

    std::cout << std::endl;
    std::cout << "Executed program succesfully." << std::endl;
    getchar();
    Cleanup(context, commandQueue, program, kernel, memObjects);

    return 0;
}

关键有两点：搭建环境、算矩阵幂

编译： nvcc -o hello HelloWorld.cpp -l OpenCL

流程图：

加速比曲线：
（1）异构GPU与多核CPU对比
曲线一：暴力算法

曲线二：高效算法：

（2）暴力算法和高效算法对比（异构GPU）