2. 编写第一个OpenCL内核

任务说明

本案例目标是在Mali GPU上运行一个OpenCL向量加法内核，验证计算结果与CPU参考结果一致。你需要完成：编写.cl内核文件、编写主机端C程序完成数据初始化与内核调度、编译并运行、对比结果。

1. 内核代码 (vector_add.cl)

__kernel void vector_add(__global const float* a,
                         __global const float* b,
                         __global float* c,
                         unsigned int n)
{
    int i = get_global_id(0);
    if (i < n) {
        c[i] = a[i] + b[i];
    }
}

2. 主机端代码 (main.c)

#include <stdio.h>
#include <stdlib.h>
#include <CL/cl.h>

#define VEC_SIZE 1024

int main(void) {
    cl_platform_id platform;
    cl_device_id device;
    cl_context context;
    cl_command_queue queue;
    cl_program program;
    cl_kernel kernel;
    cl_mem buf_a, buf_b, buf_c;
    float *h_a, *h_b, *h_c, *h_ref;
    size_t bytes = VEC_SIZE * sizeof(float);
    int err;

    // 1. 获取平台与设备（优先Mali GPU）
    err = clGetPlatformIDs(1, &platform, NULL);
    err |= clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 1, &device, NULL);
    if (err != CL_SUCCESS) {
        printf("无法获取OpenCL平台/GPU设备\n");
        return -1;
    }

    // 2. 创建上下文与命令队列
    context = clCreateContext(NULL, 1, &device, NULL, NULL, &err);
    queue = clCreateCommandQueue(context, device, 0, &err);

    // 3. 分配主机内存并初始化
    h_a = (float*)malloc(bytes);
    h_b = (float*)malloc(bytes);
    h_c = (float*)malloc(bytes);
    h_ref = (float*)malloc(bytes);
    for (int i = 0; i < VEC_SIZE; i++) {
        h_a[i] = (float)i * 0.5f;
        h_b[i] = (float)i * 1.5f;
        h_ref[i] = h_a[i] + h_b[i];  // CPU参考结果
    }

    // 4. 创建GPU缓冲区
    buf_a = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
                           bytes, h_a, &err);
    buf_b = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
                           bytes, h_b, &err);
    buf_c = clCreateBuffer(context, CL_MEM_WRITE_ONLY, bytes, NULL, &err);

    // 5. 编译内核
    FILE *fp = fopen("vector_add.cl", "r");
    char *src = (char*)malloc(65536);
    size_t src_len = fread(src, 1, 65536, fp);
    fclose(fp);
    program = clCreateProgramWithSource(context, 1, (const char**)&src,
                                        &src_len, &err);
    err = clBuildProgram(program, 1, &device, NULL, NULL, NULL);
    if (err != CL_SUCCESS) {
        char log[4096];
        clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG,
                              4096, log, NULL);
        printf("内核编译失败:\n%s\n", log);
        return -1;
    }
    free(src);

    // 6. 设置内核参数并执行
    kernel = clCreateKernel(program, "vector_add", &err);
    clSetKernelArg(kernel, 0, sizeof(cl_mem), &buf_a);
    clSetKernelArg(kernel, 1, sizeof(cl_mem), &buf_b);
    clSetKernelArg(kernel, 2, sizeof(cl_mem), &buf_c);
    clSetKernelArg(kernel, 3, sizeof(unsigned int), &(unsigned int){VEC_SIZE});

    size_t global = VEC_SIZE;
    clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &global, NULL, 0, NULL, NULL);
    clFinish(queue);

    // 7. 读取结果并验证
    clEnqueueReadBuffer(queue, buf_c, CL_TRUE, 0, bytes, h_c, 0, NULL, NULL);
    int errors = 0;
    for (int i = 0; i < VEC_SIZE; i++) {
        if (h_c[i] != h_ref[i]) {
            errors++;
            if (errors < 5)
                printf("索引%d: GPU=%f CPU=%f\n", i, h_c[i], h_ref[i]);
        }
    }
    if (errors == 0)
        printf("验证通过：%d个元素全部正确\n", VEC_SIZE);
    else
        printf("验证失败：%d个元素不匹配\n", errors);

    // 8. 清理资源
    clReleaseMemObject(buf_a);
    clReleaseMemObject(buf_b);
    clReleaseMemObject(buf_c);
    clReleaseKernel(kernel);
    clReleaseProgram(program);
    clReleaseCommandQueue(queue);
    clReleaseContext(context);
    free(h_a); free(h_b); free(h_c); free(h_ref);
    return 0;
}

3. 编译与运行

# 安装Mali驱动与OpenCL头文件后执行：
gcc -o vector_add main.c -lOpenCL
./vector_add

4. 验证方法

程序运行后，主机端会逐元素比较GPU计算结果与CPU参考值。若所有1024个元素完全一致，输出“验证通过”；否则输出不匹配的索引及对应值。判定标准：errors == 0，即GPU输出与CPU逐元素计算结果完全一致。