diff options
| author | Joel Liang <joel.liang@arm.com> | 2017-11-10 09:59:19 +0800 |
|---|---|---|
| committer | Anthony Barbier <anthony.barbier@arm.com> | 2018-11-02 16:35:24 +0000 |
| commit | f1f3ebd517089e934cf3f06e64d90619a395ad87 (patch) | |
| tree | 8dac05909b5f522a1c78e0ac4423cb6f65254391 /src/core/CL/cl_kernels | |
| parent | 283c1790da45ab562ecfb2aa7741297191886d85 (diff) | |
| download | ComputeLibrary-f1f3ebd517089e934cf3f06e64d90619a395ad87.tar.gz | |
APPBROWSER-298, APPBROWSER-306: Reimplement the common code of compute shader
The new common code of compute shader is in file helpers_cs.h
Rewrite the direct_convolution1x1.cs and softmax_layer.cs to use the new common code.
It will also remove the dependence of the token pasting operator (##).
We'll remove the "##" support after we rewrite all of the compute shader code.
Change-Id: Icd8553ef6b61ad484a8507590ac8ed499bd47061
Reviewed-on: http://mpd-gerrit.cambridge.arm.com/95455
Tested-by: Kaizen <jeremy.johnson+kaizengerrit@arm.com>
Reviewed-by: Georgios Pinitas <georgios.pinitas@arm.com>
Reviewed-by: Frank Lei <frank.lei@arm.com>
(cherry picked from commit 0a4f83570d261f839d9866b68979efe8d7a95883)
Reviewed-on: http://mpd-gerrit.cambridge.arm.com/95601
Reviewed-by: Jim He <jim.he@arm.com>
Diffstat (limited to 'src/core/CL/cl_kernels')
5 files changed, 8 insertions, 8 deletions
diff --git a/src/core/CL/cl_kernels/direct_convolution1x1.cl b/src/core/CL/cl_kernels/direct_convolution1x1.cl index 484bc35ef1..817c261ba2 100644 --- a/src/core/CL/cl_kernels/direct_convolution1x1.cl +++ b/src/core/CL/cl_kernels/direct_convolution1x1.cl @@ -153,7 +153,7 @@ inline VEC_DATA_TYPE(DATA_TYPE, 8) extract_input_stride3_8(__global const DATA_T * @param[in] dst_stride_z Stride of the destination tensor in Z dimension (in bytes) * @param[in] dst_step_z dst_stride_z * number of elements along Z processed per workitem(in bytes) * @param[in] dst_offset_first_element_in_bytes The offset of the first element in the destination tensor - * @param[out] weights_ptr Pointer to the weights tensor. Supported data types: same as @p src_ptr + * @param[in] weights_ptr Pointer to the weights tensor. Supported data types: same as @p src_ptr * @param[in] weights_stride_x Stride of the weights tensor in X dimension (in bytes) * @param[in] weights_step_x weights_stride_x * number of elements along X processed per workitem(in bytes) * @param[in] weights_stride_y Stride of the weights tensor in Y dimension (in bytes) @@ -241,7 +241,7 @@ __kernel void direct_convolution1x1( * @param[in] dst_stride_z Stride of the destination tensor in Z dimension (in bytes) * @param[in] dst_step_z dst_stride_z * number of elements along Z processed per workitem(in bytes) * @param[in] dst_offset_first_element_in_bytes The offset of the first element in the destination tensor - * @param[out] weights_ptr Pointer to the weights tensor. Supported data types: same as @p src_ptr + * @param[in] weights_ptr Pointer to the weights tensor. Supported data types: same as @p src_ptr * @param[in] weights_stride_x Stride of the weights tensor in X dimension (in bytes) * @param[in] weights_step_x weights_stride_x * number of elements along X processed per workitem(in bytes) * @param[in] weights_stride_y Stride of the weights tensor in Y dimension (in bytes) diff --git a/src/core/CL/cl_kernels/direct_convolution3x3.cl b/src/core/CL/cl_kernels/direct_convolution3x3.cl index e6e3007c95..a7abc9ff1d 100644 --- a/src/core/CL/cl_kernels/direct_convolution3x3.cl +++ b/src/core/CL/cl_kernels/direct_convolution3x3.cl @@ -102,7 +102,7 @@ MULQ_SAT_IMPL(qs32x8, qs32x8) * @param[in] dst_stride_z Stride of the destination tensor in Z dimension (in bytes) * @param[in] dst_step_z dst_stride_z * number of elements along Z processed per workitem(in bytes) * @param[in] dst_offset_first_element_in_bytes The offset of the first element in the destination tensor - * @param[out] weights_ptr Pointer to the weights tensor. Supported data types: same as @p src_ptr + * @param[in] weights_ptr Pointer to the weights tensor. Supported data types: same as @p src_ptr * @param[in] weights_stride_x Stride of the weights tensor in X dimension (in bytes) * @param[in] weights_step_x weights_stride_x * number of elements along X processed per workitem(in bytes) * @param[in] weights_stride_y Stride of the weights tensor in Y dimension (in bytes) @@ -198,7 +198,7 @@ __kernel void direct_convolution3x3( * @param[in] dst_stride_z Stride of the destination tensor in Z dimension (in bytes) * @param[in] dst_step_z dst_stride_z * number of elements along Z processed per workitem(in bytes) * @param[in] dst_offset_first_element_in_bytes The offset of the first element in the destination tensor - * @param[out] weights_ptr Pointer to the weights tensor. Supported data types: same as @p src_ptr + * @param[in] weights_ptr Pointer to the weights tensor. Supported data types: same as @p src_ptr * @param[in] weights_stride_x Stride of the weights tensor in X dimension (in bytes) * @param[in] weights_step_x weights_stride_x * number of elements along X processed per workitem(in bytes) * @param[in] weights_stride_y Stride of the weights tensor in Y dimension (in bytes) diff --git a/src/core/CL/cl_kernels/direct_convolution5x5.cl b/src/core/CL/cl_kernels/direct_convolution5x5.cl index 12cf0fb68e..e678f6f51b 100644 --- a/src/core/CL/cl_kernels/direct_convolution5x5.cl +++ b/src/core/CL/cl_kernels/direct_convolution5x5.cl @@ -91,7 +91,7 @@ * @param[in] dst_stride_z Stride of the destination tensor in Z dimension (in bytes) * @param[in] dst_step_z dst_stride_z * number of elements along Z processed per workitem(in bytes) * @param[in] dst_offset_first_element_in_bytes The offset of the first element in the destination tensor - * @param[out] weights_ptr Pointer to the weights tensor. Supported data types: same as @p src_ptr + * @param[in] weights_ptr Pointer to the weights tensor. Supported data types: same as @p src_ptr * @param[in] weights_stride_x Stride of the weights tensor in X dimension (in bytes) * @param[in] weights_step_x weights_stride_x * number of elements along X processed per workitem(in bytes) * @param[in] weights_stride_y Stride of the weights tensor in Y dimension (in bytes) @@ -197,7 +197,7 @@ __kernel void direct_convolution5x5( * @param[in] dst_stride_z Stride of the destination tensor in Z dimension (in bytes) * @param[in] dst_step_z dst_stride_z * number of elements along Z processed per workitem(in bytes) * @param[in] dst_offset_first_element_in_bytes The offset of the first element in the destination tensor - * @param[out] weights_ptr Pointer to the weights tensor. Supported data types: same as @p src_ptr + * @param[in] weights_ptr Pointer to the weights tensor. Supported data types: same as @p src_ptr * @param[in] weights_stride_x Stride of the weights tensor in X dimension (in bytes) * @param[in] weights_step_x weights_stride_x * number of elements along X processed per workitem(in bytes) * @param[in] weights_stride_y Stride of the weights tensor in Y dimension (in bytes) diff --git a/src/core/CL/cl_kernels/direct_convolution_1x1_3x3_5x5_quantized.cl b/src/core/CL/cl_kernels/direct_convolution_1x1_3x3_5x5_quantized.cl index 7a860f2008..c94f81e390 100644 --- a/src/core/CL/cl_kernels/direct_convolution_1x1_3x3_5x5_quantized.cl +++ b/src/core/CL/cl_kernels/direct_convolution_1x1_3x3_5x5_quantized.cl @@ -168,7 +168,7 @@ inline uchar8 extract_input_stride3(__global const uchar *input_pixel) * @param[in] dst_stride_z Stride of the destination tensor in Z dimension (in bytes) * @param[in] dst_step_z dst_stride_z * number of elements along Z processed per workitem(in bytes) * @param[in] dst_offset_first_element_in_bytes The offset of the first element in the destination tensor - * @param[out] weights_ptr Pointer to the weights tensor. Supported data types: same as @p weights_ptr + * @param[in] weights_ptr Pointer to the weights tensor. Supported data types: same as @p weights_ptr * @param[in] weights_stride_x Stride of the weights tensor in X dimension (in bytes) * @param[in] weights_step_x weights_stride_x * number of elements along X processed per workitem(in bytes) * @param[in] weights_stride_y Stride of the weights tensor in Y dimension (in bytes) diff --git a/src/core/CL/cl_kernels/gemv.cl b/src/core/CL/cl_kernels/gemv.cl index 76128f7033..3e38c735fe 100644 --- a/src/core/CL/cl_kernels/gemv.cl +++ b/src/core/CL/cl_kernels/gemv.cl @@ -35,7 +35,7 @@ * @param[in] src_stride_z Stride of the source tensor in Z dimension (in bytes) * @param[in] src_step_z src_stride_z * number of elements along Y processed per workitem(in bytes) * @param[in] src_offset_first_element_in_bytes The offset of the first element in the source tensor - * @param[out] weights_ptr Pointer to the weights tensor. Same as @p src_ptr + * @param[in] weights_ptr Pointer to the weights tensor. Same as @p src_ptr * @param[in] weights_stride_x Stride of the weights tensor in X dimension (in bytes) * @param[in] weights_step_x weights_stride_x * number of elements along X processed per workitem(in bytes) * @param[in] weights_stride_y Stride of the weights tensor in Y dimension (in bytes) |