OpenCL 存储器次序的验证

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

//kernel.cl
__global volatile atomic_int globalAtom = ATOMIC_VAR_INIT(0);   // 全局原子对象
__kernel void memoryOrderTest01(__global int *dst)
{
    __local volatile atomic_int localAtom;                      // 本地原子对象   
    atomic_init(&localAtom, 0);
    const int gid = get_global_id(0);
    work_group_barrier(CLK_LOCAL_MEM_FENCE);
    if (gid == 0)                                               // 0 号工作项尝试写入 1
    {
        atomic_store_explicit(&localAtom, 1, memory_order_seq_cst, memory_scope_work_group);
        atomic_store_explicit(&globalAtom, 1, 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 == 64)
    {
        int a, count;
        for (a = 0, count = 1; a == 0 && count < 10000; count++)
            a = atomic_load_explicit(&localAtom, memory_order_seq_cst, memory_scope_work_group);
        dst[0] = !!a;
        dst[2] = count;
        for (count = 1; a == 0 && count < 10000; count++)
            a = atomic_load_explicit(&globalAtom, memory_order_seq_cst, memory_scope_device);
        dst[1] = !!a;
    }
    work_group_barrier(0);// 必须添加，将 0 号工作项的副作用暴露给其他工作项
}

__kernel void memoryOrderTest02(__global int *dst)
{
    __local volatile atomic_int localAtom;
    atomic_init(&localAtom, 0);
    const int gid = get_global_id(0);
    work_group_barrier(CLK_LOCAL_MEM_FENCE);
    if (gid == 0)
    {
        atomic_store(&localAtom, 1);
        atomic_store(&globalAtom, 1);
    }
    //atomic_work_item_fence(CLK_LOCAL_MEM_FENCE, memory_order_acq_rel, memory_scope_work_group);
    if (gid == 64)
    {
        int a, count;
        for (a = 0, count = 1; a == 0 && count < 10000; count++)
            a = atomic_load(&localAtom);
        dst[0] = !!a;
        dst[2] = count;
        for (count = 1; a == 0 && count < 10000; count++)
            a = atomic_load(&globalAtom);
        dst[1] = !!a;
    }
    work_group_barrier(0);
}

//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
", kernelPath);
    fopen_s(&fp, kernelPath, "rb");
    if (!fp)
    {
        printf("Open kernel file failed
");
        exit(-1);
    }
    if (fseek(fp, 0, SEEK_END) != 0)
    {
        printf("Seek end of file faildd
");
        exit(-1);
    }
    if ((size = ftell(fp)) < 0)
    {
        printf("Get file position failed
");
        exit(-1);
    }
    rewind(fp);
    if ((source = (char *)malloc(size + 1)) == NULL)
    {
        printf("Allocate space failed
");
        exit(-1);
    }
    fread(source, 1, size, fp);
    fclose(fp);
    source[size] = '';
    return size + 1;
}

int main()
{
    const int nElement = 2048, dataSize = nElement * sizeof(float);
    int i, host[nElement] = { 0 };
    char info[1024];

    // 初始化平台
    cl_int status;
    cl_platform_id platform;
    clGetPlatformIDs(1, &platform, NULL);
    cl_device_id device[2];
    clGetDeviceIDs(platform, CL_DEVICE_TYPE_ALL, 2, device, NULL);
    cl_context_properties contextProp[] = { CL_CONTEXT_PLATFORM,(cl_context_properties)(platform), 0 };
    cl_context context = clCreateContext(contextProp, 1, device, NULL, contextProp, &status);
    cl_command_queue_properties queueProp = 0;// useless
    cl_command_queue queue = clCreateCommandQueueWithProperties(context, device[0], 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, 1, &source, &sourceLength, &status);
    status = clBuildProgram(program, 1, device, "-cl-std=CL2.0", NULL, NULL);
    if (status)
    {
        clGetProgramBuildInfo(program, device[0], CL_PROGRAM_BUILD_LOG, 1024, info, NULL);
        printf("Build log:
%s
", info);
    }
    cl_kernel kernel = clCreateKernel(program, "memoryOrderTest", &status);            
    clSetKernelArg(kernel, 0, sizeof(cl_mem), buffer);
    size_t globalSize = nElement, localSize = 256;
    clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &globalSize, &localSize, 0, NULL, NULL);
    clFinish(queue);

    clEnqueueReadBuffer(queue, buffer, CL_TRUE, 0, dataSize, host, 0, NULL, NULL);
    
    printf("Local memory result: %d, global memory result: %d, waiting count: %d
", host[0], host[1], host[2]);
       
    clReleaseContext(context);
    clReleaseCommandQueue(queue);
    clReleaseProgram(program);    
    clReleaseKernel(kernel);
    clReleaseMemObject(buffer);
    getchar();
    return 0;
}