LongLong's Blog

#include <OpenCL/opencl.h>

#ifdef __APPLE__
#include <OpenCL/opencl.h>
#else
#include <CL/cl.h>
#endif

#include <OpenCL/opencl.h>
#include <stdio.h>
#include <stdlib.h>
#define KERNEL(...)#__VA_ARGS__
//Kernel代码，计算乘积矩阵i行j列的值
const char* KernelSource = KERNEL(
    __kernel void matrix_mul(int m, int p, int n, __global float* A, __global float* B, __global float* C) {\n
        int i = get_global_id(0);\n
        int j = get_global_id(1);\n
        float sum = 0;\n
        for (int k = 0; k < p; k++) {\n
            sum += A[i * m + k] * B[p * k + j];\n
        }\n
        C[i * m + j] = sum;\n
    }
);

int main() {
    int err;
    cl_device_id device_id;
    cl_context context;
    cl_command_queue commands;
    cl_program program;
    cl_kernel kernel;

    int gpu = 1;
    //获取计算设备，这里获取GPU设备
    err = clGetDeviceIDs(NULL, gpu ? CL_DEVICE_TYPE_GPU : CL_DEVICE_TYPE_CPU, 1, &device_id, NULL);
    if (err != CL_SUCCESS) {
        return EXIT_FAILURE;
    }
    //初始化上下文和命令队列
    context = clCreateContext(0, 1, &device_id, NULL, NULL, &err);
    if (!context) {
        return EXIT_FAILURE;
    }
    commands = clCreateCommandQueue(context, device_id, 0, &err);
    if (!commands) {
        return EXIT_FAILURE;
    }

    //编译Kernel代码
    program = clCreateProgramWithSource(context, 1, (const char**) & KernelSource, NULL, &err);
    if (!program) {
        return EXIT_FAILURE;
    }

    err = clBuildProgram(program, 0, NULL, NULL, NULL, NULL);
    if (err != CL_SUCCESS) {
        return EXIT_FAILURE;
    }

    kernel = clCreateKernel(program, "matrix_mul", &err);
    if (!kernel || err != CL_SUCCESS) {
        return EXIT_FAILURE;
    }
    int count = 800;
    int m = count;
    int p = count;
    int n = count;

    cl_mem input1;
    cl_mem input2;
    cl_mem output;

    //生成两个随机矩阵
    float A[count][count];
    float B[count][count];
    for (int i = 0; i < count; i++) {
        for (int j = 0; j < count; j++) {
            A[i][j] = (float) rand() / RAND_MAX;
            B[i][j] = (float) rand() / RAND_MAX;
        }
    }

    //创建将输入的两个矩阵和乘积结果矩阵的入缓冲区，并将随机生成的数据写入缓冲区
    input1 = clCreateBuffer(context,  CL_MEM_READ_ONLY,  sizeof(float) * m * p, NULL, NULL);
    input2 = clCreateBuffer(context,  CL_MEM_READ_ONLY,  sizeof(float) * p * n, NULL, NULL);
    output = clCreateBuffer(context,  CL_MEM_WRITE_ONLY,  sizeof(float) * m * n, NULL, NULL);
    if (!input1 || !input2 || !output) {
        return EXIT_FAILURE;
    }
    err = 0;
    err = clEnqueueWriteBuffer(commands, input1, CL_TRUE, 0, m * p * sizeof(float), A, 0, NULL, NULL);
    err |= clEnqueueWriteBuffer(commands, input2, CL_TRUE, 0, p * n * sizeof(float), B, 0, NULL, NULL);
    if (err != CL_SUCCESS) {
        return EXIT_FAILURE;
    }

    //设置Kernel参数
    err = 0;
    err  = clSetKernelArg(kernel, 0, sizeof(cl_mem), &m);
    err |= clSetKernelArg(kernel, 1, sizeof(cl_mem), &p);
    err |= clSetKernelArg(kernel, 2, sizeof(cl_mem), &n);
    err |= clSetKernelArg(kernel, 3, sizeof(cl_mem), &input1);
    err |= clSetKernelArg(kernel, 4, sizeof(cl_mem), &input2);
    err |= clSetKernelArg(kernel, 5, sizeof(cl_mem), &output);

    if (err != CL_SUCCESS) {
        return EXIT_FAILURE;
    }

    //设置计算任务，并开始计算
    //其中global中设置的两个变量对应Kernel中get_global_id中获取的两个值
    size_t local[2] = {20, 20};
    size_t global[2] = {count, count};

    err = clEnqueueNDRangeKernel(commands, kernel, 2, NULL, global, local, 0, NULL, NULL);
    if (err != CL_SUCCESS) {
        printf("%d\n", err);
        return EXIT_FAILURE;
    }

    //等待计算结果并将结果读入数组
    float C[m][n];
    clFinish(commands);

    err = clEnqueueReadBuffer(commands, output, CL_TRUE, 0, sizeof(float) * m * n, C, 0, NULL, NULL);
    if (err != CL_SUCCESS) {
        return EXIT_FAILURE;
    }

    //清理释放资源
    clReleaseMemObject(input1);
    clReleaseMemObject(input2);
    clReleaseMemObject(output);
    clReleaseProgram(program);
    clReleaseKernel(kernel);
    clReleaseCommandQueue(commands);
    clReleaseContext(context);

    //输出结果
    for (int i = 0; i < m; i++) {
        for (int j = 0; j < n; j++) {
            printf("%f\t", C[i][j]);
        }
        printf("\n");
    }
    return 0;
}