clBuildProgram实现分析

<h1 align="center">调试选项</h1>

1. clBuildProgram() 被调用
   ↓
2. 计算程序哈希: KKILOIDNDGMMLILKELBBFGLPCBNNGOPGKKBPB
   ↓
3. 检查缓存: ~/.cache/pocl/kcache/LF/KKILOIDNDGMMLILKELBBFGLPCBNNGOPGKKBPB/
   ↓
4. 缓存不存在 → 开始编译
   ↓
5. Clang 编译 OpenCL C → LLVM IR
   源代码: __kernel void vector_add(...)
   输出: program.bc (LLVM 位码)
   ↓
6. 保存 program.bc 到缓存
   ↓
7. clEnqueueNDRangeKernel() 被调用时:
   ↓
8. 计算内核哈希 + 工作组配置
   ↓
9. 编译 program.bc → vector_add.so
   - 应用工作组专用化
   - 生成 x86-64 机器码
   - 输出 ELF 共享库
   ↓
10. dlopen() 加载 vector_add.so
    ↓
11. dlsym() 查找 "_pocl_kernel_vector_add_workgroup"
    ↓
12. pthread 调度器执行内核
    - 创建工作线程
    - 每个线程调用内核函数
    - 并行处理工作组

CL_API_ENTRY cl_int CL_API_CALL
POname(clBuildProgram)(cl_program program,
                       cl_uint num_devices,
                       const cl_device_id *device_list,
                       const char *options,
                       void (CL_CALLBACK *pfn_notify)(...),
                       void *user_data)
{
  // 调用统一的编译和链接函数
  return compile_and_link_program(
      1,              // compile_program = 1
      1,              // link_program = 1
      program,
      num_devices, device_list, options,
      0, NULL, NULL, 0, NULL,
      pfn_notify, user_data);
}

int compile_and_link_program(...)
{
  // 1. 解析编译选项
  process_options(options, ...);

  // 2. 为每个设备编译
  for (device_i = 0; device_i < program->num_devices; ++device_i)
    {
      cl_device_id device = program->devices[device_i];

      // 3. 从源代码编译
      if (program->source) {
        error = device->ops->build_source(
            program, device_i, 
            num_input_headers, input_headers,
            header_include_names, link_program);
      }

      // 4. 或从二进制/SPIR-V编译
      else {
        error = device->ops->build_binary(...);
      }
    }

  // 5. 设置内核元数据
  setup_kernel_metadata(program);
  setup_device_kernel_hashes(program);

  return CL_SUCCESS;
}

int pocl_driver_build_source(cl_program program, cl_uint device_i, ...)
{
  // 调用 LLVM 编译管线
  return pocl_llvm_build_program(program, device_i,
                                  num_input_headers,
                                  input_headers,
                                  header_include_names,
                                  link_builtin_lib);
}

int pocl_llvm_build_program(cl_program program, unsigned device_i, ...)
{
  // 1. 准备 Clang 编译器参数
  std::stringstream ss;
  ss << "-xcl ";                    // OpenCL C 语言
  ss << "-cl-kernel-arg-info ";     // 内核参数信息
  ss << "-D__OPENCL_VERSION__=" << device->version_as_int;
  ss << user_options;               // 用户选项

  // 2. 创建 Clang 编译器实例
  CompilerInstance CI;
  CompilerInvocation &pocl_build = CI.getInvocation();

  // 3. 将源代码写入临时文件
  pocl_cache_tempname(tempfile, ".cl", &fd);
  write(fd, program->source, strlen(program->source));

  // 4. 编译 OpenCL C → LLVM IR (.bc)
  EmitLLVMOnlyAction EmitLLVMAction(&llvm_ctx->Context);
  CI.ExecuteAction(EmitLLVMAction);

  // 5. 获取 LLVM Module
  mod = EmitLLVMAction.takeModule().release();

  // 6. 链接设备专用库（如 libclc）
  pocl_llvm_link_builtin_libs(program, device_i, mod);

  // 7. 运行优化 Pass
  pocl_llvm_optimize_program(mod, device);

  // 8. 保存 program.bc 到缓存
  pocl_cache_program_bc_path(program_bc_path, program, device_i);
  WriteBitcodeToFile(*mod, program_bc_path);

  return CL_SUCCESS;
}

~/.cache/pocl/kcache/
├── LF/                              # 哈希值的前2个字符
│   └── KKILOIDNDGMMLILKELBBFGLPCBNNGOPGKKBPB/  # 程序构建哈希 (SHA1)
│       ├── program.bc               # ← 编译后的 LLVM 位码
│       ├── last_accessed            # 最后访问时间
│       └── vector_add/              # 内核名称
│           └── 64-1-1-goffs0-smallgrid/  # 工作组配置
│               └── vector_add.so    # ← 最终的本地二进制！

// 来自 lib/CL/devices/basic/basic.c
char* pocl_basic_build_hash(cl_device_id device)
{
  char* res = calloc(1000, sizeof(char));
  snprintf(res, 1000, "cpu-minimal-%s-%s", 
           HOST_DEVICE_BUILD_HASH,  // 设备特征哈希
           device->llvm_cpu);        // CPU 型号
  return res;
}

64-1-1           → 本地工作组大小 (local_size[0]=64, [1]=1, [2]=1)
goffs0           → 全局偏移量为0
smallgrid        → 小网格优化（max_grid_width < 设备限制）