diff options
| author | Georgios Pinitas <georgios.pinitas@arm.com> | 2018-05-14 12:00:05 +0100 |
|---|---|---|
| committer | Anthony Barbier <anthony.barbier@arm.com> | 2018-11-02 16:52:19 +0000 |
| commit | 8422558b946dc783f21241ed1226e601d314bd8b (patch) | |
| tree | bd79b5a987dfb58b3dc4bf5297c25a6b1f346148 /src/core | |
| parent | 393fa4c87c84356132303170d1b9ce9a45b3c3bf (diff) | |
| download | ComputeLibrary-8422558b946dc783f21241ed1226e601d314bd8b.tar.gz | |
COMPMID-1150 : (OCLGrind) Kernel compilation error and assertion
-Multiple definitions of COLS_MTX_B in gemm.cl one for FP32 and one for
FP16.
-GEMMTranspose1xWKernel invalid check fro small window sizes.
Change-Id: I9c7ddd3577aec9afc702731ca27a1e10d6eddb81
Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/131023
Tested-by: Jenkins <bsgcomp@arm.com>
Reviewed-by: Gian Marco Iodice <gianmarco.iodice@arm.com>
Reviewed-by: Anthony Barbier <anthony.barbier@arm.com>
Diffstat (limited to 'src/core')
| -rw-r--r-- | src/core/CL/cl_kernels/gemm.cl | 7 | ||||
| -rw-r--r-- | src/core/CL/kernels/CLGEMMTranspose1xWKernel.cpp | 5 |
2 files changed, 7 insertions, 5 deletions
diff --git a/src/core/CL/cl_kernels/gemm.cl b/src/core/CL/cl_kernels/gemm.cl index 7215f5811f..69bc09f023 100644 --- a/src/core/CL/cl_kernels/gemm.cl +++ b/src/core/CL/cl_kernels/gemm.cl @@ -543,6 +543,9 @@ __kernel void gemm_mm_interleaved_transposed_f32_bifrost(IMAGE_DECLARATION(src0) vstore4((float4)(c30, c31, c32, c33), 0, (__global float *)(dst_addr + 3 * dst_stride_y)); } +// Undefine local defines +#undef COLS_MTX_B + #if defined(ARM_COMPUTE_OPENCL_FP16_ENABLED) /** This OpenCL kernel computes the matrix multiplication between matrix A (src0) and matrix B (src1) * Matrix A and matrix B must be reshaped respectively with @ref gemm_interleave4x4_16bit and @ref gemm_transpose1x8 before running the matrix multiplication @@ -879,6 +882,10 @@ __kernel void gemm_mm_interleaved_transposed_f16_bifrost(IMAGE_DECLARATION(src0) vstore8(c20, 0, (__global half *)(dst_addr + 2 * dst_stride_y)); vstore8(c30, 0, (__global half *)(dst_addr + 3 * dst_stride_y)); } + +// Undefine local defines +#undef COLS_MTX_B + #endif // defined(ARM_COMPUTE_OPENCL_FP16_ENABLED) #if defined(FIXED_POINT_POSITION) diff --git a/src/core/CL/kernels/CLGEMMTranspose1xWKernel.cpp b/src/core/CL/kernels/CLGEMMTranspose1xWKernel.cpp index a9618609e6..05a20fd604 100644 --- a/src/core/CL/kernels/CLGEMMTranspose1xWKernel.cpp +++ b/src/core/CL/kernels/CLGEMMTranspose1xWKernel.cpp @@ -71,11 +71,6 @@ std::pair<Status, Window> validate_and_configure_window(ITensorInfo *input, ITen // Configure kernel window Window win = calculate_max_window(*input, Steps(num_elems_processed_per_iteration)); - if((win.x().end() / scale_x) == 0) - { - return std::make_pair(ARM_COMPUTE_CREATE_ERROR(ErrorCode::RUNTIME_ERROR, "Transposed shape would be 0 in the second dimension"), win); - } - AccessWindowHorizontal input_access(input, 0, num_elems_processed_per_iteration); // Output tensor auto inizialitation if not yet initialized |