第10章 Compiler - 编译器测试

Compiler 测试套件验证 OpenCL 编译器的核心功能，包括在线编译、链接、编译选项、预处理器、错误处理等。这是确保 OpenCL 程序能够正确编译和执行的关键测试。

10.1 编译器测试概述

10.1.1 OpenCL 编译模型

OpenCL 支持两种编译模型：

1. 传统编译模型（OpenCL 1.x）

源代码 → clBuildProgram → 可执行程序 → clCreateKernel → 内核对象

2. 分离编译链接模型（OpenCL 2.0+）

源代码 → clCompileProgram → 编译对象 (.o)
                              ↓
多个编译对象 → clLinkProgram → 可执行程序 → clCreateKernel → 内核对象

10.1.2 测试分类

测试类别	测试数量	主要功能
程序加载测试	7	源代码加载、多字符串源、程序信息查询
编译选项测试	5	优化选项、宏定义、包含目录、非规格化数缓存
预处理器测试	4	#define、#include、#line、#pragma
版本与扩展测试	3	OpenCL C 版本、扩展宏、特性宏
分离编译链接测试	24	独立编译、静态链接、嵌入式头文件、多文件链接
编译器卸载测试	7	编译器资源卸载、多线程卸载

10.1.3 核心测试列表

test_definition test_list[] = {
    // 程序加载
    ADD_TEST(load_program_source),
    ADD_TEST(load_multistring_source),
    ADD_TEST(load_two_kernel_source),
    ADD_TEST(load_null_terminated_source),
    ADD_TEST(get_program_source),
    ADD_TEST(get_program_build_info),
    ADD_TEST(get_program_info),
    
    // 编译与构建
    ADD_TEST(large_compile),
    ADD_TEST(async_build),
    
    // 编译选项
    ADD_TEST(options_build_optimizations),
    ADD_TEST(options_build_macro),
    ADD_TEST(options_build_macro_existence),
    ADD_TEST(options_include_directory),
    ADD_TEST(options_denorm_cache),
    
    // 预处理器
    ADD_TEST(preprocessor_define_udef),
    ADD_TEST(preprocessor_include),
    ADD_TEST(preprocessor_line_error),
    ADD_TEST(preprocessor_pragma),
    
    // 版本与扩展
    ADD_TEST(opencl_c_versions),
    ADD_TEST(compiler_defines_for_extensions),
    ADD_TEST(image_macro),
    
    // 分离编译与链接
    ADD_TEST(simple_compile_only),
    ADD_TEST(simple_link_only),
    ADD_TEST(execute_after_simple_compile_and_link),
    ADD_TEST(two_file_link),
    ADD_TEST(multiple_files),
    ADD_TEST(multiple_libraries),
    
    // 编译器卸载
    ADD_TEST(unload_valid),
    ADD_TEST(unload_repeated),
    ADD_TEST(unload_build_threaded),
};

---

10.2 在线编译测试

10.2.1 基本程序编译

OpenCL 程序从源代码编译的基本流程：

// 1. 创建程序对象
cl_program program = clCreateProgramWithSource(
    context,
    1,                    // 源字符串数量
    &kernel_source,       // 源代码
    NULL,                 // 自动检测长度
    &error);

// 2. 构建程序
error = clBuildProgram(
    program,
    1,                    // 设备数量
    &device,              // 目标设备
    NULL,                 // 编译选项
    NULL,                 // 回调函数
    NULL);                // 回调数据

// 3. 检查构建状态
if (error != CL_SUCCESS) {
    // 获取构建日志
    size_t log_size;
    clGetProgramBuildInfo(program, device,
                         CL_PROGRAM_BUILD_LOG,
                         0, NULL, &log_size);
    
    char *build_log = (char*)malloc(log_size);
    clGetProgramBuildInfo(program, device,
                         CL_PROGRAM_BUILD_LOG,
                         log_size, build_log, NULL);
    
    log_error("Build failed:\n%s\n", build_log);
    free(build_log);
}

// 4. 创建内核
cl_kernel kernel = clCreateKernel(program, "kernel_name", &error);

10.2.2 多字符串源代码

OpenCL 允许程序由多个字符串组成：

const char *source_strings[] = {
    "// 第一部分：常量定义\n"
    "#define ARRAY_SIZE 1024\n"
    "\n",
    
    "// 第二部分：辅助函数\n"
    "float helper_function(float x) {\n"
    "    return x * 2.0f;\n"
    "}\n"
    "\n",
    
    "// 第三部分：内核函数\n"
    "__kernel void process(__global float *data) {\n"
    "    int id = get_global_id(0);\n"
    "    data[id] = helper_function(data[id]);\n"
    "}\n"
};

// 使用多个字符串创建程序
cl_program program = clCreateProgramWithSource(
    context,
    3,                    // 三个源字符串
    source_strings,
    NULL,
    &error);

测试要点：

• 验证多字符串源代码能够正确拼接
• 确保跨字符串的函数调用正常工作
• 检查预处理器指令在多字符串中的行为

10.2.3 异步编译

OpenCL 支持异步编译，允许在编译过程中执行其他操作：

// 异步编译回调函数
void CL_CALLBACK build_callback(cl_program program, void *user_data) {
    int *build_complete = (int*)user_data;
    *build_complete = 1;
    log_info("Build completed asynchronously\n");
}

// 启动异步编译
int build_complete = 0;
error = clBuildProgram(
    program,
    1, &device,
    NULL,
    build_callback,      // 回调函数
    &build_complete);    // 回调数据

// 等待编译完成
while (!build_complete) {
    // 可以执行其他操作
    usleep(1000);
}

// 或使用事件等待
clWaitForEvents(1, &build_event);

10.2.4 大规模程序编译

测试编译器处理大型程序的能力：

// 生成包含大量代码行的程序
char large_kernel_source[MAX_SIZE];
char *ptr = large_kernel_source;

ptr += sprintf(ptr, "__kernel void large_test(__global float *data) {\n");
ptr += sprintf(ptr, "    int id = get_global_id(0);\n");
ptr += sprintf(ptr, "    float temp = data[id];\n");

// 生成 1000 行计算语句
for (int i = 0; i < 1000; i++) {
    ptr += sprintf(ptr, "    temp = temp * 1.001f + 0.001f;\n");
}

ptr += sprintf(ptr, "    data[id] = temp;\n");
ptr += sprintf(ptr, "}\n");

// 编译大型程序
cl_program program = clCreateProgramWithSource(context, 1,
                                              (const char**)&large_kernel_source,
                                              NULL, &error);
error = clBuildProgram(program, 1, &device, NULL, NULL, NULL);

测试验证：

• 编译器能够处理大型程序
• 编译时间在合理范围内
• 编译后的程序能够正确执行

---

10.3 编译选项测试

10.3.1 优化选项

OpenCL 提供多种编译优化选项：

选项	说明
-cl-opt-disable	禁用所有优化
-cl-mad-enable	允许 a*b+c 融合为 mad 指令
-cl-no-signed-zeros	允许忽略有符号零
-cl-unsafe-math-optimizations	允许不安全的数学优化
-cl-finite-math-only	假设参数和结果不是 NaN 或 Inf
-cl-fast-relaxed-math	快速宽松数学模式
-cl-single-precision-constant	单精度常量
-cl-denorms-are-zero	非规格化数视为零

测试示例：

const char *optimization_options[] = {
    "-cl-opt-disable",
    "-cl-mad-enable",
    "-cl-fast-relaxed-math",
    NULL
};

for (int i = 0; optimization_options[i] != NULL; i++) {
    log_info("Testing option: %s\n", optimization_options[i]);
    
    error = clBuildProgram(program, 1, &device,
                          optimization_options[i],
                          NULL, NULL);
    
    if (error != CL_SUCCESS) {
        log_error("Failed to build with option %s\n",
                 optimization_options[i]);
        return -1;
    }
}

10.3.2 宏定义选项

通过编译选项定义宏：

// 内核源代码使用宏
const char *kernel_with_macro =
"__kernel void test(__global int *output) {\n"
"    output[0] = VALUE;\n"
"}\n";

// 方式1：单个宏定义
error = clBuildProgram(program, 1, &device,
                      "-DVALUE=42",
                      NULL, NULL);

// 方式2：多个宏定义
error = clBuildProgram(program, 1, &device,
                      "-DWIDTH=1024 -DHEIGHT=768",
                      NULL, NULL);

// 方式3：宏展开为表达式
error = clBuildProgram(program, 1, &device,
                      "-DARRAY_SIZE=(1024*1024)",
                      NULL, NULL);

// 验证宏是否生效
cl_int result;
// ... 执行内核并读取 result
if (result != 42) {
    log_error("Macro definition failed\n");
}

宏存在性测试：

const char *existence_test =
"__kernel void test(__global int *output) {\n"
"#ifdef TEST_MACRO\n"
"    output[0] = 1;  // 宏已定义\n"
"#else\n"
"    output[0] = 0;  // 宏未定义\n"
"#endif\n"
"}\n";

// 测试：定义宏
clBuildProgram(program1, 1, &device, "-DTEST_MACRO", NULL, NULL);
// 执行后 output[0] 应该为 1

// 测试：不定义宏
clBuildProgram(program2, 1, &device, NULL, NULL, NULL);
// 执行后 output[0] 应该为 0

10.3.3 包含目录选项

使用 -I 选项指定头文件搜索路径：

// 创建头文件：./include/common.h
const char *header_content =
"#ifndef COMMON_H\n"
"#define COMMON_H\n"
"#define MAGIC_NUMBER 42\n"
"#endif\n";

// 内核源代码包含头文件
const char *kernel_with_include =
"#include \"common.h\"\n"
"__kernel void test(__global int *output) {\n"
"    output[0] = MAGIC_NUMBER;\n"
"}\n";

// 使用 -I 指定包含目录
char build_options[256];
sprintf(build_options, "-I%s/include", test_directory);

error = clBuildProgram(program, 1, &device,
                      build_options,
                      NULL, NULL);

多级目录测试：

// 目录结构：
// ./include/
//   ├── level1.h
//   └── subdir/
//       └── level2.h

// level1.h 包含 level2.h
"#include \"subdir/level2.h\"\n"

// 编译选项
"-I./include"

10.3.4 警告与错误控制

// -w：抑制所有警告
clBuildProgram(program, 1, &device, "-w", NULL, NULL);

// -Werror：将警告视为错误
clBuildProgram(program, 1, &device, "-Werror", NULL, NULL);

// 组合使用
clBuildProgram(program, 1, &device, "-Werror -cl-opt-disable", NULL, NULL);

---

10.4 预处理器测试

10.4.1 #define 和 #undef

const char *define_test =
"#define VALUE 100\n"
"__kernel void test1(__global int *out) {\n"
"    out[0] = VALUE;\n"  // 应该是 100
"}\n"
"#undef VALUE\n"
"#define VALUE 200\n"
"__kernel void test2(__global int *out) {\n"
"    out[0] = VALUE;\n"  // 应该是 200
"}\n";

测试要点：

• #define 正确定义宏
• #undef 正确取消宏定义
• 宏重定义行为
• 宏展开的递归处理

10.4.2 #include 指令

// 测试相对路径包含
"#include \"header.h\"\n"

// 测试绝对路径包含
"#include \"/path/to/header.h\"\n"

// 测试嵌套包含
// file1.h 包含 file2.h
// file2.h 包含 file3.h

// 测试循环包含保护
// header.h:
"#ifndef HEADER_H\n"
"#define HEADER_H\n"
"// 内容\n"
"#endif\n"

10.4.3 #pragma 指令

// 测试 #pragma OPENCL EXTENSION
const char *extension_pragma =
"#pragma OPENCL EXTENSION cl_khr_fp64 : enable\n"
"__kernel void test(__global double *data) {\n"
"    data[0] = 3.14159;\n"
"}\n";

// 测试 #pragma unroll (OpenCL 2.0+)
const char *unroll_pragma =
"__kernel void test(__global int *data) {\n"
"    #pragma unroll 4\n"
"    for (int i = 0; i < 16; i++) {\n"
"        data[i] = i * 2;\n"
"    }\n"
"}\n";

// 测试未知 pragma（应该被忽略）
"#pragma unknown_directive\n"

10.4.4 #line 和 #error

// 测试 #line 指令
const char *line_test =
"#line 100 \"virtual_file.cl\"\n"
"__kernel void test() {\n"
"    int x = UNDEFINED;  // 错误应该报告在 virtual_file.cl:101\n"
"}\n";

// 测试 #error 指令
const char *error_test =
"#ifndef REQUIRED_MACRO\n"
"#error \"REQUIRED_MACRO must be defined\"\n"
"#endif\n";

// 编译应该失败，并输出错误消息

---

10.5 版本与扩展测试

10.5.1 OpenCL C 版本宏

OpenCL C 预定义版本宏：

宏名称	OpenCL 版本	宏值
__OPENCL_VERSION__	所有版本	版本号（如 120, 200, 300）
__OPENCL_C_VERSION__	OpenCL 1.1+	C 语言版本号
CL_VERSION_1_0	1.0+	100
CL_VERSION_1_1	1.1+	110
CL_VERSION_1_2	1.2+	120
CL_VERSION_2_0	2.0+	200
CL_VERSION_3_0	3.0+	300

测试示例：

const char *version_test =
"__kernel void test(__global int *output) {\n"
"#ifdef __OPENCL_VERSION__\n"
"    output[0] = __OPENCL_VERSION__;\n"
"#else\n"
"    output[0] = 0;\n"
"#endif\n"
"#ifdef CL_VERSION_2_0\n"
"    output[1] = 1;  // 支持 OpenCL 2.0\n"
"#else\n"
"    output[1] = 0;  // 不支持\n"
"#endif\n"
"}\n";

10.5.2 扩展宏定义

设备支持的扩展会自动定义相应的宏：

// 查询设备支持的扩展
char extensions[4096];
clGetDeviceInfo(device, CL_DEVICE_EXTENSIONS,
               sizeof(extensions), extensions, NULL);

// 测试扩展宏
const char *extension_macro_test =
"__kernel void test(__global int *output) {\n"
"    int idx = 0;\n"
"#ifdef cl_khr_fp64\n"
"    output[idx++] = 1;  // 支持 双精度\n"
"#endif\n"
"#ifdef cl_khr_3d_image_writes\n"
"    output[idx++] = 2;  // 支持 3D 图像写入\n"
"#endif\n"
"#ifdef cl_khr_gl_sharing\n"
"    output[idx++] = 3;  // 支持 OpenGL 共享\n"
"#endif\n"
"}\n";

10.5.3 图像支持宏

const char *image_macro_test =
"__kernel void test(__global int *output) {\n"
"#ifdef __IMAGE_SUPPORT__\n"
"    output[0] = 1;  // 设备支持图像\n"
"#else\n"
"    output[0] = 0;  // 不支持图像\n"
"#endif\n"
"}\n";

// 验证宏与设备能力一致
cl_bool image_support;
clGetDeviceInfo(device, CL_DEVICE_IMAGE_SUPPORT,
               sizeof(cl_bool), &image_support, NULL);

// 执行内核并检查结果
// output[0] 应该与 image_support 一致

10.5.4 特性宏（OpenCL 3.0+）

OpenCL 3.0 引入可选特性宏：

// __opencl_c_3d_image_writes
// __opencl_c_atomic_order_acq_rel
// __opencl_c_atomic_scope_device
// __opencl_c_generic_address_space
// __opencl_c_pipes
// __opencl_c_work_group_collective_functions

const char *feature_macro_test =
"__kernel void test(__global int *output) {\n"
"#ifdef __opencl_c_atomic_order_acq_rel\n"
"    output[0] = 1;\n"
"#endif\n"
"#ifdef __opencl_c_generic_address_space\n"
"    output[1] = 1;\n"
"#endif\n"
"}\n";

---

10.6 分离编译与链接

10.6.1 基本编译与链接流程

// 步骤1：编译源代码到目标文件
cl_program compiled_program = clCompileProgram(
    program,
    1, &device,
    NULL,                    // 编译选项
    0, NULL, NULL,           // 头文件
    NULL, NULL);             // 回调

// 步骤2：链接目标文件到可执行程序
cl_program linked_program = clLinkProgram(
    context,
    1, &device,
    NULL,                    // 链接选项
    1, &compiled_program,    // 输入程序
    NULL, NULL,              // 回调
    &error);

// 步骤3：创建内核
cl_kernel kernel = clCreateKernel(linked_program, "kernel_name", &error);

10.6.2 多文件编译

// 文件1：辅助函数
const char *helper_source =
"float helper(float x) {\n"
"    return x * 2.0f;\n"
"}\n";

// 文件2：主内核（使用辅助函数）
const char *kernel_source =
"extern float helper(float x);\n"
"__kernel void main_kernel(__global float *data) {\n"
"    int id = get_global_id(0);\n"
"    data[id] = helper(data[id]);\n"
"}\n";

// 分别编译
cl_program helper_program = /* 创建并编译 helper_source */;
cl_program kernel_program = /* 创建并编译 kernel_source */;

// 链接到一起
cl_program programs[] = { helper_program, kernel_program };
cl_program linked = clLinkProgram(context, 1, &device, NULL,
                                 2, programs,
                                 NULL, NULL, &error);

10.6.3 静态库

// 创建静态库
cl_program library = clLinkProgram(
    context,
    1, &device,
    "-create-library",       // 创建库而不是可执行程序
    1, &compiled_program,
    NULL, NULL,
    &error);

// 使用库链接程序
cl_program programs[] = { library, main_program };
cl_program executable = clLinkProgram(context, 1, &device, NULL,
                                     2, programs,
                                     NULL, NULL, &error);

10.6.4 嵌入式头文件

// 创建嵌入式头文件
const char *header_name = "embedded_header.h";
const char *header_source =
"#define CONSTANT_VALUE 42\n"
"typedef struct { int x, y; } Point;\n";

cl_program header_program = clCreateProgramWithSource(
    context, 1, &header_source, NULL, &error);

// 编译时包含嵌入式头文件
error = clCompileProgram(
    main_program,
    1, &device,
    NULL,
    1,                      // 头文件数量
    &header_program,        // 头文件程序
    &header_name,           // 头文件名称
    NULL, NULL);

---

10.7 编译器卸载测试

10.7.1 基本卸载测试

// 编译并使用程序
cl_program program = /* 创建和构建程序 */;
cl_kernel kernel = clCreateKernel(program, "test", &error);

// 执行内核
/* ... */

// 清理
clReleaseKernel(kernel);
clReleaseProgram(program);

// 卸载编译器（释放编译器资源）
error = clUnloadPlatformCompiler(platform);
test_error(error, "Failed to unload compiler");

10.7.2 多次卸载测试

// 测试重复卸载
for (int i = 0; i < 10; i++) {
    error = clUnloadPlatformCompiler(platform);
    if (error != CL_SUCCESS) {
        log_error("Unload failed at iteration %d\n", i);
        return -1;
    }
}

10.7.3 卸载后重新编译

// 1. 编译程序
cl_program program1 = /* 构建程序 */;

// 2. 卸载编译器
clUnloadPlatformCompiler(platform);

// 3. 尝试重新编译（应该成功）
cl_program program2 = clCreateProgramWithSource(context, 1, &source, NULL, &error);
error = clBuildProgram(program2, 1, &device, NULL, NULL, NULL);

// 编译器应该能够重新初始化并正常工作

10.7.4 多线程卸载测试

验证在多线程环境下卸载编译器的安全性。

---

10.8 编译错误处理

10.8.1 获取构建日志

int get_build_log(cl_program program, cl_device_id device)
{
    size_t log_size;
    cl_int error;
    
    // 获取日志大小
    error = clGetProgramBuildInfo(program, device,
                                 CL_PROGRAM_BUILD_LOG,
                                 0, NULL, &log_size);
    if (error != CL_SUCCESS)
        return error;
    
    // 分配缓冲区
    char *build_log = (char*)malloc(log_size + 1);
    if (!build_log)
        return CL_OUT_OF_HOST_MEMORY;
    
    // 获取日志内容
    error = clGetProgramBuildInfo(program, device,
                                 CL_PROGRAM_BUILD_LOG,
                                 log_size, build_log, NULL);
    if (error != CL_SUCCESS) {
        free(build_log);
        return error;
    }
    
    build_log[log_size] = '\0';
    log_info("Build log:\n%s\n", build_log);
    
    free(build_log);
    return CL_SUCCESS;
}

10.8.2 构建状态查询

// 查询构建状态
cl_build_status build_status;
clGetProgramBuildInfo(program, device,
                     CL_PROGRAM_BUILD_STATUS,
                     sizeof(cl_build_status),
                     &build_status, NULL);

switch (build_status) {
    case CL_BUILD_SUCCESS:
        log_info("Build succeeded\n");
        break;
    case CL_BUILD_ERROR:
        log_error("Build failed\n");
        get_build_log(program, device);
        break;
    case CL_BUILD_IN_PROGRESS:
        log_info("Build in progress\n");
        break;
    case CL_BUILD_NONE:
        log_info("Build not started\n");
        break;
}

10.8.3 常见编译错误

错误类型	示例	解决方法
语法错误	缺少分号	检查语法
未定义标识符	使用未声明的变量	声明变量或包含头文件
类型不匹配	float x = 1;	使用 1.0f
扩展未启用	使用 double 但未启用 cl_khr_fp64	添加 #pragma OPENCL EXTENSION
包含文件未找到	#include "missing.h"	使用 -I 指定路径

---

10.9 编译器测试最佳实践

10.9.1 测试策略

1. 渐进式测试：从简单程序开始，逐步增加复杂度
2. 覆盖所有选项：测试所有编译选项组合
3. 错误路径测试：验证错误处理是否正确
4. 性能测试：测量编译时间和编译后性能
5. 跨平台测试：在不同设备和平台上测试

10.9.2 调试技巧

// 1. 启用详细输出
"-cl-nv-verbose"         // NVIDIA
"-cl-amd-verbose"        // AMD

// 2. 保存中间文件
"-save-temps"            // 保存临时文件

// 3. 生成汇编代码
"-cl-nv-asm"            // NVIDIA PTX

10.9.3 注意事项

• 编译选项顺序：某些选项可能相互影响
• 设备差异：不同设备可能支持不同的特性
• 版本兼容性：确保代码兼容目标 OpenCL 版本
• 内存限制：大型程序可能超出编译器限制
• 线程安全：多线程编译时需要注意同步

---

本章小结

第10章详细介绍了 Compiler 测试套件的核心内容：

• 在线编译：程序创建、构建、异步编译、大规模程序处理
• 编译选项：优化选项、宏定义、包含目录、警告控制
• 预处理器：#define、#include、#pragma、#line、#error 指令
• 版本与扩展：版本宏、扩展宏、图像宏、特性宏
• 分离编译链接：独立编译、多文件链接、静态库、嵌入式头文件
• 编译器卸载：资源释放、重复卸载、多线程安全
• 错误处理：构建日志、状态查询、常见错误

编译器测试确保 OpenCL 编译器能够正确处理各种源代码，支持标准规定的所有特性，并提供清晰的错误诊断信息。