▶ 按书上写的管道的代码,需要使用 OpenCL2.0 的平台和设备,目前编译不通过,暂时不知道是什么问题,先把代码堆上来,以后换了新的设备再说
● 程序主要功能:用主机上的数组 srcHost 创建设备缓冲区 src,调用核函数 pipeProducer 将 src 分装到管道中,再调用核函数 pipeConsumer 将管道中的数据读到设备缓冲区 dst 中,最后拷贝回主机数组 dstHost 中检查结果。
● 代码
//pipe.cl
__kernel void pipeProducer(__global float *src, __write_only pipe float outPipe)
{
int gid = get_global_id();
float srcPipe = src[gid];
reserve_id_t resID = reserve_write_pipe(outPipe, );
if (is_valid_reserve_id(resID))
{
if (write_pipe(outPipe, resID, , &srcPipe) != )
return;
commit_write_pipe(outPipe, resID);
}
} __kernel void pipeConsumer(__global float *dst, __read_only pipe float inPipe)
{
int gid = get_global_id();
float dstPipe;
reserve_id_t resID = reserve_read_pipe(inPipe, );
if (is_valid_reserve_id(resID))
{
if (read_pipe(inPipe, resID, , &dstPipe) != )
return;
commit_read_pipe(inPipe, resID);
}
dst[gid] = dstPipe;
}
//main.c
#include <stdio.h>
#include <stdlib.h>
#include <cl.h> const char *sourceCode = "D:/Code/pipe.cl"; char* readSource(const char* kernelPath)// 读取文本文件,存储为 char *
{
FILE *fp;
char *source;
long int size;
//printf("readSource, Program file: %s\n", kernelPath);
fopen_s(&fp, kernelPath, "rb");
if (!fp)
{
printf("Open kernel file failed\n");
exit(-);
}
if (fseek(fp, , SEEK_END) != )
{
printf("Seek end of file faildd\n");
exit(-);
}
if ((size = ftell(fp)) < )
{
printf("Get file position failed\n");
exit(-);
}
rewind(fp);
if ((source = (char *)malloc(size + )) == NULL)
{
printf("Allocate space failed\n");
exit(-);
}
fread(source, , size, fp);
fclose(fp);
source[size] = '\0';
return source;
} int main()
{
const int nPacket = , dataSize = nPacket * sizeof(float);
char info[] = { };
int i; // 初始化平台
cl_int status;
cl_platform_id platform;
clGetPlatformIDs(, &platform, NULL);
cl_device_id device;
clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, , &device, NULL);
cl_context_properties contextProp[] = { CL_CONTEXT_PLATFORM,(cl_context_properties)(platform), };
cl_context context = clCreateContext(contextProp, , &device, NULL, contextProp, &status);
cl_command_queue queue = clCreateCommandQueueWithProperties(context, device, NULL, &status);
cl_event eventProducer, eventConsumer; const char* source = readSource(sourceCode);
cl_program program = clCreateProgramWithSource(context, , &source, NULL, &status);
status = clBuildProgram(program, , &device, "-w -g –cl-std=CL2.0", NULL, NULL); clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, , info, NULL);
printf("Build log:\n%s\n", info); cl_kernel kernelProducer = clCreateKernel(program, "pipeProducer", &status);
cl_kernel kernelConsumer = clCreateKernel(program, "pipeConsumer", &status);
size_t globalSize = nPacket, localSize = ; float *srcHost = (float *)malloc(dataSize);
float *dstHost = (float *)malloc(dataSize);
for (i = ; i < nPacket; srcHost[i] = i, dstHost[i] = 0.0f, i++); cl_mem src, dst;
src = clCreateBuffer(context, CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR, dataSize, srcHost, &status);
dst = clCreateBuffer(context, CL_MEM_WRITE_ONLY, dataSize, NULL, &status); cl_mem pipe = clCreatePipe(context, CL_MEM_HOST_NO_ACCESS, sizeof(float), nPacket, NULL, &status); clSetKernelArg(kernelProducer, , sizeof(cl_mem),src);
clSetKernelArg(kernelProducer, , sizeof(cl_mem), &pipe); clSetKernelArg(kernelProducer, , sizeof(cl_mem), dst);
clSetKernelArg(kernelProducer, , sizeof(cl_mem), &pipe); clEnqueueNDRangeKernel(queue, kernelProducer, , NULL, &globalSize, &localSize, , NULL, &eventProducer);
clEnqueueNDRangeKernel(queue, kernelConsumer, , NULL, &globalSize, &localSize, , &eventProducer, &eventConsumer);
clEnqueueReadBuffer(queue, dst, CL_TRUE, dataSize, dataSize, dstHost, , &eventConsumer, NULL);
clFinish(queue); for (i = ; i < nPacket; i++)
{
if (dstHost[i] != i)
break;
}
printf("Output is %s.\n", (i == nPacket) ? "correct" : "incorrect"); free(srcHost);
free(dstHost);
clReleaseContext(context);
clReleaseCommandQueue(queue);
clReleaseProgram(program);
clReleaseKernel(kernelProducer);
clReleaseKernel(kernelConsumer);
getchar();
return ;
}
● 输出结果
■ 使用编译参数 "-w -g –cl-std=CL2.0" 时返回 status 为 -43(CL_INVALID_BUILD_OPTIONS),不使用参数 "–cl-std=CL2.0" 的情况下返回 -11(CL_BUILD_PROGRAM_FAILURE),麻烦的是调用函数 clGetProgramBuildInfo 查询编译日志 info 始终都是空的,不知道出了什么问题。
■ 转机,去掉了修饰符 __write_only 和 __read_only(只用于图像类型的缓冲区),返回 status 为 -11,至少报错信息有了:【identifier "pipe" is undefined】和【invalid combination of type specifiers】(指在 float 上)
● 后续代码,但是上述代码都编译不了,下面的也暂时没用。(1)使用局部内存来统一工作组的写入
//pipe2.cl
__kernel void pipeProducer(__global float *src, __write_only pipe float outPipe)
{
int gid = get_global_id(), lid = get_local_id();
__local reserve_id_t resID;
if (lid == )
resID = reserve_write_pipe(outPipe, get_local_size()); // 工作组中首个工作项一次预定多个管道位置
barrier(CLK_LOCAL_MEM_FENCE); float srcPipe = src[id];
if (is_valid_reserve_id(resID))
{
if (write_pipe(outPipe, resID, lid, &srcPipe) != ) // 每个工作项写入预定的位置
return;
commit_write_pipe(outPipe, resID);
}
} __kernel void pipeConsumer(__global float *dst, __read_only pipe float inPipe)
{
int gid = get_global_id(), lid = get_local_id();
__local reserve_id_t resID;
if (lid == )
resID = reserve_read_pipe(inPipe, get_local_size());
barrier(CLK_LOCAL_MEM_FENCE); float dstPipe;
if (is_valid_reserve_id(resID))
{
if (read_pipe(inPipe, resID, lid, &dstPipe) != )
return;
commit_read_pipe(inPipe, resID);
}
dst[gid] = dstPipe;
}
● (2)使用工作组管道操作简化上述代码(只是干掉了一个 if 和一个同步)
//pipe3.cl
__kernel void pipeProducer(__global float *src, __write_only pipe float outPipe)
{
int gid = get_global_id(), lid = get_local_id();
__local reserve_id_t resID = work_group_reserve_write_pipe(outPipe, get_local_size());// 自带分支和同步 float srcPipe = src[id];
if (is_valid_reserve_id(resID))
{
if (write_pipe(outPipe, resID, lid, &srcPipe) != )
return;
commit_write_pipe(outPipe, resID);
}
} __kernel void pipeConsumer(__global float *dst, __read_only pipe float inPipe)
{
int gid = get_global_id(), lid = get_local_id();
__local reserve_id_t resID = work_group_reserve_read_pipe(inPipe, get_local_size()); float dstPipe;
if (is_valid_reserve_id(resID))
{
if (read_pipe(inPipe, resID, lid, &dstPipe) != )
return;
commit_read_pipe(inPipe, resID);
}
dst[gid] = dstPipe;
}
● 书上原本的主函数的内容(关于数据缓冲区的部分),是用虚拟内存写的,由于办公室的电脑不支持,上面的代码中被我换成了普通缓冲区
float *src = (float *)clSVMAlloc(context, CL_MEM_READ_WRITE | CL_MEM_SVM_FINE_GRAIN_BUFFER, dataSize, );
float *dst = (float *)clSVMAlloc(context, CL_MEM_READ_WRITE | CL_MEM_SVM_FINE_GRAIN_BUFFER, dataSize, );
if (src == NULL || dst == NULL)
{
printf("clSVMAlloc failed!\n");
getchar();
return ;
} clEnqueueSVMMap(queue, CL_TRUE, CL_MAP_WRITE, src, dataSize, , NULL, NULL);
for (i = ; i < nPacket; i++)
src[i] = i, dst[i] = 0.0f;
clEnqueueSVMUnmap(queue, src, , NULL, NULL); cl_mem pipe = clCreatePipe(context, CL_MEM_HOST_NO_ACCESS, sizeof(float), nPacket, NULL, &status); clSetKernelArgSVMPointer(kernelProducer, , src);
clSetKernelArg(kernelProducer, , sizeof(cl_mem), &pipe); clSetKernelArgSVMPointer(kernelProducer, , dst);
clSetKernelArg(kernelProducer, , sizeof(cl_mem), &pipe); clEnqueueNDRangeKernel(queue, kernelProducer, , NULL, &globalSize, &localSize, , NULL, &eventProducer);
clEnqueueNDRangeKernel(queue, kernelConsumer, , NULL, &globalSize, &localSize, , &eventProducer,NULL);
clFinish(queue); clEnqueueSVMMap(queue, CL_TRUE, CL_MAP_READ, dst, dataSize, , NULL, NULL);
for (i = ; i < nPacket; i++)
{
if (dst[i] != i)
break;
}
printf("Output is %s.\n", (i == nPacket) ? "correct" : "incorrect");
clEnqueueSVMUnmap(queue, dst, , NULL, NULL);