[Beignet] [PATCH 10/18] GBE: Disable SPF and use JMPI + IF/ENDIF to handle each blocks.
Zhigang Gong
zhigang.gong at linux.intel.com
Tue Apr 1 00:56:00 PDT 2014
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
More information about the Beignet
mailing list