A file holds a data structure that is written and modified by number of users in a distributed manner. Multiple users on multiple computers use Read-Modify-Write cycle provided resource is available else use use modify once before exit. Write necessary Program using OpenCL.

PROGRAM

b13.c

#include <stdio.h>

#include <stdlib.h>

#include <CL/cl.h>

#define SRC_SIZE (0x100000)

     

int main()

{

    cl_device_id device_id = NULL;

    cl_context context = NULL;

    cl_command_queue command_queue = NULL;

    cl_program program = NULL;

    cl_kernel kernel = NULL;

    cl_platform_id platform_id = NULL;

    cl_uint num_devices, num_platforms;

    cl_mem memobj1 = NULL, memobj2 = NULL, memobj3 = NULL;

    size_t g=16, l=1;

     

    FILE *fp = fopen("./b13.cl", "r");

    if(!fp) 

    {

        printf("Failed to load kernel.\n");

        exit(1);

    }

    char * src = (char*)malloc(SRC_SIZE);

    size_t src_size = fread(src, 1, SRC_SIZE, fp);

    fclose(fp);

    

    clGetPlatformIDs(1, &platform_id, &num_platforms);

    clGetDeviceIDs(platform_id, CL_DEVICE_TYPE_DEFAULT, 1, &device_id, &num_devices);

    context = clCreateContext(NULL, 1, &device_id, NULL, NULL, NULL); 

    command_queue = clCreateCommandQueue(context, device_id, 0, NULL); 

    memobj1 = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(int), NULL, NULL);

    memobj2 = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(int), NULL, NULL);

    memobj3 = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(int), NULL, NULL);

    

    program = clCreateProgramWithSource(context, 1, (const char **)&src, (const size_t *)&src_size, NULL); 

    clBuildProgram(program, 1, &device_id, NULL, NULL, NULL);

    kernel = clCreateKernel(program, "b13", NULL);   

    clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&memobj1);

    clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *)&memobj2);

    clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *)&memobj3); 

    clEnqueueNDRangeKernel(command_queue, kernel, 1, NULL, &g, &l, 0, NULL, NULL);

      

    clFlush(command_queue);

    clFinish(command_queue);

    clReleaseKernel(kernel);

    clReleaseProgram(program);

    clReleaseMemObject(memobj1);

    clReleaseMemObject(memobj2);

    clReleaseMemObject(memobj3);

    clReleaseCommandQueue(command_queue);

    clReleaseContext(context);

    free(src);

    return 0;

}

b13.cl

#pragma OPENCL EXTENSION cl_khr_global_int32_base_atomics : enable

#pragma OPENCL EXTENSION cl_khr_local_int32_base_atomics : enable

#pragma OPENCL EXTENSION cl_khr_global_int32_extended_atomics : enable

#pragma OPENCL EXTENSION cl_khr_local_int32_extended_atomics : enable

void GetSem(__global int * sem) 

{

    int occupied = atom_xchg(sem, 1);

    while(occupied > 0)

        occupied = atom_xchg(sem, 1);

}

void ReleaseSem(__global int * sem)

{

    int prevVal = atom_xchg(sem, 0);

}

__kernel void b13(__global int * sem, __global int * x, __global int * lock)

{

    int i = get_global_id(0);

    if(i%2 == 0)

    {

        GetSem(&sem[0]);

        *x = i;

        *lock = 1;

        printf("Kernel %d setting value of x: %d\n", i, *x);

        ReleaseSem(&sem[0]);

    }

    else

    {

        while((*lock)!=1)

            printf("Kernel %d waiting for first write\n",i);

        GetSem(&sem[0]);

        printf("Kernel %d reading value of x: %d\n", i, *x); 

        ReleaseSem(&sem[0]);  

    }

}