diff --git a/mace/ops/eltwise_test.cc b/mace/ops/eltwise_test.cc index 47b96bd06412cc789a5e2a9ea900edc396d6be65..0ab4a80470785b0b7c3f81a5ec7fc71396273c53 100644 --- a/mace/ops/eltwise_test.cc +++ b/mace/ops/eltwise_test.cc @@ -803,6 +803,28 @@ TEST_F(EltwiseOpTest, TensorGeneralBroadcastGPU) { TensorGeneralBroadcastEltwise( ops::EltwiseType::SQR_DIFF, {1, 1, 2, 3}, {1, 2, 3, 4, 5, 6}, {1, 1, 2, 1}, {1, 2}, {1, 1, 2, 3}, {0, 1, 4, 4, 9, 16}); + + TensorGeneralBroadcastEltwise( + ops::EltwiseType::SUM, {1, 1, 2, 5}, {0, 1, 2, 3, 4, 5, 6, 7, 8, 9}, + {1, 1, 2, 1}, {1, 2}, {1, 1, 2, 5}, {1, 2, 3, 4, 5, 7, 8, 9, 10, 11}); + TensorGeneralBroadcastEltwise( + ops::EltwiseType::SUB, {1, 1, 2, 5}, {0, 1, 2, 3, 4, 5, 6, 7, 8, 9}, + {1, 1, 2, 1}, {1, 2}, {1, 1, 2, 5}, {-1, 0, 1, 2, 3, 3, 4, 5, 6, 7}); + TensorGeneralBroadcastEltwise( + ops::EltwiseType::PROD, {1, 1, 2, 5}, {0, 1, 2, 3, 4, 5, 6, 7, 8, 9}, + {1, 1, 2, 1}, {1, 2}, {1, 1, 2, 5}, {0, 1, 2, 3, 4, 10, 12, 14, 16, 18}); + TensorGeneralBroadcastEltwise( + ops::EltwiseType::DIV, {1, 1, 2, 5}, {0, 1, 2, 3, 4, 5, 6, 7, 8, 9}, + {1, 1, 2, 1}, {1, 2}, {1, 1, 2, 5}, {0, 1, 2, 3, 4, 2.5, 3, 3.5, 4, 4.5}); + TensorGeneralBroadcastEltwise( + ops::EltwiseType::MIN, {1, 1, 2, 5}, {0, 1, 2, 3, 4, 5, 6, 7, 8, 9}, + {1, 1, 2, 1}, {3, 4}, {1, 1, 2, 5}, {0, 1, 2, 3, 3, 4, 4, 4, 4, 4}); + TensorGeneralBroadcastEltwise( + ops::EltwiseType::MAX, {1, 1, 2, 5}, {0, 1, 2, 3, 4, 5, 6, 7, 8, 9}, + {1, 1, 2, 1}, {3, 4}, {1, 1, 2, 5}, {3, 3, 3, 3, 4, 5, 6, 7, 8, 9}); + TensorGeneralBroadcastEltwise( + ops::EltwiseType::SQR_DIFF, {1, 1, 2, 5}, {0, 1, 2, 3, 4, 5, 6, 7, 8, 9}, + {1, 1, 2, 1}, {2, 3}, {1, 1, 2, 5}, {4, 1, 0, 1, 4, 4, 9, 16, 25, 36}); } TEST_F(EltwiseOpTest, QuantizedSum) { diff --git a/mace/ops/opencl/cl/eltwise.cl b/mace/ops/opencl/cl/eltwise.cl index 5c38a7dc505150ac6f663f4e508e42093e1c0ec6..dec31c56b1d9fb4cd09aeb2936c31ddffaccc9df 100644 --- a/mace/ops/opencl/cl/eltwise.cl +++ b/mace/ops/opencl/cl/eltwise.cl @@ -36,7 +36,7 @@ __kernel void eltwise(OUT_OF_RANGE_PARAMS #elif INPUT_TYPE == 3 DATA_TYPE4 in1 = READ_IMAGET(input1, SAMPLER, (int2)(chan_idx, 0)); #elif INPUT_TYPE == 4 - DATA_TYPE4 tmp = READ_IMAGET(input1, SAMPLER, (int2)(pos, hb)); + DATA_TYPE4 tmp = READ_IMAGET(input1, SAMPLER, (int2)(width_idx, hb)); DATA_TYPE4 in1 = (DATA_TYPE4)(tmp.x, tmp.x, tmp.x, tmp.x); #else DATA_TYPE4 in1 = READ_IMAGET(input1, SAMPLER, (int2)(pos, hb));