OpenCL 存储器次序的验证

▶ 《OpenCL异构并行编程实战》P224 的代码，先放上来，坐等新设备到了再执行

 //kernel.cl

 __global volatile atomic_int globalAtom = ATOMIC_VAR_INIT();   // 全局原子对象

 __kernel void memoryOrderTest01(__global int *dst)

 {

     __local volatile atomic_int localAtom;                      // 本地原子对象

     atomic_init(&localAtom, );

     const int gid = get_global_id();

     work_group_barrier(CLK_LOCAL_MEM_FENCE);

     if (gid == )                                               // 0 号工作项尝试写入 1

     {

         atomic_store_explicit(&localAtom, , memory_order_seq_cst, memory_scope_work_group);

         atomic_store_explicit(&globalAtom, , memory_order_seq_cst, memory_scope_device);

     }

     //atomic_work_item_fence(CLK_LOCAL_MEM_FENCE, memory_order_acq_rel, memory_scope_work_group);

     if (gid == )

     {

         int a, count;

         for (a = , count = ; a ==  && count < ; count++)

             a = atomic_load_explicit(&localAtom, memory_order_seq_cst, memory_scope_work_group);

         dst[] = !!a;

         dst[] = count;

         for (count = ; a ==  && count < ; count++)

             a = atomic_load_explicit(&globalAtom, memory_order_seq_cst, memory_scope_device);

         dst[] = !!a;

     }

     work_group_barrier();// 必须添加，将 0 号工作项的副作用暴露给其他工作项

 }

 __kernel void memoryOrderTest02(__global int *dst)

 {

     __local volatile atomic_int localAtom;

     atomic_init(&localAtom, );

     const int gid = get_global_id();

     work_group_barrier(CLK_LOCAL_MEM_FENCE);

     if (gid == )

     {

         atomic_store(&localAtom, );

         atomic_store(&globalAtom, );

     }

     //atomic_work_item_fence(CLK_LOCAL_MEM_FENCE, memory_order_acq_rel, memory_scope_work_group);

     if (gid == )

     {

         int a, count;

         for (a = , count = ; a ==  && count < ; count++)

             a = atomic_load(&localAtom);

         dst[] = !!a;

         dst[] = count;

         for (count = ; a ==  && count < ; count++)

             a = atomic_load(&globalAtom);

         dst[] = !!a;

     }

     work_group_barrier();

 }

 //main.c

 #include <stdio.h>

 #include <stdlib.h>

 #include <cl.h>

 const char *sourceCode = "D:/Code/kernel.cl";

 int readSource(const char* kernelPath, char *source)// 读取文本文件，存储为 char *，返回代码长度

 {

     FILE *fp;

     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 size + ;

 }

 int main()

 {

     const int nElement = , dataSize = nElement * sizeof(float);

     int i, host[nElement] = {  };

     char info[];

     // 初始化平台

     cl_int status;

     cl_platform_id platform;

     clGetPlatformIDs(, &platform, NULL);

     cl_device_id device[];

     clGetDeviceIDs(platform, CL_DEVICE_TYPE_ALL, , 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_properties queueProp = ;// useless

     cl_command_queue queue = clCreateCommandQueueWithProperties(context, device[], NULL, &status);

     cl_mem buffer = clCreateBuffer(context, CL_MEM_WRITE_ONLY, dataSize, NULL, &status);

     char *source;

     size_t sourceLength = readSource(sourceCode, source);

     cl_program program = clCreateProgramWithSource(context, , &source, &sourceLength, &status);

     status = clBuildProgram(program, , device, "-cl-std=CL2.0", NULL, NULL);

     if (status)

     {

         clGetProgramBuildInfo(program, device[], CL_PROGRAM_BUILD_LOG, , info, NULL);

         printf("Build log:\n%s\n", info);

     }

     cl_kernel kernel = clCreateKernel(program, "memoryOrderTest", &status);

     clSetKernelArg(kernel, , sizeof(cl_mem), buffer);

     size_t globalSize = nElement, localSize = ;

     clEnqueueNDRangeKernel(queue, kernel, , NULL, &globalSize, &localSize, , NULL, NULL);

     clFinish(queue);

     clEnqueueReadBuffer(queue, buffer, CL_TRUE, , dataSize, host, , NULL, NULL);

     printf("Local memory result: %d, global memory result: %d, waiting count: %d\n", host[], host[], host[]);

     clReleaseContext(context);

     clReleaseCommandQueue(queue);

     clReleaseProgram(program);

     clReleaseKernel(kernel);

     clReleaseMemObject(buffer);

     getchar();

     return ;

 }