[Beignet] [PATCH 3/5] Remove intel_gpgpu_check_binded_buf_address()

[Zhigang Gong]

This assertion is just to make sure we will not get a NULL pointer
for a normal buffer. The OpenCL spec doesn't give a very specific
statement about a NULL buffer object. But it does allow to pass
a NULL to a buffer object. Thus some one may implement the following
kernel:

__kernel foo( global uint * input, global uint *output1, global uint *output2)
{
  ...
  if (output1)
    output1[get_global_id(0)] = result0;
  if (output2)
    output2[get_global_id(1)] = result1;
}

If we pass in a NULL output1 which should be a normal allocated buffer,
it breaks the above code.

Without PPGTT, this assertion works fine till now. If with PPGTT,
we may hit this assertion, but we can't just simply remove this assertion.

Do you have good suggestion to always avoid allocate 0 offset for a
valid buffer object?

All the other patches in this series LGTM, I will push them latter.
Let's defer this one to next week.

Thanks,
Zhigang Gong.

On Thu, Oct 23, 2014 at 03:19:24PM +0800, Zhenyu Wang wrote:
> On recent kernel with full PPGTT support, we can possibly bind buffer
> offset with 0, but intel_gpgpu_check_binded_buf_address() always thinks
> it's invalid, which is not true. So simply remove the check.
> 
> Signed-off-by: Zhenyu Wang <zhenyuw at linux.intel.com>
> ---
>  src/intel/intel_gpgpu.c | 9 ---------
>  1 file changed, 9 deletions(-)
> 
> diff --git a/src/intel/intel_gpgpu.c b/src/intel/intel_gpgpu.c
> index 6cd73d6..b7958d5 100644
> --- a/src/intel/intel_gpgpu.c
> +++ b/src/intel/intel_gpgpu.c
> @@ -694,14 +694,6 @@ intel_gpgpu_batch_reset(intel_gpgpu_t *gpgpu, size_t sz)
>  {
>    return intel_batchbuffer_reset(gpgpu->batch, sz);
>  }
> -/* check we do not get a 0 starting address for binded buf */
> -static void
> -intel_gpgpu_check_binded_buf_address(intel_gpgpu_t *gpgpu)
> -{
> -  uint32_t i;
> -  for (i = 0; i < gpgpu->binded_n; ++i)
> -    assert(gpgpu->binded_buf[i]->offset != 0);
> -}
>  
>  static void
>  intel_gpgpu_flush_batch_buffer(intel_batchbuffer_t *batch)
> @@ -717,7 +709,6 @@ intel_gpgpu_flush(intel_gpgpu_t *gpgpu)
>    if (!gpgpu->batch || !gpgpu->batch->buffer)
>      return;
>    intel_gpgpu_flush_batch_buffer(gpgpu->batch);
> -  intel_gpgpu_check_binded_buf_address(gpgpu);
>  }
>  
>  static int
> -- 
> 2.1.1
> 
> _______________________________________________
> Beignet mailing list
> Beignet at lists.freedesktop.org
> http://lists.freedesktop.org/mailman/listinfo/beignet