OpenCL 管道

▶ 按書上寫的管道的代碼，須要使用 OpenCL2.0 的平臺和設備，目前編譯不經過，暫時不知道是什麼問題，先把代碼堆上來，之後換了新的設備再說

● 程序主要功能：用主機上的數組 srcHost 建立設備緩衝區 src，調用核函數 pipeProducer 將 src 分裝到管道中，再調用核函數 pipeConsumer 將管道中的數據讀到設備緩衝區 dst 中，最後拷貝回主機數組 dstHost 中檢查結果。

● 代碼

 1 //pipe.cl
 2 __kernel void pipeProducer(__global float *src, __write_only pipe float outPipe)
 3 {
 4     int gid = get_global_id(0);
 5     float srcPipe = src[gid];
 6     reserve_id_t resID = reserve_write_pipe(outPipe, 1);
 7     if (is_valid_reserve_id(resID))
 8     {
 9         if (write_pipe(outPipe, resID, 0, &srcPipe) != 0)
10             return;
11         commit_write_pipe(outPipe, resID);
12     }
13 }
14 
15 __kernel void pipeConsumer(__global float *dst, __read_only pipe float inPipe)
16 {
17     int gid = get_global_id(0);
18     float dstPipe;
19     reserve_id_t resID = reserve_read_pipe(inPipe, 1);
20     if (is_valid_reserve_id(resID))
21     {
22         if (read_pipe(inPipe, resID, 0, &dstPipe) != 0)
23             return;
24         commit_read_pipe(inPipe, resID);
25     }
26     dst[gid] = dstPipe;
27 }

  1 //main.c
  2 #include <stdio.h>  
  3 #include <stdlib.h>  
  4 #include <cl.h>
  5 
  6 const char *sourceCode = "D:/Code/pipe.cl";
  7 
  8 char* readSource(const char* kernelPath)// 讀取文本文件，存儲爲 char *
  9 {
 10     FILE *fp;
 11     char *source;
 12     long int size;
 13     //printf("readSource, Program file: %s\n", kernelPath);
 14     fopen_s(&fp, kernelPath, "rb");
 15     if (!fp)
 16     {
 17         printf("Open kernel file failed\n");
 18         exit(-1);
 19     }
 20     if (fseek(fp, 0, SEEK_END) != 0)
 21     {
 22         printf("Seek end of file faildd\n");
 23         exit(-1);
 24     }
 25     if ((size = ftell(fp)) < 0)
 26     {
 27         printf("Get file position failed\n");
 28         exit(-1);
 29     }
 30     rewind(fp);
 31     if ((source = (char *)malloc(size + 1)) == NULL)
 32     {
 33         printf("Allocate space failed\n");
 34         exit(-1);
 35     }
 36     fread(source, 1, size, fp);
 37     fclose(fp);
 38     source[size] = '\0';
 39     return source;
 40 }
 41 
 42 int main()
 43 {
 44     const int nPacket = 1024, dataSize = nPacket * sizeof(float);
 45     char info[1024] = { 0 };
 46     int i;
 47 
 48     // 初始化平臺
 49     cl_int status;    
 50     cl_platform_id platform;
 51     clGetPlatformIDs(1, &platform, NULL);    
 52     cl_device_id device;
 53     clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 1, &device, NULL);
 54     cl_context_properties contextProp[] = { CL_CONTEXT_PLATFORM,(cl_context_properties)(platform), 0 };
 55     cl_context context = clCreateContext(contextProp, 1, &device, NULL, contextProp, &status);
 56     cl_command_queue queue = clCreateCommandQueueWithProperties(context, device, NULL, &status);    
 57     cl_event eventProducer, eventConsumer; 
 58 
 59     const char* source = readSource(sourceCode);
 60     cl_program program = clCreateProgramWithSource(context, 1, &source, NULL, &status);    
 61     status = clBuildProgram(program, 1, &device, "-w -g –cl-std=CL2.0", NULL, NULL);
 62 
 63     clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, 1024, info, NULL);
 64     printf("Build log:\n%s\n", info);
 65 
 66     cl_kernel kernelProducer = clCreateKernel(program, "pipeProducer", &status);
 67     cl_kernel kernelConsumer = clCreateKernel(program, "pipeConsumer", &status);
 68     size_t globalSize = nPacket, localSize = 128;
 69 
 70     float *srcHost = (float *)malloc(dataSize);
 71     float *dstHost = (float *)malloc(dataSize);
 72     for (i = 0; i < nPacket; srcHost[i] = i, dstHost[i] = 0.0f, i++);
 73         
 74     cl_mem src, dst;
 75     src = clCreateBuffer(context, CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR, dataSize, srcHost, &status);
 76     dst = clCreateBuffer(context, CL_MEM_WRITE_ONLY, dataSize, NULL, &status);
 77    
 78     cl_mem pipe = clCreatePipe(context, CL_MEM_HOST_NO_ACCESS, sizeof(float), nPacket, NULL, &status);
 79 
 80     clSetKernelArg(kernelProducer, 0, sizeof(cl_mem),src);
 81     clSetKernelArg(kernelProducer, 1, sizeof(cl_mem), &pipe);
 82 
 83     clSetKernelArg(kernelProducer, 0, sizeof(cl_mem), dst);
 84     clSetKernelArg(kernelProducer, 1, sizeof(cl_mem), &pipe);
 85     
 86     clEnqueueNDRangeKernel(queue, kernelProducer, 1, NULL, &globalSize, &localSize, 0, NULL, &eventProducer);
 87     clEnqueueNDRangeKernel(queue, kernelConsumer, 1, NULL, &globalSize, &localSize, 1, &eventProducer, &eventConsumer);        
 88     clEnqueueReadBuffer(queue, dst, CL_TRUE, dataSize, dataSize, dstHost, 1, &eventConsumer, NULL);
 89     clFinish(queue);
 90 
 91     for (i = 0; i < nPacket; i++)
 92     {
 93         if (dstHost[i] != i)
 94             break;
 95     }
 96     printf("Output is %s.\n", (i == nPacket) ? "correct" : "incorrect");
 97 
 98     free(srcHost);
 99     free(dstHost);
100     clReleaseContext(context);    
101     clReleaseCommandQueue(queue);
102     clReleaseProgram(program);
103     clReleaseKernel(kernelProducer);
104     clReleaseKernel(kernelConsumer);
105     getchar();
106     return 0;
107 }

● 輸出結果

■ 使用編譯參數 "-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）使用局部內存來統一工做組的寫入

 1 //pipe2.cl
 2 __kernel void pipeProducer(__global float *src, __write_only pipe float outPipe)
 3 {
 4     int gid = get_global_id(0), lid = get_local_id(0);
 5     __local reserve_id_t resID;
 6     if (lid == 0)
 7         resID = reserve_write_pipe(outPipe, get_local_size(0)); // 工做組中首個工做項一次預約多個管道位置
 8     barrier(CLK_LOCAL_MEM_FENCE);
 9 
10     float srcPipe = src[id];
11     if (is_valid_reserve_id(resID))
12     {
13         if (write_pipe(outPipe, resID, lid, &srcPipe) != 0)     // 每一個工做項寫入預約的位置
14             return;
15         commit_write_pipe(outPipe, resID);
16     }
17 }
18 
19 __kernel void pipeConsumer(__global float *dst, __read_only pipe float inPipe)
20 {
21     int gid = get_global_id(0), lid = get_local_id(0);    
22     __local reserve_id_t resID;
23     if (lid == 0)
24         resID = reserve_read_pipe(inPipe, get_local_size(0));
25     barrier(CLK_LOCAL_MEM_FENCE);
26     
27     float dstPipe;
28     if (is_valid_reserve_id(resID))
29     {
30         if (read_pipe(inPipe, resID, lid, &dstPipe) != 0)
31             return;
32         commit_read_pipe(inPipe, resID);
33     }
34     dst[gid] = dstPipe;
35 }

● （2）使用工做組管道操做簡化上述代碼（只是幹掉了一個 if 和一個同步）

 1 //pipe3.cl
 2 __kernel void pipeProducer(__global float *src, __write_only pipe float outPipe)
 3 {
 4     int gid = get_global_id(0), lid = get_local_id(0);
 5     __local reserve_id_t resID = work_group_reserve_write_pipe(outPipe, get_local_size(0));// 自帶分支和同步
 6 
 7     float srcPipe = src[id];
 8     if (is_valid_reserve_id(resID))
 9     {
10         if (write_pipe(outPipe, resID, lid, &srcPipe) != 0)
11             return;
12         commit_write_pipe(outPipe, resID);
13     }
14 }
15 
16 __kernel void pipeConsumer(__global float *dst, __read_only pipe float inPipe)
17 {
18     int gid = get_global_id(0), lid = get_local_id(0);    
19     __local reserve_id_t resID = work_group_reserve_read_pipe(inPipe, get_local_size(0));    
20     
21     float dstPipe;
22     if (is_valid_reserve_id(resID))
23     {
24         if (read_pipe(inPipe, resID, lid, &dstPipe) != 0)
25             return;
26         commit_read_pipe(inPipe, resID);
27     }
28     dst[gid] = dstPipe;
29 }

 ● 書上本來的主函數的內容（關於數據緩衝區的部分），是用虛擬內存寫的，因爲辦公室的電腦不支持，上面的代碼中被我換成了普通緩衝區

 1     float *src = (float *)clSVMAlloc(context, CL_MEM_READ_WRITE | CL_MEM_SVM_FINE_GRAIN_BUFFER, dataSize, 0);
 2     float *dst = (float *)clSVMAlloc(context, CL_MEM_READ_WRITE | CL_MEM_SVM_FINE_GRAIN_BUFFER, dataSize, 0);
 3     if (src == NULL || dst == NULL)
 4     {
 5         printf("clSVMAlloc failed!\n");
 6         getchar();
 7         return 0;
 8     }
 9 
10     clEnqueueSVMMap(queue, CL_TRUE, CL_MAP_WRITE, src, dataSize, 0, NULL, NULL);
11     for (i = 0; i < nPacket; i++)
12         src[i] = i, dst[i] = 0.0f;
13     clEnqueueSVMUnmap(queue, src, 0, NULL, NULL);
14 
15     cl_mem pipe = clCreatePipe(context, CL_MEM_HOST_NO_ACCESS, sizeof(float), nPacket, NULL, &status);
16 
17     clSetKernelArgSVMPointer(kernelProducer, 0, src);
18     clSetKernelArg(kernelProducer, 1, sizeof(cl_mem), &pipe);
19 
20     clSetKernelArgSVMPointer(kernelProducer, 0, dst);
21     clSetKernelArg(kernelProducer, 1, sizeof(cl_mem), &pipe);
22     
23     clEnqueueNDRangeKernel(queue, kernelProducer, 1, NULL, &globalSize, &localSize, 0, NULL, &eventProducer);
24     clEnqueueNDRangeKernel(queue, kernelConsumer, 1, NULL, &globalSize, &localSize, 1, &eventProducer,NULL);
25     clFinish(queue);
26     
27     clEnqueueSVMMap(queue, CL_TRUE, CL_MAP_READ, dst, dataSize, 0, NULL, NULL);    
28     for (i = 0; i < nPacket; i++)
29     {
30         if (dst[i] != i)
31             break;
32     }
33     printf("Output is %s.\n", (i == nPacket) ? "correct" : "incorrect");
34     clEnqueueSVMUnmap(queue, dst, 0, NULL, NULL);