// 查询所有 OpenCL 平台
cl_uint num_platforms;
cl_int err = clGetPlatformIDs(0, NULL, &num_platforms);
if (err != CL_SUCCESS || num_platforms == 0) {
    printf("No OpenCL platforms found (error: %d)\n", err);
    return -1;
}

cl_platform_id *platforms = malloc(sizeof(cl_platform_id) * num_platforms);
clGetPlatformIDs(num_platforms, platforms, NULL);

char info[1024];
for (cl_uint i = 0; i < num_platforms; i++) {
    clGetPlatformInfo(platforms[i], CL_PLATFORM_NAME, sizeof(info), info, NULL);
    printf("Platform %u: %s\n", i, info);

    clGetPlatformInfo(platforms[i], CL_PLATFORM_VENDOR, sizeof(info), info, NULL);
    printf("  Vendor:  %s\n", info);

    clGetPlatformInfo(platforms[i], CL_PLATFORM_VERSION, sizeof(info), info, NULL);
    printf("  Version: %s\n", info);

    clGetPlatformInfo(platforms[i], CL_PLATFORM_EXTENSIONS, sizeof(info), info, NULL);
    printf("  Extensions: %s\n", info);
}

cl_uint num_devices;
clGetDeviceIDs(platforms[0], CL_DEVICE_TYPE_ALL, 0, NULL, &num_devices);

cl_device_id *devices = malloc(sizeof(cl_device_id) * num_devices);
clGetDeviceIDs(platforms[0], CL_DEVICE_TYPE_ALL, num_devices, devices, NULL);

for (cl_uint i = 0; i < num_devices; i++) {
    cl_device_type type;
    char name[256];
    cl_uint compute_units;
    cl_ulong global_mem, local_mem;
    size_t max_work_group;
    size_t max_work_items[3];

    clGetDeviceInfo(devices[i], CL_DEVICE_TYPE, sizeof(type), &type, NULL);
    clGetDeviceInfo(devices[i], CL_DEVICE_NAME, sizeof(name), name, NULL);
    clGetDeviceInfo(devices[i], CL_DEVICE_MAX_COMPUTE_UNITS, sizeof(compute_units), &compute_units, NULL);
    clGetDeviceInfo(devices[i], CL_DEVICE_GLOBAL_MEM_SIZE, sizeof(global_mem), &global_mem, NULL);
    clGetDeviceInfo(devices[i], CL_DEVICE_LOCAL_MEM_SIZE, sizeof(local_mem), &local_mem, NULL);
    clGetDeviceInfo(devices[i], CL_DEVICE_MAX_WORK_GROUP_SIZE, sizeof(max_work_group), &max_work_group, NULL);
    clGetDeviceInfo(devices[i], CL_DEVICE_MAX_WORK_ITEM_SIZES, sizeof(max_work_items), max_work_items, NULL);

    printf("\nDevice %u: %s\n", i, name);
    printf("  Type:          %s\n",
           type == CL_DEVICE_TYPE_GPU ? "GPU" :
           type == CL_DEVICE_TYPE_CPU ? "CPU" :
           type == CL_DEVICE_TYPE_ACCELERATOR ? "Accelerator" : "Other");
    printf("  Compute Units: %u\n", compute_units);
    printf("  Global Memory: %lu MB\n", global_mem / (1024 * 1024));
    printf("  Local Memory:  %lu KB\n", local_mem / 1024);
    printf("  Max Work Group: %zu\n", max_work_group);
    printf("  Max Work Items: [%zu, %zu, %zu]\n",
           max_work_items[0], max_work_items[1], max_work_items[2]);
}

// 方式 1：自动选择平台上的 GPU 设备
cl_context_properties properties[] = {
    CL_CONTEXT_PLATFORM, (cl_context_properties)platforms[0],
    0  // 结束标志
};

cl_int err;
cl_context context = clCreateContextFromType(
    properties,
    CL_DEVICE_TYPE_GPU,  // 只要 GPU
    NULL,                // 回调函数
    NULL,                // 回调数据
    &err
);

if (err != CL_SUCCESS) {
    printf("Failed to create context: %d\n", err);
    return -1;
}

// 方式 2：指定具体设备
cl_context context = clCreateContext(
    properties,
    1,              // 设备数量
    &devices[0],    // 设备列表
    NULL,           // 回调函数
    NULL,           // 回调数据
    &err
);

cl_uint ctx_num_devices;
clGetContextInfo(context, CL_CONTEXT_NUM_DEVICES, sizeof(ctx_num_devices), &ctx_num_devices, NULL);
printf("Context has %u devices\n", ctx_num_devices);

size_t dev_list_size;
clGetContextInfo(context, CL_CONTEXT_DEVICES, 0, NULL, &dev_list_size);
cl_device_id *ctx_devices = malloc(dev_list_size);
clGetContextInfo(context, CL_CONTEXT_DEVICES, dev_list_size, ctx_devices, NULL);

// OpenCL 1.2 方式
cl_command_queue queue = clCreateCommandQueue(
    context,
    devices[0],
    CL_QUEUE_PROFILING_ENABLE,  // 启用性能分析
    &err
);

// OpenCL 2.0+ 方式
cl_queue_properties props[] = {
    CL_QUEUE_PROPERTIES, CL_QUEUE_PROFILING_ENABLE,
    0
};
cl_command_queue queue = clCreateCommandQueueWithProperties(
    context, devices[0], props, &err
);

// vector_add.cl - OpenCL C 内核源码
__kernel void vector_add(
    __global const float *a,
    __global const float *b,
    __global float *out,
    const int n)
{
    int gid = get_global_id(0);
    if (gid < n) {
        out[gid] = a[gid] + b[gid];
    }
}

// 读取源码
const char *source = readFile("vector_add.cl");
size_t source_len = strlen(source);

// 创建程序对象
cl_program program = clCreateProgramWithSource(
    context, 1, &source, &source_len, &err
);

// 编译
err = clBuildProgram(program, 1, &devices[0], "-cl-fast-relaxed-math", NULL, NULL);
if (err != CL_SUCCESS) {
    // 获取编译日志
    size_t log_size;
    clGetProgramBuildInfo(program, devices[0], CL_PROGRAM_BUILD_LOG, 0, NULL, &log_size);
    char *log = malloc(log_size);
    clGetProgramBuildInfo(program, devices[0], CL_PROGRAM_BUILD_LOG, log_size, log, NULL);
    printf("Build log:\n%s\n", log);
    free(log);
    return -1;
}

// 创建内核对象
cl_kernel kernel = clCreateKernel(program, "vector_add", &err);

const int N = 1024 * 1024;
float *host_a = malloc(N * sizeof(float));
float *host_b = malloc(N * sizeof(float));
float *host_out = malloc(N * sizeof(float));

// 填充数据
for (int i = 0; i < N; i++) {
    host_a[i] = (float)i;
    host_b[i] = (float)(i * 2);
}

// 创建设备缓冲区
cl_mem buf_a = clCreateBuffer(context, CL_MEM_READ_ONLY,  N * sizeof(float), NULL, &err);
cl_mem buf_b = clCreateBuffer(context, CL_MEM_READ_ONLY,  N * sizeof(float), NULL, &err);
cl_mem buf_out = clCreateBuffer(context, CL_MEM_WRITE_ONLY, N * sizeof(float), NULL, &err);

// 写入数据到设备
clEnqueueWriteBuffer(queue, buf_a, CL_TRUE, 0, N * sizeof(float), host_a, 0, NULL, NULL);
clEnqueueWriteBuffer(queue, buf_b, CL_TRUE, 0, N * sizeof(float), host_b, 0, NULL, NULL);

// 设置内核参数
clSetKernelArg(kernel, 0, sizeof(cl_mem), &buf_a);
clSetKernelArg(kernel, 1, sizeof(cl_mem), &buf_b);
clSetKernelArg(kernel, 2, sizeof(cl_mem), &buf_out);
clSetKernelArg(kernel, 3, sizeof(int), &N);

// 执行内核
size_t global_size = N;
size_t local_size = 256;
clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &global_size, &local_size, 0, NULL, NULL);

// 读回结果
clEnqueueReadBuffer(queue, buf_out, CL_TRUE, 0, N * sizeof(float), host_out, 0, NULL, NULL);

// 验证
for (int i = 0; i < 10; i++) {
    printf("out[%d] = %.1f (expected %.1f)\n", i, host_out[i], host_a[i] + host_b[i]);
}

// vector_add.c - 完整的 OpenCL 向量加法
#define CL_TARGET_OPENCL_VERSION 120
#include <CL/cl.h>
#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#include <math.h>

const char *kernel_source =
"__kernel void vector_add(\n"
"    __global const float *a,\n"
"    __global const float *b,\n"
"    __global float *out,\n"
"    const int n)\n"
"{\n"
"    int gid = get_global_id(0);\n"
"    if (gid < n) {\n"
"        out[gid] = a[gid] + b[gid];\n"
"    }\n"
"}\n";

int main() {
    const int N = 1024 * 1024;
    cl_int err;

    // 1. 获取平台和设备
    cl_platform_id platform;
    cl_device_id device;
    clGetPlatformIDs(1, &platform, NULL);
    clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 1, &device, NULL);

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

    // 3. 编译内核
    cl_program program = clCreateProgramWithSource(context, 1, &kernel_source, NULL, &err);
    clBuildProgram(program, 1, &device, NULL, NULL, NULL);
    cl_kernel kernel = clCreateKernel(program, "vector_add", &err);

    // 4. 分配主机内存
    float *a = malloc(N * sizeof(float));
    float *b = malloc(N * sizeof(float));
    float *out = malloc(N * sizeof(float));
    for (int i = 0; i < N; i++) { a[i] = i; b[i] = i * 2; }

    // 5. 创建设备缓冲区并传输数据
    cl_mem buf_a = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
                                   N * sizeof(float), a, &err);
    cl_mem buf_b = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
                                   N * sizeof(float), b, &err);
    cl_mem buf_out = clCreateBuffer(context, CL_MEM_WRITE_ONLY, N * sizeof(float), NULL, &err);

    // 6. 设置参数并执行
    clSetKernelArg(kernel, 0, sizeof(cl_mem), &buf_a);
    clSetKernelArg(kernel, 1, sizeof(cl_mem), &buf_b);
    clSetKernelArg(kernel, 2, sizeof(cl_mem), &buf_out);
    clSetKernelArg(kernel, 3, sizeof(int), &N);

    size_t global_size = N;
    size_t local_size = 256;
    clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &global_size, &local_size, 0, NULL, NULL);
    clFinish(queue);

    // 7. 读回结果
    clEnqueueReadBuffer(queue, buf_out, CL_TRUE, 0, N * sizeof(float), out, 0, NULL, NULL);

    // 8. 验证
    int errors = 0;
    for (int i = 0; i < N; i++) {
        if (fabs(out[i] - (a[i] + b[i])) > 1e-5) errors++;
    }
    printf("Completed with %d errors out of %d elements\n", errors, N);

    // 9. 清理
    clReleaseMemObject(buf_a);
    clReleaseMemObject(buf_b);
    clReleaseMemObject(buf_out);
    clReleaseKernel(kernel);
    clReleaseProgram(program);
    clReleaseCommandQueue(queue);
    clReleaseContext(context);
    free(a); free(b); free(out);
    return 0;
}

gcc vector_add.c -o vector_add -lOpenCL -lm
./vector_add
# 预期输出: Completed with 0 errors out of 1048576 elements

// 错误码宏
#define CL_CHECK(err, msg) do { \
    if ((err) != CL_SUCCESS) { \
        fprintf(stderr, "OpenCL Error (%d) at %s:%d: %s\n", err, __FILE__, __LINE__, msg); \
        exit(1); \
    } \
} while(0)

// 使用
cl_int err;
cl_mem buf = clCreateBuffer(context, CL_MEM_READ_WRITE, size, NULL, &err);
CL_CHECK(err, "clCreateBuffer failed");