opencl-minimal-demo

kernel.cl

__kernel void add_one(__global int* data) {
    int idx = get_global_id(0);
    data[idx] += 1;
}

main.c

// generate by chatgpt
#define CL_TARGET_OPENCL_VERSION 300

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

#define MAX_SOURCE_SIZE (0x100000)

// 错误检查函数
void checkError(cl_int err, const char *operation) {
  if (err != CL_SUCCESS) {
    fprintf(stderr, "Error during operation '%s': %d\n", operation, err);
    exit(1);
  }
}

// 数据初始化函数
void initData(int **data, size_t size) {
  *data = (int *)malloc(sizeof(int) * size);
  for (size_t i = 0; i < size; i++) {
    (*data)[i] = i;
  }
}

// 打印数组内容
void printArray(const char *message, int *data, size_t size) {
  printf("%s: ", message);
  for (size_t i = 0; i < size; i++) {
    printf("%d ", data[i]);
    if (i >= 9) { // 仅打印前10个元素
      printf("... ");
      break;
    }
  }
  printf("\n");
}

// 设置OpenCL环境
void setupOpenCL(cl_platform_id *platform_id, cl_device_id *device_id,
                 cl_context *context, cl_command_queue *command_queue) {
  cl_int ret;
  ret = clGetPlatformIDs(1, platform_id, NULL);
  checkError(ret, "Getting platform IDs");

  ret =
      clGetDeviceIDs(*platform_id, CL_DEVICE_TYPE_DEFAULT, 1, device_id, NULL);
  checkError(ret, "Getting device IDs");

  *context = clCreateContext(NULL, 1, device_id, NULL, NULL, &ret);
  checkError(ret, "Creating context");

  // 使用 clCreateCommandQueueWithProperties 替代已弃用的 clCreateCommandQueue
  cl_command_queue_properties properties[] = {CL_QUEUE_PROPERTIES,
                                              CL_QUEUE_PROFILING_ENABLE, 0};
  *command_queue = clCreateCommandQueueWithProperties(*context, *device_id,
                                                      properties, &ret);
  checkError(ret, "Creating command queue with properties");

  // 输出设备信息
  char device_name[128];
  ret = clGetDeviceInfo(*device_id, CL_DEVICE_NAME, 128, device_name, NULL);
  checkError(ret, "Getting device name");
  printf("Using device: %s\n", device_name);
}

// 创建和编译OpenCL程序
void createAndCompileProgram(const char *fileName, cl_context context,
                             cl_device_id device_id, cl_program *program) {
  cl_int ret;
  FILE *fp;
  char *source_str;
  size_t source_size;

  fp = fopen(fileName, "r");
  if (!fp) {
    fprintf(stderr, "Failed to load kernel.\n");
    exit(1);
  }
  source_str = (char *)malloc(MAX_SOURCE_SIZE);
  source_size = fread(source_str, 1, MAX_SOURCE_SIZE, fp);
  fclose(fp);

  *program = clCreateProgramWithSource(context, 1, (const char **)&source_str,
                                       (const size_t *)&source_size, &ret);
  checkError(ret, "Creating program");

  ret = clBuildProgram(*program, 1, &device_id, NULL, NULL, NULL);
  checkError(ret, "Building program");

  free(source_str);
}

// 运行OpenCL内核
void runKernel(cl_program program, cl_command_queue command_queue,
               cl_mem data_mem_obj, size_t list_size, cl_device_id device_id) {
  cl_int ret;
  cl_kernel kernel = clCreateKernel(program, "add_one", &ret);
  checkError(ret, "Creating kernel");

  // 查询设备的最大工作组大小
  size_t max_work_group_size;
  ret = clGetKernelWorkGroupInfo(kernel, device_id, CL_KERNEL_WORK_GROUP_SIZE,
                                 sizeof(max_work_group_size),
                                 &max_work_group_size, NULL);
  checkError(ret, "Getting kernel work group info");

  // 确保局部工作组大小不大于最大值
  size_t local_item_size =
      (max_work_group_size > 64) ? 64 : max_work_group_size;

  ret = clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&data_mem_obj);
  checkError(ret, "Setting kernel arguments");

  size_t global_item_size = list_size;
  // 确保全局工作组大小是局部工作组大小的整数倍
  if (global_item_size % local_item_size != 0) {
    global_item_size =
        (global_item_size / local_item_size + 1) * local_item_size;
  }
  printf("global_item_size=%d\n", global_item_size);

  ret =
      clEnqueueNDRangeKernel(command_queue, kernel, 1, NULL, &global_item_size,
                             &local_item_size, 0, NULL, NULL);
  checkError(ret, "Enqueuing kernel");

  ret = clFinish(command_queue);
  checkError(ret, "Finishing command queue");

  clReleaseKernel(kernel);
}

// 清理OpenCL资源
void cleanup(cl_program program, cl_mem data_mem_obj,
             cl_command_queue command_queue, cl_context context, int *data) {
  clReleaseProgram(program);
  clReleaseMemObject(data_mem_obj);
  clReleaseCommandQueue(command_queue);
  clReleaseContext(context);
  free(data);
}

int main() {
  // 数据初始化
  const size_t LIST_SIZE = 1000;
  int *data;
  initData(&data, LIST_SIZE);

  // 打印原始数组值
  printArray("Original array", data, LIST_SIZE);

  // 设置OpenCL环境
  cl_platform_id platform_id = NULL;
  cl_device_id device_id = NULL;
  cl_context context = NULL;
  cl_command_queue command_queue = NULL;
  setupOpenCL(&platform_id, &device_id, &context, &command_queue);

  // 创建内存缓冲区
  cl_int ret;
  cl_mem data_mem_obj = clCreateBuffer(context, CL_MEM_READ_WRITE,
                                       LIST_SIZE * sizeof(int), NULL, &ret);
  checkError(ret, "Creating buffer");

  ret = clEnqueueWriteBuffer(command_queue, data_mem_obj, CL_TRUE, 0,
                             LIST_SIZE * sizeof(int), data, 0, NULL, NULL);
  checkError(ret, "Writing to buffer");

  // 创建和编译程序
  cl_program program = NULL;
  createAndCompileProgram("kernel.cl", context, device_id, &program);

  // 运行内核
  runKernel(program, command_queue, data_mem_obj, LIST_SIZE, device_id);

  // 将结果从缓冲区传回到主机内存
  ret = clEnqueueReadBuffer(command_queue, data_mem_obj, CL_TRUE, 0,
                            LIST_SIZE * sizeof(int), data, 0, NULL, NULL);
  checkError(ret, "Reading from buffer");

  // 打印修改后的数组值
  printArray("Modified array", data, LIST_SIZE);

  // 清理资源
  cleanup(program, data_mem_obj, command_queue, context, data);

  return 0;
}

test.sh

#!/usr/bin/env bash

gcc -o my_program main.c -lOpenCL
./my_program

# use clinfo can check opencl devices