diff options
| author | Gian Marco Iodice <gianmarco.iodice@arm.com> | 2017-10-06 15:44:27 +0100 |
|---|---|---|
| committer | Anthony Barbier <anthony.barbier@arm.com> | 2018-11-02 16:35:24 +0000 |
| commit | 744b5edd1e7eedab8ac52a8cea33bf62fb95affc (patch) | |
| tree | 6fa2b15cfc9d0c78ab835026dc5f7c9aeebaf82d | |
| parent | ff421f2100e0e9e532f5fe78585300546af61690 (diff) | |
| download | ComputeLibrary-744b5edd1e7eedab8ac52a8cea33bf62fb95affc.tar.gz | |
COMPMID-606 - Fix for S8 failures
Added volatile to the for loop counter in direct_convolution1x1.cl, direct_convolution3x3.cl and direct_convolution5x5.cl
This fix seems to solve the problem on Samsung S8 about the mismatches and clEnqueueMapBuffer.
Change-Id: I51687ec94bb897af2698ceab1133c988821e4c4e
Reviewed-on: http://mpd-gerrit.cambridge.arm.com/90601
Reviewed-by: Anthony Barbier <anthony.barbier@arm.com>
Tested-by: Kaizen <jeremy.johnson+kaizengerrit@arm.com>
| -rw-r--r-- | src/core/CL/cl_kernels/direct_convolution1x1.cl | 2 | ||||
| -rw-r--r-- | src/core/CL/cl_kernels/direct_convolution3x3.cl | 2 | ||||
| -rw-r--r-- | src/core/CL/cl_kernels/direct_convolution5x5.cl | 2 |
3 files changed, 3 insertions, 3 deletions
diff --git a/src/core/CL/cl_kernels/direct_convolution1x1.cl b/src/core/CL/cl_kernels/direct_convolution1x1.cl index fb516ddd6e..7b73b85eac 100644 --- a/src/core/CL/cl_kernels/direct_convolution1x1.cl +++ b/src/core/CL/cl_kernels/direct_convolution1x1.cl @@ -191,7 +191,7 @@ __kernel void direct_convolution1x1( weights.ptr += z_index * weights_stride_w; - for(int d = 0; d < WEIGHTS_DEPTH; ++d) + for(volatile int d = 0; d < WEIGHTS_DEPTH; ++d) { DATA_TYPE weight = *(__global DATA_TYPE *)weights.ptr; VEC_DATA_TYPE(DATA_TYPE, 8) diff --git a/src/core/CL/cl_kernels/direct_convolution3x3.cl b/src/core/CL/cl_kernels/direct_convolution3x3.cl index d094eca115..1420d7c873 100644 --- a/src/core/CL/cl_kernels/direct_convolution3x3.cl +++ b/src/core/CL/cl_kernels/direct_convolution3x3.cl @@ -138,7 +138,7 @@ __kernel void direct_convolution3x3( const int kernel_index = get_global_id(2); weights_addr += kernel_index * weights_stride_w; - for(int d = 0; d < WEIGHTS_DEPTH; ++d) + for(volatile int d = 0; d < WEIGHTS_DEPTH; ++d) { CONVOLUTION1x3(pixels0, (__global DATA_TYPE *)(src_addr + 0 * src_stride_y), (__global DATA_TYPE *)(weights_addr + 0 * weights_stride_y)); CONVOLUTION1x3(pixels0, (__global DATA_TYPE *)(src_addr + 1 * src_stride_y), (__global DATA_TYPE *)(weights_addr + 1 * weights_stride_y)); diff --git a/src/core/CL/cl_kernels/direct_convolution5x5.cl b/src/core/CL/cl_kernels/direct_convolution5x5.cl index 496da97a09..6fdd019a14 100644 --- a/src/core/CL/cl_kernels/direct_convolution5x5.cl +++ b/src/core/CL/cl_kernels/direct_convolution5x5.cl @@ -127,7 +127,7 @@ __kernel void direct_convolution5x5( const int kernel_index = get_global_id(2); weights_addr += kernel_index * weights_stride_w; - for(int d = 0; d < WEIGHTS_DEPTH; ++d) + for(volatile int d = 0; d < WEIGHTS_DEPTH; ++d) { CONVOLUTION1x5(pixels0, (__global DATA_TYPE *)src_addr, (__global DATA_TYPE *)weights_addr); CONVOLUTION1x5(pixels0, (__global DATA_TYPE *)(src_addr + 1 * src_stride_y), (__global DATA_TYPE *)(weights_addr + 1 * weights_stride_y)); |