[Beignet] [PATCH] Remove global offset need divide by local size restriction.

[Zhigang Gong]

You need to focus on the specific uint64 instruction in
reg_insn_selection.cpp, and analyze the src/dst register carefully, 

Two possible root causes there:
1. Wrongly set a register as a source but you will touch it in that
instruction.
2. Incorrectly set the register type, for example, you want to use it as 64
bit but you only set it as DWORD.

And that's all I can tell now.


-----Original Message-----
From: beignet-bounces+zhigang.gong=linux.intel.com at lists.freedesktop.org
[mailto:beignet-bounces+zhigang.gong=linux.intel.com at lists.freedesktop.org]
On Behalf Of Xing, Homer
Sent: Thursday, September 26, 2013 1:34 PM
To: Zhigang Gong; Yang, Rong R
Cc: beignet at lists.freedesktop.org
Subject: Re: [Beignet] [PATCH] Remove global offset need divide by local
size restriction.

I can take this bug. But currently I don't have enough information.

Zhigang, do you know more details about how to find dependency bug in code
path?

-----Original Message-----
From: beignet-bounces+homer.xing=intel.com at lists.freedesktop.org
[mailto:beignet-bounces+homer.xing=intel.com at lists.freedesktop.org] On
Behalf Of Zhigang Gong
Sent: Thursday, September 26, 2013 10:43 AM
To: Yang, Rong R
Cc: beignet at lists.freedesktop.org
Subject: Re: [Beignet] [PATCH] Remove global offset need divide by local
size restriction.

Hi guys,

I just pushed this patch, and it causes a regression as below:

compiler_abs_diff_long16:
  compiler_abs_diff_long16()    [FAILED]
    Error: !memcmp(buf_data[2], cpu_diff, sizeof(T) * n)
  at file /home/gongzg/git/fdo/beignet/utests/compiler_abs_diff.cpp,
function compiler_abs_diff_with_type, line 177

I took a quick look at it and found the root cause of this regression is not
in this patch. As when I disable the OCL_PRE_ALLOC_INSN, it works fine.
For me, the most possible root cause is there is a dependency bug within the
abs_diff_long code path.

I will take vacation from tomorrow and will back to work at 8/10. If anybody
in the list can take this bug, I will appreciate. Thanks

- Zhigang

On Mon, Sep 23, 2013 at 02:04:08PM +0800, Yang Rong wrote:
> Set to global offset to 0 in walker, and add global offset when
get_global_id.
> 
> Signed-off-by: Yang Rong <rong.r.yang at intel.com>
> ---
>  backend/src/ocl_stdlib.tmpl.h | 2 +-
>  src/cl_api.c                  | 7 -------
>  src/intel/intel_gpgpu.c       | 6 +++---
>  3 files changed, 4 insertions(+), 11 deletions(-)
> 
> diff --git a/backend/src/ocl_stdlib.tmpl.h 
> b/backend/src/ocl_stdlib.tmpl.h index 9b76ba1..4921ee4 100644
> --- a/backend/src/ocl_stdlib.tmpl.h
> +++ b/backend/src/ocl_stdlib.tmpl.h
> @@ -588,7 +588,7 @@ DECL_PUBLIC_WORK_ITEM_FN(get_num_groups, 1) #undef 
> DECL_PUBLIC_WORK_ITEM_FN
>  
>  INLINE uint get_global_id(uint dim) {
> -  return get_local_id(dim) + get_local_size(dim) * get_group_id(dim);
> +  return get_local_id(dim) + get_local_size(dim) * get_group_id(dim)
> + + get_global_offset(dim);
>  }
>  
>  
> //////////////////////////////////////////////////////////////////////
> /////// diff --git a/src/cl_api.c b/src/cl_api.c index 
> b60160b..c19b80a 100644
> --- a/src/cl_api.c
> +++ b/src/cl_api.c
> @@ -2264,19 +2264,12 @@ clEnqueueNDRangeKernel(cl_command_queue
command_queue,
>      goto error;
>    }
>  
> -  /* Check offset values. We add a non standard restriction. The offsets
must
> -   * also be evenly divided by the local sizes
> -   */
>    if (global_work_offset != NULL)
>      for (i = 0; i < work_dim; ++i) {
>        if (UNLIKELY(~0LL - global_work_offset[i] > global_work_size[i])) {
>          err = CL_INVALID_GLOBAL_OFFSET;
>          goto error;
>        }
> -      if (UNLIKELY(local_work_size != NULL && global_work_offset[i] %
local_work_size[i])) {
> -        err = CL_INVALID_GLOBAL_OFFSET;
> -        goto error;
> -      }
>      }
>  
>    /* Local sizes must be non-null and divide global sizes */ diff 
> --git a/src/intel/intel_gpgpu.c b/src/intel/intel_gpgpu.c index 
> 7b82b76..44f44ef 100644
> --- a/src/intel/intel_gpgpu.c
> +++ b/src/intel/intel_gpgpu.c
> @@ -886,11 +886,11 @@ intel_gpgpu_walker(intel_gpgpu_t *gpgpu,
>      OUT_BATCH(gpgpu->batch, (1 << 30) | (thread_n-1)); /* SIMD16 | thread
max */
>    else
>      OUT_BATCH(gpgpu->batch, (0 << 30) | (thread_n-1)); /* SIMD8  | 
> thread max */
> -  OUT_BATCH(gpgpu->batch, global_wk_off[0]);
> +  OUT_BATCH(gpgpu->batch, 0);
>    OUT_BATCH(gpgpu->batch, global_wk_dim[0]);
> -  OUT_BATCH(gpgpu->batch, global_wk_off[1]);
> +  OUT_BATCH(gpgpu->batch, 0);
>    OUT_BATCH(gpgpu->batch, global_wk_dim[1]);
> -  OUT_BATCH(gpgpu->batch, global_wk_off[2]);
> +  OUT_BATCH(gpgpu->batch, 0);
>    OUT_BATCH(gpgpu->batch, global_wk_dim[2]);
>    OUT_BATCH(gpgpu->batch, right_mask);
>    OUT_BATCH(gpgpu->batch, ~0x0);                     /* we always set
height as 1, so set bottom mask as all 1*/
> --
> 1.8.1.2
> 
> _______________________________________________
> Beignet mailing list
> Beignet at lists.freedesktop.org
> http://lists.freedesktop.org/mailman/listinfo/beignet
_______________________________________________
Beignet mailing list
Beignet at lists.freedesktop.org
http://lists.freedesktop.org/mailman/listinfo/beignet
_______________________________________________
Beignet mailing list
Beignet at lists.freedesktop.org
http://lists.freedesktop.org/mailman/listinfo/beignet