diff --git a/mace/kernels/opencl/cl/conv_2d_1x1.cl b/mace/kernels/opencl/cl/conv_2d_1x1.cl index 6007cdf6ddc183cc7e8ba8b258079f4637a32ec2..c9d49525723285a25410b2a976d0b95de0319c05 100644 --- a/mace/kernels/opencl/cl/conv_2d_1x1.cl +++ b/mace/kernels/opencl/cl/conv_2d_1x1.cl @@ -1,29 +1,5 @@ #include -__kernel void conv_2d_1x1_naive(__global const float *input, /* n, c, h, w */ - __global const float *filter, /* o, i, kh, kw */ - __global const float *bias, /* o */ - __global float *output, /* n, c, h, w */ - __private const int in_chan_num) { - const int batch = get_global_id(0); - const int channel = get_global_id(1); - const int channels = get_global_size(1); - const int pixel = get_global_id(2); - const int pixels = get_global_size(2); - - float *output_ptr = output + (batch * channels + channel) * pixels; - output_ptr[pixel] = bias[channel]; - - for (int inc = 0; inc < in_chan_num; ++inc) { - const float *input_ptr = input + (batch * in_chan_num + inc) * pixels + pixel; - const float weights = filter[channel * in_chan_num + inc]; - float in = input_ptr[0]; - float out = output_ptr[0]; - out += in * weights; - output_ptr[0] = out; - } -} - #define vec_conv_2d_1x1_s1 \ VEC_DATA_TYPE(DATA_TYPE,4) in0 = vload4(0, input_ptr); \ VEC_DATA_TYPE(DATA_TYPE,4) in1 = vload4(0, input_ptr + in_pixel); \ diff --git a/mace/ops/ops_test_util.h b/mace/ops/ops_test_util.h index 16b1158850f83043d828a1a93ae677f5b0c29ae1..bbe2d1ab37109cdae5618fee35080a89be5ae951 100644 --- a/mace/ops/ops_test_util.h +++ b/mace/ops/ops_test_util.h @@ -107,32 +107,6 @@ class OpsTestNet { memcpy(input_data, data.data(), data.size() * sizeof(T)); } - template - void AddInputImageFromArray(const std::string &name, - const std::vector &shape, - const std::vector &data) { - Tensor *input = - ws_.CreateTensor(name, GetDeviceAllocator(D), DataTypeToEnum::v()); - std::vector image_shape; - input->ResizeImage(shape, image_shape); - Tensor::MappingGuard input_mapper(input); - T *input_data = input->mutable_data(); - MACE_CHECK(static_cast(input->size()) == data.size()); - const T *data_ptr = data.data(); - const int type_size = sizeof(T); - const int row_pitch = shape[3]; - auto mapped_image_pitch = input_mapper.mapped_image_pitch(); - const size_t slice_size = mapped_image_pitch[1] / sizeof(T); - for (int c = 0; c < shape[0] * shape[1]; ++c) { - T *input_ptr = input_data + c * slice_size; - for (int h = 0; h < shape[2]; ++h) { - memcpy(input_ptr, data_ptr, row_pitch * type_size); - input_ptr += mapped_image_pitch[0] / sizeof(T); - data_ptr += row_pitch; - } - } - } - template void AddRepeatedInput(const std::string &name, const std::vector &shape,