[Beignet] Pointer load/store - obsolete documentation?

Rebecca N. Palmer rebecca_palmer at zoho.com
Thu Jun 9 21:40:09 UTC 2016


The documentation says "an unsupported feature which is to store/load 
pointers to/from memory[...]We plan to fix it in next major release 
1.1.0" (docs/Beignet.mdwn:216-219)

Given that 1.1.0 has now been released, and does appear to fix this 
(creating an array of pointers works in 1.1.2 but not in 1.0.3; 
https://cgit.freedesktop.org/beignet/commit/?id=48e22ae4536fbf944f39ab13e8f1133d3f5edcc3), 
should this note now be removed?

(This is the test case I used: would you like it making into a utest?)
/*output in 1.1.2:
built 0 built 0 kernel created 00 returned (7, 8, 7)
1 returned (2, 7, 6)
2 returned (0, 1, 0)

output in 1.0.3:
store i32 7, i32 addrspace(1)* %26, align 4, !tbaa !11
built 0 Illegal pointer which is not from a valid memory space.
Aborting...
*/

#include <stdio.h>
#include <CL/cl.h>
int main(int argc,char** argv)
{
   cl_platform_id platforms[4];
   cl_uint num_platforms;
   cl_device_id device;
   cl_int status, ret;
   cl_context ctx;
   cl_command_queue queue;
   cl_program program,program2;
   cl_kernel kernel;
   cl_mem buffer[3];
   cl_event kernel_finished;
   FILE *f;
   size_t n = 3,program_length;
   int i;
   cl_int test_data[3][3] = {{3, 8, 5},{2, 4, 6},{0, 1, 0}};
   const char *kernel_source="__kernel void pointer_store_load"
   "(__global int *p0, __global int *p1, __global int *p2)"
   "{__global int * __private a[2];int i;"
   "i=get_global_id(0);a[0]=p0;a[1]=p1;a[p2[i]][i]=7;}";
   char *build_log;
   ret = 2;
   build_log=calloc(1001,1);
   clGetPlatformIDs(4,platforms,&num_platforms);
   for(i=0;i<num_platforms;i++){
   status=clGetDeviceIDs(platforms[i],CL_DEVICE_TYPE_GPU,
   1,&device,NULL);
   if (status == CL_SUCCESS) {
   break;
   }
   }
   if (status != CL_SUCCESS) {
   printf("Device not found");
   exit(1);
   }
   ctx = clCreateContext(NULL, 1, &device, NULL, NULL,
    &status);
   if (status == CL_SUCCESS) {
     queue = clCreateCommandQueue(ctx, device, 0,
      &status);
     if (status == CL_SUCCESS) {
       program = clCreateProgramWithSource(ctx, 1,
        &kernel_source, NULL, &status);
       if (status == CL_SUCCESS) {
         status = clCompileProgram(program, 1, &device,
          "", 0, NULL, NULL, NULL, NULL);
         clGetProgramBuildInfo(program,device,
         CL_PROGRAM_BUILD_LOG,1000,build_log,NULL);
         printf("built %i %s",status,build_log);
         program2 = clLinkProgram(ctx, 1, &device,
          "", 1,&program, NULL, NULL, NULL);
         clGetProgramBuildInfo(program,device,
         CL_PROGRAM_BUILD_LOG,1000,build_log,NULL);
         printf("built %i %s",status,build_log);
         if (status == CL_SUCCESS) {
           kernel = clCreateKernel(program2, "pointer_store_load",
            &status);
           printf("kernel created %i",status);
           if (status == CL_SUCCESS) {
             for(i=0;i<3;i++){
             buffer[i] = clCreateBuffer(ctx, CL_MEM_COPY_HOST_PTR,
              4*n, test_data[i], &status);
             }
             if (status == CL_SUCCESS) {
             for(i=0;i<3;i++){
               status = clSetKernelArg(kernel, i,
                sizeof(cl_mem), &buffer[i]);
               }
               if (status == CL_SUCCESS) {
                 status = clEnqueueNDRangeKernel(queue, kernel, 1,
                  NULL, &n, &n, 0, NULL, &kernel_finished);
                 if (status == CL_SUCCESS) {
                 for(i=0;i<3;i++){
                   status = clEnqueueReadBuffer(queue, buffer[i],
                   CL_TRUE, 0, n*4, test_data[i], 1,
                    &kernel_finished, NULL);
                   }
                   if (status == CL_SUCCESS) {
                   for(i=0;i<3;i++){
                     printf("%i returned (%i, %i, %i)\n", i,
                     test_data[i][0], test_data[i][1],
                      test_data[i][2]);
                     }
                   }
                 }
               }
             }
             for(i=0;i<3;i++){
             clReleaseMemObject(buffer[i]);
             }
           }
           clReleaseKernel(kernel);
         }
       }
       clReleaseProgram(program);
     }
     clReleaseCommandQueue(queue);
   }
   clReleaseContext(ctx);
}




More information about the Beignet mailing list