[Beignet] CL_MEM_USE_HOST_PTR usage in clCreateBuffer API

[Zhigang Gong]

On Tue, Jul 23, 2013 at 01:54:25AM -0700, Edward Ching wrote:
> I tried the CL_MEM_USE_HOST_PTR option from one of the recent Beignet
> checkins using a simple test case but it failed. Did I misunderstand the
> usage of CL_MEM_USE_HOST_TR?  Same test cased worked OK when using Intel
> OpenCL SDK 2012 (Windows) for IVB. The test case patch is attached and also
> cut n pasted here.

It seems that your usage of CL_MEM_USE_HOST_PTR is not fully comply with the spec.

Please refer the following statement I just copied from the spec:

  If the buffer object is created with CL_MEM_USE_HOST_PTR set in mem_flags, the following 
  will be true: 
    The host_ptr specified in clCreateBuffer is guaranteed to contain the latest bits in the 
    region being mapped when the clEnqueueMapBuffer command has completed. 
    The pointer value returned by clEnqueueMapBuffer will be derived from the host_ptr 
    specified when the buffer object is created. 

So before you access the out buffer, you need to call clEnqueueMapBuffer.
Actually, this is a little unclear in the spec. You can refer some discussion at
the following link for more details:
http://www.khronos.org/message_boards/showthread.php/6912-Clarify-CL_MEM_USE_HOST_PTR

> 
> /Ed
> 
> 
> 1) test_zcopy.cl:
> 
> __kernel void
> test_zcopy(__global int *src, __global int *dst, int n, int num_threads)
> {
> int tid, pos;
>  tid = get_global_id(0);
> for (pos=tid; pos < n; pos+=num_threads) {
> dst[pos] = src[pos];
> }
> }
> 
> 2) compiler_zcopy.cpp:
> 
> #include "utest_helper.hpp"
> #include <sys/time.h>
> #include <malloc.h>
> #include <string.h>
> 
> #define T_GET(t)        gettimeofday(&t, NULL);
> #define T_LAPSE(t1, t2) \
>   ((t2.tv_sec+t2.tv_usec*0.000001) - (t1.tv_sec+t1.tv_usec*0.000001))
> #define ALIGN 4096
> 
> static void compiler_zcopy(void)
> {
>   const size_t n = 32*1024*1024;
>   struct timeval t1, t2;
>   float t_fin;
>   int *in, *out;
> 
>   in  = (int *)memalign(ALIGN, n*sizeof(int));
>   out = (int *)memalign(ALIGN, n*sizeof(int));
>   OCL_ASSERT(in);
>   OCL_ASSERT(out);
>   for (int i=0; i<n; ++i) in[i] = i;
>   bzero(out, n);
> 
>   // Setup kernel and buffers
>   OCL_CREATE_KERNEL("test_zcopy");
>   OCL_CREATE_BUFFER(buf[0], CL_MEM_USE_HOST_PTR, n * sizeof(int), in);
>   OCL_CREATE_BUFFER(buf[1], CL_MEM_USE_HOST_PTR, n * sizeof(int), out);
> 
>   // Run the kernel
>   locals[0]  = 64;
>   globals[0] = 32 * locals[0];
>   OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
>   OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]);
>   OCL_SET_ARG(2, sizeof(int), &n);
>   OCL_SET_ARG(3, sizeof(int), &globals[0]);
> 
>   T_GET(t1);
>   OCL_NDRANGE(1);
>   OCL_FINISH();
>   T_GET(t2);
>   t_fin = T_LAPSE(t1, t2);
>   printf("test_zcopy: lapse = %f\n",t_fin);
>   for (int i=0; i<n; ++i) OCL_ASSERT((out[i]==i));
> }
> 
> MAKE_UTEST_FROM_FUNCTION(compiler_zcopy);


> _______________________________________________
> Beignet mailing list
> Beignet at lists.freedesktop.org
> http://lists.freedesktop.org/mailman/listinfo/beignet