[Beignet] [PATCH 10/18] GBE: Disable SPF and use JMPI + IF/ENDIF to handle each blocks.

[Zhigang Gong]

I found some garbage hided in this patch, I marked them as below, and please
ignore them and continue to review the rest part.
Thanks.
 
On Fri, Mar 28, 2014 at 03:10:48PM +0800, Zhigang Gong wrote:
> When enable SPF (single program flow), we always need to use f0
> as the predication of almost each instruction. This bring some
> trouble when we want to get tow levels mask mechanism, for an
> example the SEL instruction, and some BOOL operations. We
> have to use more than one instructions to do that and simply
> introduce 100% of overhead of those instructions.
> 
> Signed-off-by: Zhigang Gong <zhigang.gong at intel.com>
> ---
> diff --git a/src/cl_api.c b/src/cl_api.c
> index 9638994..b572500 100644
> --- a/src/cl_api.c
> +++ b/src/cl_api.c
> @@ -2526,6 +2526,7 @@ clEnqueueNDRangeKernel(cl_command_queue  command_queue,
>  
>      err = cl_command_queue_flush(command_queue);
>    }
> +  clFinish(command_queue);

This above line should be removed, it is for debug purpose only. 

The following changes in the utests are also for debug purpose and should be removed.

> diff --git a/utests/compiler_long_cmp.cpp b/utests/compiler_long_cmp.cpp
> index 35d4c4f..f901fdf 100644
> --- a/utests/compiler_long_cmp.cpp
> +++ b/utests/compiler_long_cmp.cpp
> @@ -45,6 +45,7 @@ void compiler_long_cmp(void)
>      int64_t *dest = (int64_t *)buf_data[2];
>      int64_t x = (src1[i] < src2[i]) ? 3 : 4;
>      OCL_ASSERT(x == dest[i]);
> +    //printf("%d %ld  %ld \n", i, dest[i], x);
>    }
>    OCL_UNMAP_BUFFER(2);
>    OCL_DESTROY_KERNEL_KEEP_PROGRAM(true);
> diff --git a/utests/compiler_unstructured_branch0.cpp b/utests/compiler_unstructured_branch0.cpp
> index 128a53e..1a371e9 100644
> --- a/utests/compiler_unstructured_branch0.cpp
> +++ b/utests/compiler_unstructured_branch0.cpp
> @@ -27,7 +27,6 @@ static void compiler_unstructured_branch0(void)
>      OCL_ASSERT(((int32_t*)buf_data[1])[i] == 2);
>    for (uint32_t i = 16; i < 32; ++i)
>      OCL_ASSERT(((int32_t*)buf_data[1])[i] == 1);
> -
>    // Second control flow
>    for (uint32_t i = 0; i < n; ++i) ((int32_t*)buf_data[0])[i] = -2;
>    OCL_UNMAP_BUFFER(0);
> @@ -36,8 +35,7 @@ static void compiler_unstructured_branch0(void)
>    OCL_MAP_BUFFER(0);
>    OCL_MAP_BUFFER(1);
>    for (uint32_t i = 0; i < 32; ++i)
> -    OCL_ASSERT(((int32_t*)buf_data[1])[i] == 1);
> -
> +   OCL_ASSERT(((int32_t*)buf_data[1])[i] == 1);
>    // Third control flow
>    for (uint32_t i = 0; i < 8; ++i) ((int32_t*)buf_data[0])[i] = 2;
>    OCL_UNMAP_BUFFER(0);
> diff --git a/utests/compiler_unstructured_branch1.cpp b/utests/compiler_unstructured_branch1.cpp
> index 6021f5b..fb24cec 100644
> --- a/utests/compiler_unstructured_branch1.cpp
> +++ b/utests/compiler_unstructured_branch1.cpp
> @@ -25,7 +25,6 @@ static void compiler_unstructured_branch1(void)
>    OCL_MAP_BUFFER(1);
>    for (uint32_t i = 0; i < n; ++i)
>      OCL_ASSERT(((int32_t*)buf_data[1])[i] == 2);
> -
>    // Second control flow
>    for (uint32_t i = 0; i < n; ++i) ((int32_t*)buf_data[0])[i] = -2;
>    OCL_UNMAP_BUFFER(0);
> @@ -34,7 +33,7 @@ static void compiler_unstructured_branch1(void)
>    OCL_MAP_BUFFER(0);
>    OCL_MAP_BUFFER(1);
>    for (uint32_t i = 0; i < n; ++i)
> -    OCL_ASSERT(((uint32_t*)buf_data[1])[i] == 3);
> +   OCL_ASSERT(((uint32_t*)buf_data[1])[i] == 3);
>  
>    // Third control flow
>    for (uint32_t i = 0; i < 8; ++i) ((int32_t*)buf_data[0])[i] = 2;
> diff --git a/utests/compiler_unstructured_branch2.cpp b/utests/compiler_unstructured_branch2.cpp
> index d61c6b5..68c7448 100644
> --- a/utests/compiler_unstructured_branch2.cpp
> +++ b/utests/compiler_unstructured_branch2.cpp
> @@ -23,6 +23,7 @@ static void compiler_unstructured_branch2(void)
>    // First control flow
>    OCL_MAP_BUFFER(0);
>    OCL_MAP_BUFFER(1);
> +#if 1
>    for (uint32_t i = 0; i < n; ++i)
>      OCL_ASSERT(((int32_t*)buf_data[1])[i] == 12);
>  
> @@ -35,7 +36,7 @@ static void compiler_unstructured_branch2(void)
>    OCL_MAP_BUFFER(1);
>    for (uint32_t i = 0; i < n; ++i)
>      OCL_ASSERT(((int32_t*)buf_data[1])[i] == -6);
> -
> +#endif
>    // Third control flow
>    for (uint32_t i = 0; i < 8; ++i) ((int32_t*)buf_data[0])[i] = 2;
>    for (uint32_t i = 8; i < n; ++i) ((int32_t*)buf_data[0])[i] = -2;
> @@ -45,9 +46,16 @@ static void compiler_unstructured_branch2(void)
>    OCL_MAP_BUFFER(0);
>    OCL_MAP_BUFFER(1);
>    for (uint32_t i = 0; i < 8; ++i)
> +  {
> +    //printf("%d: %d %d\n", i, ((int32_t*)buf_data[1])[i], 12);
>      OCL_ASSERT(((int32_t*)buf_data[1])[i] == 12);
> +  }
>    for (uint32_t i = 8; i < n; ++i)
> +  {
> +    //printf("%d: %d %d\n", i, ((int32_t*)buf_data[1])[i], -6);
>      OCL_ASSERT(((int32_t*)buf_data[1])[i] == -6);
> +  }
> +  //exit(0);
>  
>    // Fourth control flow
>    for (uint32_t i = 0; i < 4; ++i) ((int32_t*)buf_data[0])[i] = 1;
> -- 
> 1.8.3.2
> 
> _______________________________________________
> Beignet mailing list
> Beignet at lists.freedesktop.org
> http://lists.freedesktop.org/mailman/listinfo/beignet