创建和构建程序
要创建程序对象，可以传入OpenCL C源代码文本，或者利用程序二进制码来创建。由OpenCL C源代码创建程序对象是开发人员创建程序对象的一般做法。OpenCL C程序的源代码放在一个外部文件中（例如，就像我们的示例代码中的.cl文件），应用程序将使用clcreateProgramwithsource()函数由源代码创建程序对象，另一种做法是由为设备预编译的二进制码创建程序对象，这个方法将在后面讨论。现在我们先来介绍如何使用clcreateProgramwithSource(从源代码创建程序对象：
cl_program clCreateProgramWithSource(cl_context context,
cl_uint count,
const char **strings,
const size_t *lengths,
cl_int *errcode_ret)
/*
context 创建对象的上下文
count strings 参数中字符串指针的数目
strings 包含count个字符串的指针。将这个参数中包含的所有字符串结合在一起，就构成了创建程序对象的完整的源代码
lengths 这是一个大小为count的数组，包含strings中各个元素的字符数。这个参数可以为NULL,在这种情况下，则认为字符串是以null终止的
errcode_ret 如果非NULL，函数返回的错误代码将由这个参数返回
*/
调用clCreateProgramwWithSource()会使用传入的源代码创建一个新的程序对象。返回值是与上下文关联的一个新程序对象。一般地，调用clCreateProgramWithSource()之后，下一步就是使用clBuildProgram()构建这个程序对象：
cl_int clBuildProgram(cl_program program,
cl_uint num_devices,
const cl_device_id *device_list,
const char *options,
void(CL_CALLBACK *pfn_notify)
(cl_program program,
void *user_data),
void *user_data)
/*
program 一个合法的程序对象
num_devices 要构建程序对象的设备数
device_list 这是一个数组，包含将构建程序的所有num_devices的设备ID.
如果device_list为NULL,则为创建程序对象的上下文中所创建的设备构建这个程序对象
options 这是一个字符串，包含程序的构建选项
pfn_notify 使用pfn_notify参数可以完成异步构建。如果pfn_notify为NULL,那么完成构建之前clBuildProgram不会返回到调用者。
不过，如果用户传入pfn_notify，那么clBuildProgram可以在完成构建之前返回，并在程序构建过程中调用pfn_notify.
这个参数的一个用法是让所有构建排队，使之异步完成，与此同时应用程序可以完成其他工作。
不过需要说明的是，即使传入了pfn_notify,OpenCL实现仍可能选择以同步方式返回。
如果确实需要异步构建应用程序，保证异步执行最可靠的方法是在单独的应用程序线程中执行构建
user_data 如果pfn_notify不为NULL,那么这是作为pfn_notify的参数传入的任意数据
*/
调用clBuildProgram()会为指定的所有设备构建程序对象（或者，如果没有指定设备列表，则为与上下文关联的所有设备构建程序对象)。这一步实际上等价于在一个C程序上调用编译器/链接器。options参数包含一个构建选项字符串，包括预处理器定义和各种优化和代码生成选项（例如，– DUSE_FEATURE=1-cl-mad-enable)。这些选项将在最后详细说明。可执行代码存储在所有目标设备的程序对象中。如果程序在所有设备上成功构建，则clBuildProgram()函数返回CL_SUCCESS；否则，会返回一个错误码。如果存在一个构建错误，则可以调用clGetProgramBuildInfo()并指定param_name为CL_PROGRAM_BUILD_LOG来查看详细的构建日志。
cl_int clGetProgramBuildInfo(cl_program program,
cl_device_id device,
cl_program_build_info param_name,
size_t param_value_size,
void *param_value,
size_t *param_value_size_ret)
/*
program 一个合法的程序对象
device 要获取相应构建信息的设备。这必须是要构建程序的设备之一。
程序要为指定的设备构建，对于不同的设备可能会有不同的错误，所以必须单独查询日志
param_name 要查询的参数。接受以下参数：
CL_PROGRAM_BUILD_STATUS(cl_build_status)返回构建的状态，可以是以下状态：
CL_BUILD_NONE: 未完成任何构建
CL_BUILD_ERROR: 最后一次构建有一个错误
CL_BUILD_SUCCESS: 最后一次构建成功
CL_BUILD_IN_PROGRESS: 还在运行一个异步构建。如果为clBuildProgram 提供了一个函数指针，则可能出现这个状态。
CL_PROGRAM_BUILD_OPTIONS(char[]): 返回一个字符串,其中包含传递给clBuildProgram的options参数。
CL_PROGRAM_BUILD_LOG(char[]): 返回一个字符串,其中包含为这个设备最后一次构建的构建日志。
param_value_size param_value的字节数,这必须足够大,足以存储指定查询的结果。param_value 存储查询结果的内存位置的指针。
param_value_size_ret 实际复制到param_value的字节数。
*/
综合以上内容，下面代码展示了如何从源代码创建一个程序对象，为所有关联的设备构建这个程序对象，并为一个设备查询构建结果。
cl_program CreateProgram(cl_context context, cl_device_id device, const char* fileName)
{
cl_int errNum;
cl_program program;
std::ifstream kernelFile(fileName, std::ios::in);
if (!kernelFile.is_open())
{
std::cerr << "Failed to open file for reading: " << fileName << std::endl;
return NULL;
}
std::ostringstream oss;
oss << kernelFile.rdbuf();
std::string srcStdStr = oss.str();
const char* srcStr = srcStdStr.c_str();
program = clCreateProgramWithSource(context, 1,
(const char**)&srcStr,
NULL, NULL);
if (program == NULL)
{
std::cerr << "Failed to create CL program from source." << std::endl;
return NULL;
}
errNum = clBuildProgram(program, 0, NULL, NULL, NULL, NULL);
if (errNum != CL_SUCCESS)
{
// Determine the reason for the error
char buildLog[16384];
clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG,
sizeof(buildLog), buildLog, NULL);
std::cerr << "Error in kernel: " << std::endl;
std::cerr << buildLog;
clReleaseProgram(program);
return NULL;
}
return program;
}
取自于https://blog.csdn.net/qq_36314864/article/details/130606515里面HelloWorld代码
程序构建选项
本节前面提到，clBuildProgram()取一个字符串参数（const char *options)，控制多种不同类型的构建选项：
1）预处理器选项。
2）浮点数选项（数学性质)。·
3）优化选项。
4）杂项选项。
与C或C++编译器类似，OpenCL有大量选项来控制程序编译的行为。OpenCL程序编译器有一个预处理器，可以在clBuildProgram()的 options参数中定义预处理器的选项。表6-1列出了可以为预处理器指定的选项。
关于定义预处理器变量，需要说明的是，对于构建程序的所有设备，程序对象的内核函数签名必须相同。例如，有以下内核源代码：
#ifdef SOME_MACRO
__kernel void my_kernel(__global const float* p)
{
//...
}
#else //!SOME_MACRO
__kernel void my_kernel(__global const int* p)
{
//...
}
#endif //!SOME_MACRO
在这个例子中，my_kernel()函数签名根据SOME_MACRO的值会有所不同（其参数是__global const float*或__global const int float*）。这本身并不成问题。不过，如果为同一个程序对象的各个设备单独调用clBuildProgram()，一旦为一个设备传人-DSOME_MACRO，如果没有为另一个设备定义SOME_MACRO，得到的内核在程序中会有不同的函数签名，这会失败。也就是说，对于构建一个程序对象的所有设备，内核函数签名必须相同。可以传入不同的预处理器指令，从而对各个设备以不同方式影响程序构建，这是可以接受的，但是不能改变内核函数签名。对于构建一个程序对象的各个设备，内核函数签名必须完全相同。
OpenCL程序编译器还有一些选项来控制浮点数学运算的行为。这些选项见表6-2。与预处理器选项类似，这些选项可以在clBuildProgram()的options参数中指定。
还可以控制OpenCL C编译器运行完成的优化
最后，列出了OpenCL C编译器接受的最后一组杂项选项。
由二进制码创建程序
除了从源代码创建程序对象外，还有一种方法是从二进制码创建程序对象。程序二进制码是对应一个指定设备的源代码的编译版本。程序二进制码的数据格式是不透明的。也就是说,二进制码的内容没有标准化的格式。一个OpenCL实现可能选择将程序的可执行版本存储在二进制码中，也可能选择存储一个中间表示，这个中间表示可以在运行时转换为可执行版本。
由于程序二进制码已经经过编译（可能部分编译为中间表示，或者完全编译为一个可执行文件），所以加载时速度会更快，需要的内存更少，从而可以减少应用程序的加载时间。使用程序二进制码的另一个好处是保护知识产权：可以在安装时生成程序二进制码，而不用将原始OpenCL C源代码存储在磁盘上。一般情况是在安装时生成程序二进制码，或者首先运行并在磁盘上存储二进制码以备以后加载。生成程序二进制码的做法是，首先使用OpenCL从源代码构建程序，然后查询得到的程序二进制码。要从一个已构建的程序得到程序的二进制码，可以使用clGetProgramInfo()：
cl_int clGetProgramInfo(cl_program program,
cl_program_info param_name,
size_t param_value_size,
void *param_value,
size_t *param_value_size_ret)
program 一个合法的程序对象
param_name 要向程序查询哪一个参数的信息，可以接受以下参数：
CL_PROGRAM_REFERENCE_COUNT(cl_uint): 程序的引用数.这可以用于标识是否存在资源泄露
CL_PROGRAM_CONTEXT(cl_context):程序关联到哪个上下文
CL_PROGRAM_NUM_DEVICES(cl_uint):程序关联的设备数
CL_PROGRAM_DEVICES(cl_device_id[]):返回一个cl_device_id数组，其中包含程序关联的设备的ID
CL_PROGRAM_SOURCE(char[]): 将创建程序所用的所有源字符串连接到一个字符串中返回。如果程序对象由二进制码创建，则不会返回任何字符。
CL_PROGRAM_BINARY_SIZES(size_t[]): 返回一个size_t数组，其大小等于与程序关联的设备数。每个元素分别是各设备二进制码的大小。
CL_PROGRAM_BINARIES(unsigned char*[]): 返回一个unsigned char*数组，其中各个元素包含相应设备的程序二进制码。各个数组的大小可以由CL_PROGRAM_BINARY_SIZES查询的结果确定。
param_value_size param_value的字节数
param_value 存储结果的位置指针。这个位置必须分配足够的字节来存储请求的结果。
param_value_size_ret 实际写至param_value的字节数
查询得到程序对象的二进制码之后，可以将这个二进制码存储在磁盘上以备将来运行。下一次运行程序时，可以使用clCreateProgramwithBinary()创建程序对象：
cl_program clCreateProgramWithBinary(cl_context context,
cl_uint num_devices,
const cl_device_id *device_list,
const unsigned char **binaries,
cl_int *binary_status,
cl_int *errcode_ret)
context 创建程序对象的上下文
num_devices 构建程序对象的设备数
device_list 这是一个数组，其中包含将构建程序的所有num_devices的设备ID.
如果device_list为NULL,将为创建程序对象的上下文中创建的所有设备构建程序对象
lengths 这是大小为count的一个数组，包含binaries中各元素的字节数
binaries 这是一个指针数组，所指向的字节包含对应各个设备的各个程序二进制码。
各个二进制码的大小必须为传入lengths中相关元素的大小
binary_status 这个数组包含是否成功加载各个设备二进制码的结果。
如果成功，各元素会设置为CL_SUCCESS;如果失败，会报告一个错误码
errode_ret 如果这个参数不为NULL,则函数返回的错误代码将由这个参数返回
HelloBinaryWorld示例展示了如何由二进制码创建一个程序。这里对之前 HelloWorld示例稍做修改。两者的区别是，这一章的HelloBinaryWorld示例会在应用首次运行时尝试获取程序二进制码，并把它存储到HelloWorld cl. bin。将来执行时，应用程序会从生成的这个二进制码加载程序。完成这个缓存的主要逻辑在代码清单6-2中给出，这个代码取自HelloBinaryWorld 的main()函数。
首次运行时，缓存程序二进制码
program = CreateProgramFromBinary(context, device, "HelloBinaryWorld.cl.bin");
if (program == NULL)
{
std::cout << "Binary not loaded, create from source..." << std::endl;
program = CreateProgram(context, device, "HelloWorld.cl");
if (program == NULL)
{
Cleanup(context, commandQueue, program, kernel, memObjects);
return 1;
}
std::cout << "Save program binary for future run..." << std::endl;
if (SaveProgramBinary(program, device, "HelloWorld.cl.bin") == false)
{
std::cerr << "Failed to write program binary" << std::endl;
Cleanup(context, commandQueue, program, kernel, memObjects);
return 1;
}
}
else
{
std::cout << "Read program from binary." << std::endl;
}
首先来看SaveProgramBinary()，这是查询和存储程序二进制码的函数。这个函数假设已经由源代码创建并构建了程序对象。SaveProgramBinary()的代码见代码清单6-3。这个函数首先调用clGetProgramInfo()查询与程序关联的设备数。接下来获取与各个设备关联的设备ID。得到设备列表之后，这个函数再获取对应各个设备的各个程序二进制码的大小以及程序二进码本身。获取所有程序二进码之后，函数会循环处理设备，找出作为SaveProgramBinary()的参数传入的设备。这个程序二进制码最终会使用fwrite()写至磁盘上的文件Helloworld.cl.bin。
bool SaveProgramBinary(cl_program program, cl_device_id device, const char* fileName)
{
cl_uint numDevices = 0;
cl_int errNum;
// 1 - Query for number of devices attached to program
errNum = clGetProgramInfo(program, CL_PROGRAM_NUM_DEVICES, sizeof(cl_uint),
&numDevices, NULL);
if (errNum != CL_SUCCESS)
{
std::cerr << "Error querying for number of devices." << std::endl;
return false;
}
// 2 - Get all of the Device IDs
cl_device_id* devices = new cl_device_id[numDevices];
errNum = clGetProgramInfo(program, CL_PROGRAM_DEVICES,
sizeof(cl_device_id) * numDevices,
devices, NULL);
if (errNum != CL_SUCCESS)
{
std::cerr << "Error querying for devices." << std::endl;
delete[] devices;
return false;
}
// 3 - Determine the size of each program binary
size_t* programBinarySizes = new size_t[numDevices];
errNum = clGetProgramInfo(program, CL_PROGRAM_BINARY_SIZES,
sizeof(size_t) * numDevices,
programBinarySizes, NULL);
if (errNum != CL_SUCCESS)
{
std::cerr << "Error querying for program binary sizes." << std::endl;
delete[] devices;
delete[] programBinarySizes;
return false;
}
unsigned char** programBinaries = new unsigned char* [numDevices];
for (cl_uint i = 0; i < numDevices; i++)
{
programBinaries[i] = new unsigned char[programBinarySizes[i]];
}
// 4 - Get all of the program binaries
errNum = clGetProgramInfo(program, CL_PROGRAM_BINARIES, sizeof(unsigned char*) * numDevices,
programBinaries, NULL);
if (errNum != CL_SUCCESS)
{
std::cerr << "Error querying for program binaries" << std::endl;
delete[] devices;
delete[] programBinarySizes;
for (cl_uint i = 0; i < numDevices; i++)
{
delete[] programBinaries[i];
}
delete[] programBinaries;
return false;
}
// 5 - Finally store the binaries for the device requested out to disk for future reading.
for (cl_uint i = 0; i < numDevices; i++)
{
// Store the binary just for the device requested. In a scenario where
// multiple devices were being used you would save all of the binaries out here.
if (devices[i] == device)
{
FILE* fp = fopen(fileName, "wb");
fwrite(programBinaries[i], 1, programBinarySizes[i], fp);
fclose(fp);
break;
}
}
// Cleanup
delete[] devices;
delete[] programBinarySizes;
for (cl_uint i = 0; i < numDevices; i++)
{
delete[] programBinaries[i];
}
delete[] programBinaries;
return true;
}
关于程序二进制码，开发人员还要了解一些重要的问题。首先，一个程序二进制码只对创建它的相应设备有效。OpenCL实现可能会选择用其二进制格式存储程序的一个中间表示，也可能存储可执行代码。这是实现所做的选择，应用程序无从得知。因此，如果假设一个二进制码可以在其他设备上工作，那么这是不安全的，除非OpenCL开发商明确做出保证。一般地，要为新设备重新编译二进制码，这对于确保兼容性非常重要。
代码清单6-4给出了一个程序二进制码的例子，这是NVIDIA OpenCL实现生成的。这个代码清单对于熟悉CUDA 的开发人员来说并不陌生。NVIDIA二进制格式采用专用PTX格式存储。Apple和AMD也采用其自己的格式存储二进制码。这些二进制码都无法做到对多个开发商兼容。尽管PTX格式恰好是可读文本，但程序二进制码完全可以是人不可读的二进制位。
![Json 数据通用提取工具，使用 HTML JavaScript [单文件离线版]](https://img-blog.csdnimg.cn/8a832c0a11bb49f48f171daeb8a48d05.png)
