From 5b2c0168340a8b4e67def8f52885ab21b4bbcadc Mon Sep 17 00:00:00 2001 From: Alexander Alekhin Date: Fri, 3 Sep 2021 02:38:53 +0000 Subject: [PATCH] dnn(ocl): avoid out of buffer access in copyWeightsSwizzled --- modules/dnn/src/opencl/conv_spatial_helper.cl | 12 ++++++++---- 1 file changed, 8 insertions(+), 4 deletions(-) diff --git a/modules/dnn/src/opencl/conv_spatial_helper.cl b/modules/dnn/src/opencl/conv_spatial_helper.cl index 33d9db57c8..660d085956 100644 --- a/modules/dnn/src/opencl/conv_spatial_helper.cl +++ b/modules/dnn/src/opencl/conv_spatial_helper.cl @@ -62,8 +62,8 @@ __kernel void TEMPLATE(copyWeightsSwizzled, Dtype) //Original location //Output location - int outputSublayer = channels / swizzleFactor; - int outputSublayerIndex = channels % swizzleFactor; + //int outputSublayer = channels / swizzleFactor; + //int outputSublayerIndex = channels % swizzleFactor; int filter = sX / (kernel_w*kernel_h*channels); int kernel_X = sX % kernel_w; @@ -73,6 +73,10 @@ __kernel void TEMPLATE(copyWeightsSwizzled, Dtype) int FP = filter / swizzleFactor; int F1 = filter % swizzleFactor; - weightOut[FP*(kernel_w*kernel_h*channels*swizzleFactor) + kernel_C*(kernel_w*kernel_h*swizzleFactor) + kernel_Y*(kernel_w*swizzleFactor) + kernel_X*swizzleFactor + F1] - = weightIn[filter*(kernel_w*kernel_h*channels) + kernel_C*(kernel_w*kernel_h) + kernel_Y*kernel_w + kernel_X]; + int idxOut = FP*(kernel_w*kernel_h*channels*swizzleFactor) + kernel_C*(kernel_w*kernel_h*swizzleFactor) + kernel_Y*(kernel_w*swizzleFactor) + kernel_X*swizzleFactor + F1; + int idxIn = filter*(kernel_w*kernel_h*channels) + kernel_C*(kernel_w*kernel_h) + kernel_Y*kernel_w + kernel_X; + + // idxIn is not valid if (filter >= outputs) - no data for these elements. Output alignment gaps are filled by zeros + Dtype v = (filter < outputs) ? weightIn[idxIn] : (Dtype)0; + weightOut[idxOut] = v; }