dnn(ocl): fix out of bounds reading in case of filter >= outputs

alalek · alalek · commit 7e83e398791c · 2021-09-03T18:33:35.000Z
diff --git a/modules/dnn/src/opencl/conv_spatial_helper.cl b/modules/dnn/src/opencl/conv_spatial_helper.cl
@@ -84,5 +84,7 @@ __kernel void TEMPLATE(copyWeightsSwizzled, Dtype)
   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;
 
-  weightOut[idxOut] = weightIn[idxIn];
+  // idxIn is not valid if (filter >= outputs) - no data for these elements. Output gaps are filled by zeros
+  Dtype v = (filter < outputs) ? weightIn[idxIn] : (Dtype)0;
+  weightOut[idxOut] = v;
 }

Original file line number	Diff line number	Diff line change
`@@ -84,5 +84,7 @@ __kernel void TEMPLATE(copyWeightsSwizzled, Dtype)`
`84`	`84`	`int idxOut = FP(kernel_wkernel_hchannelsswizzleFactor) + kernel_C(kernel_wkernel_hswizzleFactor) + kernel_Y(kernel_wswizzleFactor) + kernel_XswizzleFactor + F1;`
`85`	`85`	`int idxIn = filter(kernel_wkernel_hchannels) + kernel_C(kernel_wkernel_h) + kernel_Ykernel_w + kernel_X;`
`86`	`86`
`87`		`- weightOut[idxOut] = weightIn[idxIn];`
	`87`	`+ // idxIn is not valid if (filter >= outputs) - no data for these elements. Output gaps are filled by zeros`
	`88`	`+ Dtype v = (filter < outputs) ? weightIn[idxIn] : (Dtype)0;`
	`89`	`+ weightOut[idxOut] = v;`
`88`	`90`	`}`