diff --git a/src/layer/vulkan/convolution_vulkan.cpp b/src/layer/vulkan/convolution_vulkan.cpp index b49b6af2e..14eefe962 100644 --- a/src/layer/vulkan/convolution_vulkan.cpp +++ b/src/layer/vulkan/convolution_vulkan.cpp @@ -152,136 +152,6 @@ int Convolution_vulkan::create_pipeline(const Option& _opt) bool is_conv1x1s1d1 = kernel_w == 1 && kernel_h == 1 && stride_w == 1 && stride_h == 1 && dilation_w == 1 && dilation_h == 1; bool is_conv3x3s1d1 = kernel_w == 3 && kernel_h == 3 && stride_w == 1 && stride_h == 1 && dilation_w == 1 && dilation_h == 1; - int block_x = 0; - int block_y = 0; - Mat shape_winograd_bordered; - Mat shape_winograd_input_transformed; - Mat shape_winograd_gemm; - Mat shape_winograd_out_bordered; - Mat shape_winograd_bordered_packed; - Mat shape_winograd_input_transformed_packed; - Mat shape_winograd_gemm_packed; - Mat shape_winograd_out_bordered_packed; - if (opt.use_winograd_convolution && is_conv3x3s1d1 && num_input >= 32 && num_output >= 32 && ((elempack == 4 && out_elempack == 4) || (elempack == 8 && out_elempack == 8))) - { - if (out_shape.dims != 0) - { - int outw_bordered = (out_shape.w + 3) / 4 * 4; - int outh_bordered = (out_shape.h + 3) / 4 * 4; - - int w_bordered = outw_bordered + 2; - int h_bordered = outh_bordered + 2; - - block_x = outw_bordered / 4; - block_y = outh_bordered / 4; - - shape_winograd_bordered = Mat(w_bordered, h_bordered, shape.c, (void*)0); - shape_winograd_input_transformed = Mat(block_x * block_y, shape.c, 36, (void*)0); - shape_winograd_gemm = Mat(block_x * block_y, out_shape.c, 36, (void*)0); - shape_winograd_out_bordered = Mat(outw_bordered, outh_bordered, out_shape.c, (void*)0); - } - - if (shape_winograd_bordered.dims == 3) shape_winograd_bordered_packed = Mat(shape_winograd_bordered.w, shape_winograd_bordered.h, shape_winograd_bordered.c / elempack, (void*)0, elemsize, elempack); - - if (shape_winograd_input_transformed.dims == 3) shape_winograd_input_transformed_packed = Mat(shape_winograd_input_transformed.w, shape_winograd_input_transformed.h / elempack, 36, (void*)0, elemsize, elempack); - - if (shape_winograd_gemm.dims == 3) shape_winograd_gemm_packed = Mat(shape_winograd_gemm.w, shape_winograd_gemm.h / out_elempack, 36, (void*)0, out_elemsize, out_elempack); - - if (shape_winograd_out_bordered.dims == 3) shape_winograd_out_bordered_packed = Mat(shape_winograd_out_bordered.w, shape_winograd_out_bordered.h, shape_winograd_out_bordered.c / out_elempack, (void*)0, out_elemsize, out_elempack); - - // check blob shape - if (!vkdev->shape_support_image_storage(shape_winograd_bordered_packed) - || !vkdev->shape_support_image_storage(shape_winograd_input_transformed_packed) - || !vkdev->shape_support_image_storage(shape_winograd_gemm_packed) - || !vkdev->shape_support_image_storage(shape_winograd_out_bordered_packed)) - { - support_image_storage = false; - opt.use_image_storage = false; - } - - Mat weight_data_packed_tm(num_input / elempack, num_output / out_elempack, 36, (size_t)4 * elempack * out_elempack, elempack * out_elempack); - if (!vkdev->shape_support_image_storage(weight_data_packed_tm)) - { - support_image_storage = false; - opt.use_image_storage = false; - } - - if (vkdev->info.vendor_id() == 0x5143 && vkdev->info.api_version() < VK_MAKE_VERSION(1, 0, 66)) - { - // FIXME workaround qcom adreno image shader produce wrong result on old drivers - support_image_storage = false; - opt.use_image_storage = false; - } - } - else if (opt.use_winograd_convolution && is_conv3x3s1d1 && num_input >= 16 && num_output >= 16 && ((elempack == 4 && out_elempack == 4) || (elempack == 8 && out_elempack == 8))) - { - if (out_shape.dims != 0) - { - int outw_bordered = (out_shape.w + 1) / 2 * 2; - int outh_bordered = (out_shape.h + 1) / 2 * 2; - - int w_bordered = outw_bordered + 2; - int h_bordered = outh_bordered + 2; - - block_x = outw_bordered / 2; - block_y = outh_bordered / 2; - - shape_winograd_bordered = Mat(w_bordered, h_bordered, shape.c, (void*)0); - shape_winograd_input_transformed = Mat(block_x * block_y, shape.c, 16, (void*)0); - shape_winograd_gemm = Mat(block_x * block_y, out_shape.c, 16, (void*)0); - shape_winograd_out_bordered = Mat(outw_bordered, outh_bordered, out_shape.c, (void*)0); - } - - if (shape_winograd_bordered.dims == 3) shape_winograd_bordered_packed = Mat(shape_winograd_bordered.w, shape_winograd_bordered.h, shape_winograd_bordered.c / elempack, (void*)0, elemsize, elempack); - - if (shape_winograd_input_transformed.dims == 3) shape_winograd_input_transformed_packed = Mat(shape_winograd_input_transformed.w, shape_winograd_input_transformed.h / elempack, 16, (void*)0, elemsize, elempack); - - if (shape_winograd_gemm.dims == 3) shape_winograd_gemm_packed = Mat(shape_winograd_gemm.w, shape_winograd_gemm.h / out_elempack, 16, (void*)0, out_elemsize, out_elempack); - - if (shape_winograd_out_bordered.dims == 3) shape_winograd_out_bordered_packed = Mat(shape_winograd_out_bordered.w, shape_winograd_out_bordered.h, shape_winograd_out_bordered.c / out_elempack, (void*)0, out_elemsize, out_elempack); - - // check blob shape - if (!vkdev->shape_support_image_storage(shape_winograd_bordered_packed) - || !vkdev->shape_support_image_storage(shape_winograd_input_transformed_packed) - || !vkdev->shape_support_image_storage(shape_winograd_gemm_packed) - || !vkdev->shape_support_image_storage(shape_winograd_out_bordered_packed)) - { - support_image_storage = false; - opt.use_image_storage = false; - } - - Mat weight_data_packed_tm(num_input / elempack, num_output / out_elempack, 16, (size_t)4 * elempack * out_elempack, elempack * out_elempack); - if (!vkdev->shape_support_image_storage(weight_data_packed_tm)) - { - support_image_storage = false; - opt.use_image_storage = false; - } - - if (vkdev->info.vendor_id() == 0x5143 && vkdev->info.api_version() < VK_MAKE_VERSION(1, 0, 66)) - { - // FIXME workaround qcom adreno image shader produce wrong result on old drivers - support_image_storage = false; - opt.use_image_storage = false; - } - } - else - { - // check blob shape - if (!vkdev->shape_support_image_storage(shape_bordered_packed) || !vkdev->shape_support_image_storage(out_shape_packed)) - { - support_image_storage = false; - opt.use_image_storage = false; - } - - // check weight shape - Mat weight_data_packed(maxk, num_input / elempack, num_output / out_elempack, (void*)0, (size_t)4 * elempack * out_elempack, elempack * out_elempack); - if (!vkdev->shape_support_image_storage(weight_data_packed)) - { - support_image_storage = false; - opt.use_image_storage = false; - } - } - { padding = ncnn::create_layer(ncnn::LayerType::Padding); padding->vkdev = vkdev; @@ -344,18 +214,12 @@ int Convolution_vulkan::create_pipeline(const Option& _opt) } pipeline_convolution_1x1s1d1->create(shader_type_index, opt, specializations); } - else if (opt.use_winograd_convolution && is_conv3x3s1d1 && num_input >= 32 && num_output >= 32 && ((elempack == 4 && out_elempack == 4) || (elempack == 8 && out_elempack == 8))) + if (opt.use_winograd_convolution && is_conv3x3s1d1 && num_input >= 16 && num_output >= 16 && ((elempack == 4 && out_elempack == 4) || (elempack == 8 && out_elempack == 8))) { - // winograd43 { winograd_padding = ncnn::create_layer(ncnn::LayerType::Padding); winograd_padding->vkdev = vkdev; - winograd_padding->bottom_shapes.resize(1); - winograd_padding->bottom_shapes[0] = shape_bordered; - winograd_padding->top_shapes.resize(1); - winograd_padding->top_shapes[0] = shape_winograd_bordered; - ncnn::ParamDict pd; pd.set(0, -233); pd.set(1, -233); @@ -373,11 +237,6 @@ int Convolution_vulkan::create_pipeline(const Option& _opt) winograd_crop = ncnn::create_layer(ncnn::LayerType::Crop); winograd_crop->vkdev = vkdev; - winograd_crop->bottom_shapes.resize(1); - winograd_crop->bottom_shapes[0] = shape_winograd_out_bordered; - winograd_crop->top_shapes.resize(1); - winograd_crop->top_shapes[0] = out_shape; - ncnn::ParamDict pd; pd.set(0, -233); pd.set(1, -233); @@ -391,189 +250,283 @@ int Convolution_vulkan::create_pipeline(const Option& _opt) winograd_crop->create_pipeline(opt); } + // winograd43 { - std::vector specializations(0 + 7); - specializations[0 + 0].i = shape_winograd_bordered_packed.w; - specializations[0 + 1].i = shape_winograd_bordered_packed.h; - specializations[0 + 2].i = shape_winograd_bordered_packed.c; - specializations[0 + 3].i = shape_winograd_bordered_packed.cstep; - specializations[0 + 4].i = shape_winograd_input_transformed_packed.cstep; - specializations[0 + 5].i = block_x; - specializations[0 + 6].i = block_y; + int block_x = 0; + int block_y = 0; + Mat shape_winograd_bordered; + Mat shape_winograd_input_transformed; + Mat shape_winograd_gemm; + Mat shape_winograd_out_bordered; + Mat shape_winograd_bordered_packed; + Mat shape_winograd_input_transformed_packed; + Mat shape_winograd_gemm_packed; + Mat shape_winograd_out_bordered_packed; + + if (out_shape.dims != 0) + { + int outw_bordered = (out_shape.w + 3) / 4 * 4; + int outh_bordered = (out_shape.h + 3) / 4 * 4; - int shader_type_index = -1; - if (elempack == 4 && out_elempack == 4) shader_type_index = LayerShaderType::convolution_pack4_3x3s1d1_winograd43_transform_input; - if (elempack == 8 && out_elempack == 8) shader_type_index = LayerShaderType::convolution_pack8_3x3s1d1_winograd43_transform_input; + int w_bordered = outw_bordered + 2; + int h_bordered = outh_bordered + 2; - pipeline_convolution_3x3s1d1_winograd43_transform_input = new Pipeline(vkdev); - pipeline_convolution_3x3s1d1_winograd43_transform_input->set_local_size_xyz(4, 4, 1); - pipeline_convolution_3x3s1d1_winograd43_transform_input->create(shader_type_index, opt, specializations); - } + block_x = outw_bordered / 4; + block_y = outh_bordered / 4; - { - std::vector specializations(1 + 5); - specializations[0].i = 36; - specializations[1 + 0].i = shape_winograd_input_transformed_packed.h; - specializations[1 + 1].i = shape_winograd_input_transformed_packed.cstep; - specializations[1 + 2].i = shape_winograd_gemm_packed.w; - specializations[1 + 3].i = shape_winograd_gemm_packed.h; - specializations[1 + 4].i = shape_winograd_gemm_packed.cstep; + shape_winograd_bordered = Mat(w_bordered, h_bordered, shape.c, (void*)0); + shape_winograd_input_transformed = Mat(block_x * block_y, shape.c, 36, (void*)0); + shape_winograd_gemm = Mat(block_x * block_y, out_shape.c, 36, (void*)0); + shape_winograd_out_bordered = Mat(outw_bordered, outh_bordered, out_shape.c, (void*)0); + } - int shader_type_index = -1; - if (elempack == 4 && out_elempack == 4) shader_type_index = LayerShaderType::convolution_pack4_3x3s1d1_winograd_gemm; - if (elempack == 8 && out_elempack == 8) shader_type_index = LayerShaderType::convolution_pack8_3x3s1d1_winograd_gemm; + if (shape_winograd_bordered.dims == 3) shape_winograd_bordered_packed = Mat(shape_winograd_bordered.w, shape_winograd_bordered.h, shape_winograd_bordered.c / elempack, (void*)0, elemsize, elempack); - pipeline_convolution_3x3s1d1_winograd43_gemm = new Pipeline(vkdev); - if (opt.use_shader_local_memory) + if (shape_winograd_input_transformed.dims == 3) shape_winograd_input_transformed_packed = Mat(shape_winograd_input_transformed.w, shape_winograd_input_transformed.h / elempack, 36, (void*)0, elemsize, elempack); + + if (shape_winograd_gemm.dims == 3) shape_winograd_gemm_packed = Mat(shape_winograd_gemm.w, shape_winograd_gemm.h / out_elempack, 36, (void*)0, out_elemsize, out_elempack); + + if (shape_winograd_out_bordered.dims == 3) shape_winograd_out_bordered_packed = Mat(shape_winograd_out_bordered.w, shape_winograd_out_bordered.h, shape_winograd_out_bordered.c / out_elempack, (void*)0, out_elemsize, out_elempack); + + // check blob shape + if (!vkdev->shape_support_image_storage(shape_winograd_bordered_packed) + || !vkdev->shape_support_image_storage(shape_winograd_input_transformed_packed) + || !vkdev->shape_support_image_storage(shape_winograd_gemm_packed) + || !vkdev->shape_support_image_storage(shape_winograd_out_bordered_packed)) { - pipeline_convolution_3x3s1d1_winograd43_gemm->set_local_size_xyz(8, 8, 1); + support_image_storage = false; + opt.use_image_storage = false; } - else + + Mat weight_data_packed_tm(num_input / elempack, num_output / out_elempack, 36, (size_t)4 * elempack * out_elempack, elempack * out_elempack); + if (!vkdev->shape_support_image_storage(weight_data_packed_tm)) { - pipeline_convolution_3x3s1d1_winograd43_gemm->set_local_size_xyz(4, std::min(4, num_output / out_elempack), 4); + support_image_storage = false; + opt.use_image_storage = false; } - pipeline_convolution_3x3s1d1_winograd43_gemm->create(shader_type_index, opt, specializations); - } - - { - std::vector specializations(4 + 7); - specializations[0].i = bias_term; - specializations[1].i = activation_type; - specializations[2].f = activation_params.w >= 1 ? activation_params[0] : 0.f; - specializations[3].f = activation_params.w == 2 ? activation_params[1] : 0.f; - specializations[4 + 0].i = shape_winograd_gemm_packed.c; - specializations[4 + 1].i = shape_winograd_gemm_packed.cstep; - specializations[4 + 2].i = block_x; - specializations[4 + 3].i = block_y; - specializations[4 + 4].i = shape_winograd_out_bordered_packed.w; - specializations[4 + 5].i = shape_winograd_out_bordered_packed.h; - specializations[4 + 6].i = shape_winograd_out_bordered_packed.cstep; - - int shader_type_index = -1; - if (elempack == 4 && out_elempack == 4) shader_type_index = LayerShaderType::convolution_pack4_3x3s1d1_winograd43_transform_output; - if (elempack == 8 && out_elempack == 8) shader_type_index = LayerShaderType::convolution_pack8_3x3s1d1_winograd43_transform_output; - - pipeline_convolution_3x3s1d1_winograd43_transform_output = new Pipeline(vkdev); - pipeline_convolution_3x3s1d1_winograd43_transform_output->set_local_size_xyz(4, 4, 1); - pipeline_convolution_3x3s1d1_winograd43_transform_output->create(shader_type_index, opt, specializations); - } - } - else if (opt.use_winograd_convolution && is_conv3x3s1d1 && num_input >= 16 && num_output >= 16 && ((elempack == 4 && out_elempack == 4) || (elempack == 8 && out_elempack == 8))) - { - // winograd23 - { - winograd_padding = ncnn::create_layer(ncnn::LayerType::Padding); - winograd_padding->vkdev = vkdev; - winograd_padding->bottom_shapes.resize(1); - winograd_padding->bottom_shapes[0] = shape_bordered; - winograd_padding->top_shapes.resize(1); - winograd_padding->top_shapes[0] = shape_winograd_bordered; + if (vkdev->info.vendor_id() == 0x5143 && vkdev->info.api_version() < VK_MAKE_VERSION(1, 0, 66)) + { + // FIXME workaround qcom adreno image shader produce wrong result on old drivers + support_image_storage = false; + opt.use_image_storage = false; + } - ncnn::ParamDict pd; - pd.set(0, -233); - pd.set(1, -233); - pd.set(2, -233); - pd.set(3, -233); - pd.set(4, 0); - pd.set(5, 0.f); + { + std::vector specializations(0 + 7); + specializations[0 + 0].i = shape_winograd_bordered_packed.w; + specializations[0 + 1].i = shape_winograd_bordered_packed.h; + specializations[0 + 2].i = shape_winograd_bordered_packed.c; + specializations[0 + 3].i = shape_winograd_bordered_packed.cstep; + specializations[0 + 4].i = shape_winograd_input_transformed_packed.cstep; + specializations[0 + 5].i = block_x; + specializations[0 + 6].i = block_y; + + int shader_type_index = -1; + if (elempack == 4 && out_elempack == 4) shader_type_index = LayerShaderType::convolution_pack4_3x3s1d1_winograd43_transform_input; + if (elempack == 8 && out_elempack == 8) shader_type_index = LayerShaderType::convolution_pack8_3x3s1d1_winograd43_transform_input; + + pipeline_convolution_3x3s1d1_winograd43_transform_input = new Pipeline(vkdev); + pipeline_convolution_3x3s1d1_winograd43_transform_input->set_local_size_xyz(4, 4, 1); + pipeline_convolution_3x3s1d1_winograd43_transform_input->create(shader_type_index, opt, specializations); + } - winograd_padding->load_param(pd); + { + std::vector specializations(1 + 5); + specializations[0].i = 36; + specializations[1 + 0].i = shape_winograd_input_transformed_packed.h; + specializations[1 + 1].i = shape_winograd_input_transformed_packed.cstep; + specializations[1 + 2].i = shape_winograd_gemm_packed.w; + specializations[1 + 3].i = shape_winograd_gemm_packed.h; + specializations[1 + 4].i = shape_winograd_gemm_packed.cstep; + + int shader_type_index = -1; + if (elempack == 4 && out_elempack == 4) shader_type_index = LayerShaderType::convolution_pack4_3x3s1d1_winograd_gemm; + if (elempack == 8 && out_elempack == 8) shader_type_index = LayerShaderType::convolution_pack8_3x3s1d1_winograd_gemm; + + pipeline_convolution_3x3s1d1_winograd43_gemm = new Pipeline(vkdev); + if (opt.use_shader_local_memory) + { + pipeline_convolution_3x3s1d1_winograd43_gemm->set_local_size_xyz(8, 8, 1); + } + else + { + pipeline_convolution_3x3s1d1_winograd43_gemm->set_local_size_xyz(4, std::min(4, num_output / out_elempack), 4); + } + pipeline_convolution_3x3s1d1_winograd43_gemm->create(shader_type_index, opt, specializations); + } - winograd_padding->create_pipeline(opt); + { + std::vector specializations(4 + 7); + specializations[0].i = bias_term; + specializations[1].i = activation_type; + specializations[2].f = activation_params.w >= 1 ? activation_params[0] : 0.f; + specializations[3].f = activation_params.w == 2 ? activation_params[1] : 0.f; + specializations[4 + 0].i = shape_winograd_gemm_packed.c; + specializations[4 + 1].i = shape_winograd_gemm_packed.cstep; + specializations[4 + 2].i = block_x; + specializations[4 + 3].i = block_y; + specializations[4 + 4].i = shape_winograd_out_bordered_packed.w; + specializations[4 + 5].i = shape_winograd_out_bordered_packed.h; + specializations[4 + 6].i = shape_winograd_out_bordered_packed.cstep; + + int shader_type_index = -1; + if (elempack == 4 && out_elempack == 4) shader_type_index = LayerShaderType::convolution_pack4_3x3s1d1_winograd43_transform_output; + if (elempack == 8 && out_elempack == 8) shader_type_index = LayerShaderType::convolution_pack8_3x3s1d1_winograd43_transform_output; + + pipeline_convolution_3x3s1d1_winograd43_transform_output = new Pipeline(vkdev); + pipeline_convolution_3x3s1d1_winograd43_transform_output->set_local_size_xyz(4, 4, 1); + pipeline_convolution_3x3s1d1_winograd43_transform_output->create(shader_type_index, opt, specializations); + } } + // winograd23 { - winograd_crop = ncnn::create_layer(ncnn::LayerType::Crop); - winograd_crop->vkdev = vkdev; - - winograd_crop->bottom_shapes.resize(1); - winograd_crop->bottom_shapes[0] = shape_winograd_out_bordered; - winograd_crop->top_shapes.resize(1); - winograd_crop->top_shapes[0] = out_shape; + int block_x = 0; + int block_y = 0; + Mat shape_winograd_bordered; + Mat shape_winograd_input_transformed; + Mat shape_winograd_gemm; + Mat shape_winograd_out_bordered; + Mat shape_winograd_bordered_packed; + Mat shape_winograd_input_transformed_packed; + Mat shape_winograd_gemm_packed; + Mat shape_winograd_out_bordered_packed; + + if (out_shape.dims != 0) + { + int outw_bordered = (out_shape.w + 1) / 2 * 2; + int outh_bordered = (out_shape.h + 1) / 2 * 2; - ncnn::ParamDict pd; - pd.set(0, -233); - pd.set(1, -233); - pd.set(2, -233); - pd.set(3, 0); - pd.set(4, 0); - pd.set(5, 0); + int w_bordered = outw_bordered + 2; + int h_bordered = outh_bordered + 2; - winograd_crop->load_param(pd); + block_x = outw_bordered / 2; + block_y = outh_bordered / 2; - winograd_crop->create_pipeline(opt); - } + shape_winograd_bordered = Mat(w_bordered, h_bordered, shape.c, (void*)0); + shape_winograd_input_transformed = Mat(block_x * block_y, shape.c, 16, (void*)0); + shape_winograd_gemm = Mat(block_x * block_y, out_shape.c, 16, (void*)0); + shape_winograd_out_bordered = Mat(outw_bordered, outh_bordered, out_shape.c, (void*)0); + } - { - std::vector specializations(0 + 7); - specializations[0 + 0].i = shape_winograd_bordered_packed.w; - specializations[0 + 1].i = shape_winograd_bordered_packed.h; - specializations[0 + 2].i = shape_winograd_bordered_packed.c; - specializations[0 + 3].i = shape_winograd_bordered_packed.cstep; - specializations[0 + 4].i = shape_winograd_input_transformed_packed.cstep; - specializations[0 + 5].i = block_x; - specializations[0 + 6].i = block_y; + if (shape_winograd_bordered.dims == 3) shape_winograd_bordered_packed = Mat(shape_winograd_bordered.w, shape_winograd_bordered.h, shape_winograd_bordered.c / elempack, (void*)0, elemsize, elempack); - int shader_type_index = -1; - if (elempack == 4 && out_elempack == 4) shader_type_index = LayerShaderType::convolution_pack4_3x3s1d1_winograd23_transform_input; - if (elempack == 8 && out_elempack == 8) shader_type_index = LayerShaderType::convolution_pack8_3x3s1d1_winograd23_transform_input; + if (shape_winograd_input_transformed.dims == 3) shape_winograd_input_transformed_packed = Mat(shape_winograd_input_transformed.w, shape_winograd_input_transformed.h / elempack, 16, (void*)0, elemsize, elempack); - pipeline_convolution_3x3s1d1_winograd23_transform_input = new Pipeline(vkdev); - pipeline_convolution_3x3s1d1_winograd23_transform_input->set_local_size_xyz(8, 8, 1); - pipeline_convolution_3x3s1d1_winograd23_transform_input->create(shader_type_index, opt, specializations); - } + if (shape_winograd_gemm.dims == 3) shape_winograd_gemm_packed = Mat(shape_winograd_gemm.w, shape_winograd_gemm.h / out_elempack, 16, (void*)0, out_elemsize, out_elempack); - { - std::vector specializations(1 + 5); - specializations[0].i = 16; - specializations[1 + 0].i = shape_winograd_input_transformed_packed.h; - specializations[1 + 1].i = shape_winograd_input_transformed_packed.cstep; - specializations[1 + 2].i = shape_winograd_gemm_packed.w; - specializations[1 + 3].i = shape_winograd_gemm_packed.h; - specializations[1 + 4].i = shape_winograd_gemm_packed.cstep; + if (shape_winograd_out_bordered.dims == 3) shape_winograd_out_bordered_packed = Mat(shape_winograd_out_bordered.w, shape_winograd_out_bordered.h, shape_winograd_out_bordered.c / out_elempack, (void*)0, out_elemsize, out_elempack); - int shader_type_index = -1; - if (elempack == 4 && out_elempack == 4) shader_type_index = LayerShaderType::convolution_pack4_3x3s1d1_winograd_gemm; - if (elempack == 8 && out_elempack == 8) shader_type_index = LayerShaderType::convolution_pack8_3x3s1d1_winograd_gemm; + // check blob shape + if (!vkdev->shape_support_image_storage(shape_winograd_bordered_packed) + || !vkdev->shape_support_image_storage(shape_winograd_input_transformed_packed) + || !vkdev->shape_support_image_storage(shape_winograd_gemm_packed) + || !vkdev->shape_support_image_storage(shape_winograd_out_bordered_packed)) + { + support_image_storage = false; + opt.use_image_storage = false; + } - pipeline_convolution_3x3s1d1_winograd23_gemm = new Pipeline(vkdev); - if (opt.use_shader_local_memory) + Mat weight_data_packed_tm(num_input / elempack, num_output / out_elempack, 16, (size_t)4 * elempack * out_elempack, elempack * out_elempack); + if (!vkdev->shape_support_image_storage(weight_data_packed_tm)) { - pipeline_convolution_3x3s1d1_winograd23_gemm->set_local_size_xyz(8, 8, 1); + support_image_storage = false; + opt.use_image_storage = false; } - else + + if (vkdev->info.vendor_id() == 0x5143 && vkdev->info.api_version() < VK_MAKE_VERSION(1, 0, 66)) { - pipeline_convolution_3x3s1d1_winograd23_gemm->set_local_size_xyz(4, std::min(4, num_output / out_elempack), 4); + // FIXME workaround qcom adreno image shader produce wrong result on old drivers + support_image_storage = false; + opt.use_image_storage = false; } - pipeline_convolution_3x3s1d1_winograd23_gemm->create(shader_type_index, opt, specializations); - } - { - std::vector specializations(4 + 7); - specializations[0].i = bias_term; - specializations[1].i = activation_type; - specializations[2].f = activation_params.w >= 1 ? activation_params[0] : 0.f; - specializations[3].f = activation_params.w == 2 ? activation_params[1] : 0.f; - specializations[4 + 0].i = shape_winograd_gemm_packed.h; - specializations[4 + 1].i = shape_winograd_gemm_packed.cstep; - specializations[4 + 2].i = block_x; - specializations[4 + 3].i = block_y; - specializations[4 + 4].i = shape_winograd_out_bordered_packed.w; - specializations[4 + 5].i = shape_winograd_out_bordered_packed.h; - specializations[4 + 6].i = shape_winograd_out_bordered_packed.cstep; + { + std::vector specializations(0 + 7); + specializations[0 + 0].i = shape_winograd_bordered_packed.w; + specializations[0 + 1].i = shape_winograd_bordered_packed.h; + specializations[0 + 2].i = shape_winograd_bordered_packed.c; + specializations[0 + 3].i = shape_winograd_bordered_packed.cstep; + specializations[0 + 4].i = shape_winograd_input_transformed_packed.cstep; + specializations[0 + 5].i = block_x; + specializations[0 + 6].i = block_y; + + int shader_type_index = -1; + if (elempack == 4 && out_elempack == 4) shader_type_index = LayerShaderType::convolution_pack4_3x3s1d1_winograd23_transform_input; + if (elempack == 8 && out_elempack == 8) shader_type_index = LayerShaderType::convolution_pack8_3x3s1d1_winograd23_transform_input; + + pipeline_convolution_3x3s1d1_winograd23_transform_input = new Pipeline(vkdev); + pipeline_convolution_3x3s1d1_winograd23_transform_input->set_local_size_xyz(8, 8, 1); + pipeline_convolution_3x3s1d1_winograd23_transform_input->create(shader_type_index, opt, specializations); + } - int shader_type_index = -1; - if (elempack == 4 && out_elempack == 4) shader_type_index = LayerShaderType::convolution_pack4_3x3s1d1_winograd23_transform_output; - if (elempack == 8 && out_elempack == 8) shader_type_index = LayerShaderType::convolution_pack8_3x3s1d1_winograd23_transform_output; + { + std::vector specializations(1 + 5); + specializations[0].i = 16; + specializations[1 + 0].i = shape_winograd_input_transformed_packed.h; + specializations[1 + 1].i = shape_winograd_input_transformed_packed.cstep; + specializations[1 + 2].i = shape_winograd_gemm_packed.w; + specializations[1 + 3].i = shape_winograd_gemm_packed.h; + specializations[1 + 4].i = shape_winograd_gemm_packed.cstep; + + int shader_type_index = -1; + if (elempack == 4 && out_elempack == 4) shader_type_index = LayerShaderType::convolution_pack4_3x3s1d1_winograd_gemm; + if (elempack == 8 && out_elempack == 8) shader_type_index = LayerShaderType::convolution_pack8_3x3s1d1_winograd_gemm; + + pipeline_convolution_3x3s1d1_winograd23_gemm = new Pipeline(vkdev); + if (opt.use_shader_local_memory) + { + pipeline_convolution_3x3s1d1_winograd23_gemm->set_local_size_xyz(8, 8, 1); + } + else + { + pipeline_convolution_3x3s1d1_winograd23_gemm->set_local_size_xyz(4, std::min(4, num_output / out_elempack), 4); + } + pipeline_convolution_3x3s1d1_winograd23_gemm->create(shader_type_index, opt, specializations); + } - pipeline_convolution_3x3s1d1_winograd23_transform_output = new Pipeline(vkdev); - pipeline_convolution_3x3s1d1_winograd23_transform_output->set_local_size_xyz(8, 8, 1); - pipeline_convolution_3x3s1d1_winograd23_transform_output->create(shader_type_index, opt, specializations); + { + std::vector specializations(4 + 7); + specializations[0].i = bias_term; + specializations[1].i = activation_type; + specializations[2].f = activation_params.w >= 1 ? activation_params[0] : 0.f; + specializations[3].f = activation_params.w == 2 ? activation_params[1] : 0.f; + specializations[4 + 0].i = shape_winograd_gemm_packed.h; + specializations[4 + 1].i = shape_winograd_gemm_packed.cstep; + specializations[4 + 2].i = block_x; + specializations[4 + 3].i = block_y; + specializations[4 + 4].i = shape_winograd_out_bordered_packed.w; + specializations[4 + 5].i = shape_winograd_out_bordered_packed.h; + specializations[4 + 6].i = shape_winograd_out_bordered_packed.cstep; + + int shader_type_index = -1; + if (elempack == 4 && out_elempack == 4) shader_type_index = LayerShaderType::convolution_pack4_3x3s1d1_winograd23_transform_output; + if (elempack == 8 && out_elempack == 8) shader_type_index = LayerShaderType::convolution_pack8_3x3s1d1_winograd23_transform_output; + + pipeline_convolution_3x3s1d1_winograd23_transform_output = new Pipeline(vkdev); + pipeline_convolution_3x3s1d1_winograd23_transform_output->set_local_size_xyz(8, 8, 1); + pipeline_convolution_3x3s1d1_winograd23_transform_output->create(shader_type_index, opt, specializations); + } } } - else if (opt.use_sgemm_convolution && !is_conv1x1s1d1 && num_input >= 16 && num_output >= 16) + if (opt.use_sgemm_convolution && !is_conv1x1s1d1 && num_input >= 16 && num_output >= 16) { + // check blob shape + if (!vkdev->shape_support_image_storage(shape_bordered_packed) || !vkdev->shape_support_image_storage(out_shape_packed)) + { + support_image_storage = false; + opt.use_image_storage = false; + } + + // check weight shape + Mat weight_data_packed(maxk, num_input / elempack, num_output / out_elempack, (void*)0, (size_t)4 * elempack * out_elempack, elempack * out_elempack); + if (!vkdev->shape_support_image_storage(weight_data_packed)) + { + support_image_storage = false; + opt.use_image_storage = false; + } + { std::vector specializations(10 + 8); specializations[0].i = kernel_w; @@ -813,185 +766,187 @@ int Convolution_vulkan::upload_model(VkTransfer& cmd, const Option& opt) bool is_conv3x3s1d1 = kernel_w == 3 && kernel_h == 3 && stride_w == 1 && stride_h == 1 && dilation_w == 1 && dilation_h == 1; - if (opt.use_winograd_convolution && is_conv3x3s1d1 && num_input >= 32 && num_output >= 32 && ((elempack == 4 && out_elempack == 4) || (elempack == 8 && out_elempack == 8))) + if (opt.use_winograd_convolution && is_conv3x3s1d1 && num_input >= 16 && num_output >= 16 && ((elempack == 4 && out_elempack == 4) || (elempack == 8 && out_elempack == 8))) { // winograd43 transform kernel - Mat weight_data_tm; - weight_data_tm.create(6 * 6, num_input, num_output); - - const float ktm[6][3] = { - {1.0f, 0.0f, 0.0f}, - {-2.0f / 3, -2.0f / 3, -2.0f / 3}, - {-2.0f / 3, 2.0f / 3, -2.0f / 3}, - {1.0f / 6, 1.0f / 3, 2.0f / 3}, - {1.0f / 6, -1.0f / 3, 2.0f / 3}, - {0.0f, 0.0f, 4.0f} - }; - - #pragma omp parallel for num_threads(opt.num_threads) - for (int p = 0; p < num_output; p++) { - for (int q = 0; q < num_input; q++) - { - const float* kernel0 = (const float*)weight_data + p * num_input * 9 + q * 9; - float* kernel_tm0 = weight_data_tm.channel(p).row(q); + Mat weight_data_tm; + weight_data_tm.create(6 * 6, num_input, num_output); - // transform kernel - const float* k0 = kernel0; - const float* k1 = kernel0 + 3; - const float* k2 = kernel0 + 6; + const float ktm[6][3] = { + {1.0f, 0.0f, 0.0f}, + {-2.0f / 3, -2.0f / 3, -2.0f / 3}, + {-2.0f / 3, 2.0f / 3, -2.0f / 3}, + {1.0f / 6, 1.0f / 3, 2.0f / 3}, + {1.0f / 6, -1.0f / 3, 2.0f / 3}, + {0.0f, 0.0f, 4.0f} + }; - // h - float tmp[6][3]; - for (int i = 0; i < 6; i++) + #pragma omp parallel for num_threads(opt.num_threads) + for (int p = 0; p < num_output; p++) + { + for (int q = 0; q < num_input; q++) { - tmp[i][0] = k0[0] * ktm[i][0] + k0[1] * ktm[i][1] + k0[2] * ktm[i][2]; - tmp[i][1] = k1[0] * ktm[i][0] + k1[1] * ktm[i][1] + k1[2] * ktm[i][2]; - tmp[i][2] = k2[0] * ktm[i][0] + k2[1] * ktm[i][1] + k2[2] * ktm[i][2]; - } + const float* kernel0 = (const float*)weight_data + p * num_input * 9 + q * 9; + float* kernel_tm0 = weight_data_tm.channel(p).row(q); - // U - for (int j = 0; j < 6; j++) - { - float* tmpp = &tmp[j][0]; + // transform kernel + const float* k0 = kernel0; + const float* k1 = kernel0 + 3; + const float* k2 = kernel0 + 6; + // h + float tmp[6][3]; for (int i = 0; i < 6; i++) { - kernel_tm0[j * 6 + i] = tmpp[0] * ktm[i][0] + tmpp[1] * ktm[i][1] + tmpp[2] * ktm[i][2]; + tmp[i][0] = k0[0] * ktm[i][0] + k0[1] * ktm[i][1] + k0[2] * ktm[i][2]; + tmp[i][1] = k1[0] * ktm[i][0] + k1[1] * ktm[i][1] + k1[2] * ktm[i][2]; + tmp[i][2] = k2[0] * ktm[i][0] + k2[1] * ktm[i][1] + k2[2] * ktm[i][2]; + } + + // U + for (int j = 0; j < 6; j++) + { + float* tmpp = &tmp[j][0]; + + for (int i = 0; i < 6; i++) + { + kernel_tm0[j * 6 + i] = tmpp[0] * ktm[i][0] + tmpp[1] * ktm[i][1] + tmpp[2] * ktm[i][2]; + } } } } - } - // src = 36-inch-outch - // dst = 8a-8b-inch/8a-outch/8b-36 - Mat weight_data_tm_packed; - { - weight_data_tm_packed.create(num_input / elempack, num_output / out_elempack, 36, (size_t)4 * elempack * out_elempack, elempack * out_elempack); - - for (int k = 0; k < 36; k++) + // src = 36-inch-outch + // dst = 8a-8b-inch/8a-outch/8b-36 + Mat weight_data_tm_packed; { - float* g00 = weight_data_tm_packed.channel(k); + weight_data_tm_packed.create(num_input / elempack, num_output / out_elempack, 36, (size_t)4 * elempack * out_elempack, elempack * out_elempack); - for (int q = 0; q + (out_elempack - 1) < num_output; q += out_elempack) + for (int k = 0; k < 36; k++) { - for (int p = 0; p + (elempack - 1) < num_input; p += elempack) + float* g00 = weight_data_tm_packed.channel(k); + + for (int q = 0; q + (out_elempack - 1) < num_output; q += out_elempack) { - for (int i = 0; i < out_elempack; i++) + for (int p = 0; p + (elempack - 1) < num_input; p += elempack) { - const Mat k0 = weight_data_tm.channel(q + i); - - for (int j = 0; j < elempack; j++) + for (int i = 0; i < out_elempack; i++) { - const float* k00 = k0.row(p + j); + const Mat k0 = weight_data_tm.channel(q + i); + + for (int j = 0; j < elempack; j++) + { + const float* k00 = k0.row(p + j); - g00[0] = k00[k]; + g00[0] = k00[k]; - g00++; + g00++; + } } } } } } - } - if (support_image_storage && opt.use_image_storage) - { - cmd.record_upload(weight_data_tm_packed, weight_data_gpu_tm_image, opt); - } - else - { - cmd.record_upload(weight_data_tm_packed, weight_data_gpu_tm, opt); + if (support_image_storage && opt.use_image_storage) + { + cmd.record_upload(weight_data_tm_packed, weight_data_gpu_tm_winograd43_image, opt); + } + else + { + cmd.record_upload(weight_data_tm_packed, weight_data_gpu_tm_winograd43, opt); + } } - } - else if (opt.use_winograd_convolution && is_conv3x3s1d1 && num_input >= 16 && num_output >= 16 && ((elempack == 4 && out_elempack == 4) || (elempack == 8 && out_elempack == 8))) - { + // winograd23 transform kernel - Mat weight_data_tm; - weight_data_tm.create(4 * 4, num_input, num_output); - - // G - const float ktm[4][3] = { - {1.0f, 0.0f, 0.0f}, - {1.0f / 2, 1.0f / 2, 1.0f / 2}, - {1.0f / 2, -1.0f / 2, 1.0f / 2}, - {0.0f, 0.0f, 1.0f} - }; - - #pragma omp parallel for num_threads(opt.num_threads) - for (int p = 0; p < num_output; p++) { - for (int q = 0; q < num_input; q++) - { - const float* kernel0 = (const float*)weight_data + p * num_input * 9 + q * 9; - float* kernel_tm0 = weight_data_tm.channel(p).row(q); + Mat weight_data_tm; + weight_data_tm.create(4 * 4, num_input, num_output); - // transform kernel - const float* k0 = kernel0; - const float* k1 = kernel0 + 3; - const float* k2 = kernel0 + 6; + // G + const float ktm[4][3] = { + {1.0f, 0.0f, 0.0f}, + {1.0f / 2, 1.0f / 2, 1.0f / 2}, + {1.0f / 2, -1.0f / 2, 1.0f / 2}, + {0.0f, 0.0f, 1.0f} + }; - // h - float tmp[4][3]; - for (int i = 0; i < 4; i++) + #pragma omp parallel for num_threads(opt.num_threads) + for (int p = 0; p < num_output; p++) + { + for (int q = 0; q < num_input; q++) { - tmp[i][0] = k0[0] * ktm[i][0] + k0[1] * ktm[i][1] + k0[2] * ktm[i][2]; - tmp[i][1] = k1[0] * ktm[i][0] + k1[1] * ktm[i][1] + k1[2] * ktm[i][2]; - tmp[i][2] = k2[0] * ktm[i][0] + k2[1] * ktm[i][1] + k2[2] * ktm[i][2]; - } + const float* kernel0 = (const float*)weight_data + p * num_input * 9 + q * 9; + float* kernel_tm0 = weight_data_tm.channel(p).row(q); - // U - for (int j = 0; j < 4; j++) - { - float* tmpp = &tmp[j][0]; + // transform kernel + const float* k0 = kernel0; + const float* k1 = kernel0 + 3; + const float* k2 = kernel0 + 6; + // h + float tmp[4][3]; for (int i = 0; i < 4; i++) { - kernel_tm0[j * 4 + i] = tmpp[0] * ktm[i][0] + tmpp[1] * ktm[i][1] + tmpp[2] * ktm[i][2]; + tmp[i][0] = k0[0] * ktm[i][0] + k0[1] * ktm[i][1] + k0[2] * ktm[i][2]; + tmp[i][1] = k1[0] * ktm[i][0] + k1[1] * ktm[i][1] + k1[2] * ktm[i][2]; + tmp[i][2] = k2[0] * ktm[i][0] + k2[1] * ktm[i][1] + k2[2] * ktm[i][2]; + } + + // U + for (int j = 0; j < 4; j++) + { + float* tmpp = &tmp[j][0]; + + for (int i = 0; i < 4; i++) + { + kernel_tm0[j * 4 + i] = tmpp[0] * ktm[i][0] + tmpp[1] * ktm[i][1] + tmpp[2] * ktm[i][2]; + } } } } - } - // src = 16-inch-outch - // dst = 8a-8b-inch/8a-outch/8b-16 - Mat weight_data_tm_packed; - { - weight_data_tm_packed.create(num_input / elempack, num_output / out_elempack, 16, (size_t)4 * elempack * out_elempack, elempack * out_elempack); - - for (int k = 0; k < 16; k++) + // src = 16-inch-outch + // dst = 8a-8b-inch/8a-outch/8b-16 + Mat weight_data_tm_packed; { - float* g00 = weight_data_tm_packed.channel(k); + weight_data_tm_packed.create(num_input / elempack, num_output / out_elempack, 16, (size_t)4 * elempack * out_elempack, elempack * out_elempack); - for (int q = 0; q + (out_elempack - 1) < num_output; q += out_elempack) + for (int k = 0; k < 16; k++) { - for (int p = 0; p + (elempack - 1) < num_input; p += elempack) + float* g00 = weight_data_tm_packed.channel(k); + + for (int q = 0; q + (out_elempack - 1) < num_output; q += out_elempack) { - for (int i = 0; i < out_elempack; i++) + for (int p = 0; p + (elempack - 1) < num_input; p += elempack) { - const Mat k0 = weight_data_tm.channel(q + i); - - for (int j = 0; j < elempack; j++) + for (int i = 0; i < out_elempack; i++) { - const float* k00 = k0.row(p + j); + const Mat k0 = weight_data_tm.channel(q + i); + + for (int j = 0; j < elempack; j++) + { + const float* k00 = k0.row(p + j); - g00[0] = k00[k]; + g00[0] = k00[k]; - g00++; + g00++; + } } } } } } - } - if (support_image_storage && opt.use_image_storage) - { - cmd.record_upload(weight_data_tm_packed, weight_data_gpu_tm_image, opt); - } - else - { - cmd.record_upload(weight_data_tm_packed, weight_data_gpu_tm, opt); + if (support_image_storage && opt.use_image_storage) + { + cmd.record_upload(weight_data_tm_packed, weight_data_gpu_tm_winograd23_image, opt); + } + else + { + cmd.record_upload(weight_data_tm_packed, weight_data_gpu_tm_winograd23, opt); + } } } @@ -1121,288 +1076,295 @@ int Convolution_vulkan::forward(const VkMat& bottom_blob, VkMat& top_blob, VkCom bool is_conv1x1s1d1 = kernel_w == 1 && kernel_h == 1 && stride_w == 1 && stride_h == 1 && dilation_w == 1 && dilation_h == 1; bool is_conv3x3s1d1 = kernel_w == 3 && kernel_h == 3 && stride_w == 1 && stride_h == 1 && dilation_w == 1 && dilation_h == 1; - if (opt.use_winograd_convolution && is_conv3x3s1d1 && channels * elempack >= 32 && num_output >= 32 && ((elempack == 4 && out_elempack == 4) || (elempack == 8 && out_elempack == 8))) - { - // winograd43 - int outw_bordered = (outw + 3) / 4 * 4; - int outh_bordered = (outh + 3) / 4 * 4; - - int w_bordered = outw_bordered + 2; - int h_bordered = outh_bordered + 2; - - int block_x = outw_bordered / 4; - int block_y = outh_bordered / 4; - - // pad to 4n+2 - { - Option opt_pad = opt; - opt_pad.blob_vkallocator = opt.workspace_vkallocator; - - VkMat padding_param_blob(6, (size_t)4u, 1, opt.staging_vkallocator); - int* padding_params = padding_param_blob.mapped(); - - padding_params[0] = 0; - padding_params[1] = h_bordered - bottom_blob_bordered.h; - padding_params[2] = 0; - padding_params[3] = w_bordered - bottom_blob_bordered.w; - padding_params[4] = 0; - padding_params[5] = 0; - - std::vector padding_inputs(2); - padding_inputs[0] = bottom_blob_bordered; - padding_inputs[1] = padding_param_blob; - - std::vector padding_outputs(1); - winograd_padding->forward(padding_inputs, padding_outputs, cmd, opt_pad); - bottom_blob_bordered = padding_outputs[0]; - } - - // transform input - VkMat bottom_tm_blob; - { - bottom_tm_blob.create(block_x * block_y, channels, 36, elemsize, elempack, opt.workspace_vkallocator); - if (bottom_tm_blob.empty()) - return -100; - - std::vector bindings(2); - bindings[0] = bottom_blob_bordered; - bindings[1] = bottom_tm_blob; - - std::vector constants(7); - constants[0].i = bottom_blob_bordered.w; - constants[1].i = bottom_blob_bordered.h; - constants[2].i = bottom_blob_bordered.c; - constants[3].i = bottom_blob_bordered.cstep; - constants[4].i = bottom_tm_blob.cstep; - constants[5].i = block_x; - constants[6].i = block_y; - - VkMat dispatcher; - dispatcher.w = block_x; - dispatcher.h = block_y; - dispatcher.c = bottom_tm_blob.h; - - cmd.record_pipeline(pipeline_convolution_3x3s1d1_winograd43_transform_input, bindings, constants, dispatcher); - } - - // gemm - VkMat top_tm_blob; - { - top_tm_blob.create(block_x * block_y, num_output / out_elempack, 36, out_elemsize, out_elempack, opt.workspace_vkallocator); - if (top_tm_blob.empty()) - return -100; - - std::vector bindings(3); - bindings[0] = bottom_tm_blob; - bindings[1] = top_tm_blob; - bindings[2] = weight_data_gpu_tm; - - std::vector constants(5); - constants[0].i = bottom_tm_blob.h; - constants[1].i = bottom_tm_blob.cstep; - constants[2].i = top_tm_blob.w; - constants[3].i = top_tm_blob.h; - constants[4].i = top_tm_blob.cstep; - - VkMat dispatcher; - dispatcher.w = (top_tm_blob.w + 3) / 4; - dispatcher.h = top_tm_blob.h; - dispatcher.c = 36; - - cmd.record_pipeline(pipeline_convolution_3x3s1d1_winograd43_gemm, bindings, constants, dispatcher); - } - - // transform output - VkMat top_blob_bordered; - { - top_blob_bordered.create(outw_bordered, outh_bordered, num_output / out_elempack, out_elemsize, out_elempack, opt.blob_vkallocator); - if (top_blob_bordered.empty()) - return -100; - - std::vector bindings(3); - bindings[0] = top_tm_blob; - bindings[1] = top_blob_bordered; - bindings[2] = bias_data_gpu; - - std::vector constants(7); - constants[0].i = top_tm_blob.h; - constants[1].i = top_tm_blob.cstep; - constants[2].i = block_x; - constants[3].i = block_y; - constants[4].i = top_blob_bordered.w; - constants[5].i = top_blob_bordered.h; - constants[6].i = top_blob_bordered.cstep; - - VkMat dispatcher; - dispatcher.w = block_x; - dispatcher.h = block_y; - dispatcher.c = top_blob_bordered.c; - - cmd.record_pipeline(pipeline_convolution_3x3s1d1_winograd43_transform_output, bindings, constants, dispatcher); - } - - // crop top_blob - { - VkMat crop_param_blob(6, (size_t)4u, 1, opt.staging_vkallocator); - int* crop_params = crop_param_blob.mapped(); - - crop_params[0] = 0; - crop_params[1] = 0; - crop_params[2] = 0; - crop_params[3] = outw; - crop_params[4] = outh; - crop_params[5] = num_output; - - std::vector crop_inputs(2); - crop_inputs[0] = top_blob_bordered; - crop_inputs[1] = crop_param_blob; - - std::vector crop_outputs(1); - winograd_crop->forward(crop_inputs, crop_outputs, cmd, opt); - top_blob = crop_outputs[0]; - } - - return 0; - } if (opt.use_winograd_convolution && is_conv3x3s1d1 && channels * elempack >= 16 && num_output >= 16 && ((elempack == 4 && out_elempack == 4) || (elempack == 8 && out_elempack == 8))) { - // winograd23 - int outw_bordered = (outw + 1) / 2 * 2; - int outh_bordered = (outh + 1) / 2 * 2; - - int w_bordered = outw_bordered + 2; - int h_bordered = outh_bordered + 2; + bool pre_winograd43 = true; + if (vkdev->info.type() == 0 && (w <= 24 && h <= 24 && (w != 22 && h != 22))) + pre_winograd43 = false; + if (vkdev->info.type() != 0 && (w <= 12 && h <= 12)) + pre_winograd43 = false; - int block_x = outw_bordered / 2; - int block_y = outh_bordered / 2; - - // pad to 2n+2 + if (pre_winograd43) { - Option opt_pad = opt; - opt_pad.blob_vkallocator = opt.workspace_vkallocator; - - VkMat padding_param_blob(6, (size_t)4u, 1, opt.staging_vkallocator); - int* padding_params = padding_param_blob.mapped(); + // winograd43 + int outw_bordered = (outw + 3) / 4 * 4; + int outh_bordered = (outh + 3) / 4 * 4; - padding_params[0] = 0; - padding_params[1] = h_bordered - bottom_blob_bordered.h; - padding_params[2] = 0; - padding_params[3] = w_bordered - bottom_blob_bordered.w; - padding_params[4] = 0; - padding_params[5] = 0; - - std::vector padding_inputs(2); - padding_inputs[0] = bottom_blob_bordered; - padding_inputs[1] = padding_param_blob; + int w_bordered = outw_bordered + 2; + int h_bordered = outh_bordered + 2; - std::vector padding_outputs(1); - winograd_padding->forward(padding_inputs, padding_outputs, cmd, opt_pad); - bottom_blob_bordered = padding_outputs[0]; - } + int block_x = outw_bordered / 4; + int block_y = outh_bordered / 4; - // transform input - VkMat bottom_tm_blob; - { - bottom_tm_blob.create(block_x * block_y, channels, 16, elemsize, elempack, opt.workspace_vkallocator); - if (bottom_tm_blob.empty()) - return -100; + // pad to 4n+2 + { + Option opt_pad = opt; + opt_pad.blob_vkallocator = opt.workspace_vkallocator; + + VkMat padding_param_blob(6, (size_t)4u, 1, opt.staging_vkallocator); + int* padding_params = padding_param_blob.mapped(); + + padding_params[0] = 0; + padding_params[1] = h_bordered - bottom_blob_bordered.h; + padding_params[2] = 0; + padding_params[3] = w_bordered - bottom_blob_bordered.w; + padding_params[4] = 0; + padding_params[5] = 0; + + std::vector padding_inputs(2); + padding_inputs[0] = bottom_blob_bordered; + padding_inputs[1] = padding_param_blob; + + std::vector padding_outputs(1); + winograd_padding->forward(padding_inputs, padding_outputs, cmd, opt_pad); + bottom_blob_bordered = padding_outputs[0]; + } - std::vector bindings(2); - bindings[0] = bottom_blob_bordered; - bindings[1] = bottom_tm_blob; + // transform input + VkMat bottom_tm_blob; + { + bottom_tm_blob.create(block_x * block_y, channels, 36, elemsize, elempack, opt.workspace_vkallocator); + if (bottom_tm_blob.empty()) + return -100; + + std::vector bindings(2); + bindings[0] = bottom_blob_bordered; + bindings[1] = bottom_tm_blob; + + std::vector constants(7); + constants[0].i = bottom_blob_bordered.w; + constants[1].i = bottom_blob_bordered.h; + constants[2].i = bottom_blob_bordered.c; + constants[3].i = bottom_blob_bordered.cstep; + constants[4].i = bottom_tm_blob.cstep; + constants[5].i = block_x; + constants[6].i = block_y; + + VkMat dispatcher; + dispatcher.w = block_x; + dispatcher.h = block_y; + dispatcher.c = bottom_tm_blob.h; + + cmd.record_pipeline(pipeline_convolution_3x3s1d1_winograd43_transform_input, bindings, constants, dispatcher); + } - std::vector constants(7); - constants[0].i = bottom_blob_bordered.w; - constants[1].i = bottom_blob_bordered.h; - constants[2].i = bottom_blob_bordered.c; - constants[3].i = bottom_blob_bordered.cstep; - constants[4].i = bottom_tm_blob.cstep; - constants[5].i = block_x; - constants[6].i = block_y; + // gemm + VkMat top_tm_blob; + { + top_tm_blob.create(block_x * block_y, num_output / out_elempack, 36, out_elemsize, out_elempack, opt.workspace_vkallocator); + if (top_tm_blob.empty()) + return -100; + + std::vector bindings(3); + bindings[0] = bottom_tm_blob; + bindings[1] = top_tm_blob; + bindings[2] = weight_data_gpu_tm_winograd43; + + std::vector constants(5); + constants[0].i = bottom_tm_blob.h; + constants[1].i = bottom_tm_blob.cstep; + constants[2].i = top_tm_blob.w; + constants[3].i = top_tm_blob.h; + constants[4].i = top_tm_blob.cstep; + + VkMat dispatcher; + dispatcher.w = (top_tm_blob.w + 3) / 4; + dispatcher.h = top_tm_blob.h; + dispatcher.c = 36; + + cmd.record_pipeline(pipeline_convolution_3x3s1d1_winograd43_gemm, bindings, constants, dispatcher); + } - VkMat dispatcher; - dispatcher.w = block_x; - dispatcher.h = block_y; - dispatcher.c = bottom_tm_blob.h; + // transform output + VkMat top_blob_bordered; + { + top_blob_bordered.create(outw_bordered, outh_bordered, num_output / out_elempack, out_elemsize, out_elempack, opt.blob_vkallocator); + if (top_blob_bordered.empty()) + return -100; + + std::vector bindings(3); + bindings[0] = top_tm_blob; + bindings[1] = top_blob_bordered; + bindings[2] = bias_data_gpu; + + std::vector constants(7); + constants[0].i = top_tm_blob.h; + constants[1].i = top_tm_blob.cstep; + constants[2].i = block_x; + constants[3].i = block_y; + constants[4].i = top_blob_bordered.w; + constants[5].i = top_blob_bordered.h; + constants[6].i = top_blob_bordered.cstep; + + VkMat dispatcher; + dispatcher.w = block_x; + dispatcher.h = block_y; + dispatcher.c = top_blob_bordered.c; + + cmd.record_pipeline(pipeline_convolution_3x3s1d1_winograd43_transform_output, bindings, constants, dispatcher); + } - cmd.record_pipeline(pipeline_convolution_3x3s1d1_winograd23_transform_input, bindings, constants, dispatcher); + // crop top_blob + { + VkMat crop_param_blob(6, (size_t)4u, 1, opt.staging_vkallocator); + int* crop_params = crop_param_blob.mapped(); + + crop_params[0] = 0; + crop_params[1] = 0; + crop_params[2] = 0; + crop_params[3] = outw; + crop_params[4] = outh; + crop_params[5] = num_output; + + std::vector crop_inputs(2); + crop_inputs[0] = top_blob_bordered; + crop_inputs[1] = crop_param_blob; + + std::vector crop_outputs(1); + winograd_crop->forward(crop_inputs, crop_outputs, cmd, opt); + top_blob = crop_outputs[0]; + } } - - // gemm - VkMat top_tm_blob; + else { - top_tm_blob.create(block_x * block_y, num_output / out_elempack, 16, out_elemsize, out_elempack, opt.workspace_vkallocator); - if (top_tm_blob.empty()) - return -100; + // winograd23 + int outw_bordered = (outw + 1) / 2 * 2; + int outh_bordered = (outh + 1) / 2 * 2; - std::vector bindings(3); - bindings[0] = bottom_tm_blob; - bindings[1] = top_tm_blob; - bindings[2] = weight_data_gpu_tm; - - std::vector constants(5); - constants[0].i = bottom_tm_blob.h; - constants[1].i = bottom_tm_blob.cstep; - constants[2].i = top_tm_blob.w; - constants[3].i = top_tm_blob.h; - constants[4].i = top_tm_blob.cstep; - - VkMat dispatcher; - dispatcher.w = (top_tm_blob.w + 3) / 4; - dispatcher.h = top_tm_blob.h; - dispatcher.c = 16; - - cmd.record_pipeline(pipeline_convolution_3x3s1d1_winograd23_gemm, bindings, constants, dispatcher); - } + int w_bordered = outw_bordered + 2; + int h_bordered = outh_bordered + 2; - // transform output - VkMat top_blob_bordered; - { - top_blob_bordered.create(outw_bordered, outh_bordered, num_output / out_elempack, out_elemsize, out_elempack, opt.blob_vkallocator); - if (top_blob_bordered.empty()) - return -100; + int block_x = outw_bordered / 2; + int block_y = outh_bordered / 2; - std::vector bindings(3); - bindings[0] = top_tm_blob; - bindings[1] = top_blob_bordered; - bindings[2] = bias_data_gpu; + // pad to 2n+2 + { + Option opt_pad = opt; + opt_pad.blob_vkallocator = opt.workspace_vkallocator; + + VkMat padding_param_blob(6, (size_t)4u, 1, opt.staging_vkallocator); + int* padding_params = padding_param_blob.mapped(); + + padding_params[0] = 0; + padding_params[1] = h_bordered - bottom_blob_bordered.h; + padding_params[2] = 0; + padding_params[3] = w_bordered - bottom_blob_bordered.w; + padding_params[4] = 0; + padding_params[5] = 0; + + std::vector padding_inputs(2); + padding_inputs[0] = bottom_blob_bordered; + padding_inputs[1] = padding_param_blob; + + std::vector padding_outputs(1); + winograd_padding->forward(padding_inputs, padding_outputs, cmd, opt_pad); + bottom_blob_bordered = padding_outputs[0]; + } - std::vector constants(7); - constants[0].i = top_tm_blob.h; - constants[1].i = top_tm_blob.cstep; - constants[2].i = block_x; - constants[3].i = block_y; - constants[4].i = top_blob_bordered.w; - constants[5].i = top_blob_bordered.h; - constants[6].i = top_blob_bordered.cstep; + // transform input + VkMat bottom_tm_blob; + { + bottom_tm_blob.create(block_x * block_y, channels, 16, elemsize, elempack, opt.workspace_vkallocator); + if (bottom_tm_blob.empty()) + return -100; + + std::vector bindings(2); + bindings[0] = bottom_blob_bordered; + bindings[1] = bottom_tm_blob; + + std::vector constants(7); + constants[0].i = bottom_blob_bordered.w; + constants[1].i = bottom_blob_bordered.h; + constants[2].i = bottom_blob_bordered.c; + constants[3].i = bottom_blob_bordered.cstep; + constants[4].i = bottom_tm_blob.cstep; + constants[5].i = block_x; + constants[6].i = block_y; + + VkMat dispatcher; + dispatcher.w = block_x; + dispatcher.h = block_y; + dispatcher.c = bottom_tm_blob.h; + + cmd.record_pipeline(pipeline_convolution_3x3s1d1_winograd23_transform_input, bindings, constants, dispatcher); + } - VkMat dispatcher; - dispatcher.w = block_x; - dispatcher.h = block_y; - dispatcher.c = top_blob_bordered.c; + // gemm + VkMat top_tm_blob; + { + top_tm_blob.create(block_x * block_y, num_output / out_elempack, 16, out_elemsize, out_elempack, opt.workspace_vkallocator); + if (top_tm_blob.empty()) + return -100; + + std::vector bindings(3); + bindings[0] = bottom_tm_blob; + bindings[1] = top_tm_blob; + bindings[2] = weight_data_gpu_tm_winograd23; + + std::vector constants(5); + constants[0].i = bottom_tm_blob.h; + constants[1].i = bottom_tm_blob.cstep; + constants[2].i = top_tm_blob.w; + constants[3].i = top_tm_blob.h; + constants[4].i = top_tm_blob.cstep; + + VkMat dispatcher; + dispatcher.w = (top_tm_blob.w + 3) / 4; + dispatcher.h = top_tm_blob.h; + dispatcher.c = 16; + + cmd.record_pipeline(pipeline_convolution_3x3s1d1_winograd23_gemm, bindings, constants, dispatcher); + } - cmd.record_pipeline(pipeline_convolution_3x3s1d1_winograd23_transform_output, bindings, constants, dispatcher); - } + // transform output + VkMat top_blob_bordered; + { + top_blob_bordered.create(outw_bordered, outh_bordered, num_output / out_elempack, out_elemsize, out_elempack, opt.blob_vkallocator); + if (top_blob_bordered.empty()) + return -100; + + std::vector bindings(3); + bindings[0] = top_tm_blob; + bindings[1] = top_blob_bordered; + bindings[2] = bias_data_gpu; + + std::vector constants(7); + constants[0].i = top_tm_blob.h; + constants[1].i = top_tm_blob.cstep; + constants[2].i = block_x; + constants[3].i = block_y; + constants[4].i = top_blob_bordered.w; + constants[5].i = top_blob_bordered.h; + constants[6].i = top_blob_bordered.cstep; + + VkMat dispatcher; + dispatcher.w = block_x; + dispatcher.h = block_y; + dispatcher.c = top_blob_bordered.c; + + cmd.record_pipeline(pipeline_convolution_3x3s1d1_winograd23_transform_output, bindings, constants, dispatcher); + } - // crop top_blob - { - VkMat crop_param_blob(6, (size_t)4u, 1, opt.staging_vkallocator); - int* crop_params = crop_param_blob.mapped(); - - crop_params[0] = 0; - crop_params[1] = 0; - crop_params[2] = 0; - crop_params[3] = outw; - crop_params[4] = outh; - crop_params[5] = num_output; - - std::vector crop_inputs(2); - crop_inputs[0] = top_blob_bordered; - crop_inputs[1] = crop_param_blob; - - std::vector crop_outputs(1); - winograd_crop->forward(crop_inputs, crop_outputs, cmd, opt); - top_blob = crop_outputs[0]; + // crop top_blob + { + VkMat crop_param_blob(6, (size_t)4u, 1, opt.staging_vkallocator); + int* crop_params = crop_param_blob.mapped(); + + crop_params[0] = 0; + crop_params[1] = 0; + crop_params[2] = 0; + crop_params[3] = outw; + crop_params[4] = outh; + crop_params[5] = num_output; + + std::vector crop_inputs(2); + crop_inputs[0] = top_blob_bordered; + crop_inputs[1] = crop_param_blob; + + std::vector crop_outputs(1); + winograd_crop->forward(crop_inputs, crop_outputs, cmd, opt); + top_blob = crop_outputs[0]; + } } return 0; @@ -1591,288 +1553,295 @@ int Convolution_vulkan::forward(const VkImageMat& bottom_blob, VkImageMat& top_b bool is_conv1x1s1d1 = kernel_w == 1 && kernel_h == 1 && stride_w == 1 && stride_h == 1 && dilation_w == 1 && dilation_h == 1; bool is_conv3x3s1d1 = kernel_w == 3 && kernel_h == 3 && stride_w == 1 && stride_h == 1 && dilation_w == 1 && dilation_h == 1; - if (opt.use_winograd_convolution && is_conv3x3s1d1 && channels * elempack >= 32 && num_output >= 32 && ((elempack == 4 && out_elempack == 4) || (elempack == 8 && out_elempack == 8))) - { - // winograd43 - int outw_bordered = (outw + 3) / 4 * 4; - int outh_bordered = (outh + 3) / 4 * 4; - - int w_bordered = outw_bordered + 2; - int h_bordered = outh_bordered + 2; - - int block_x = outw_bordered / 4; - int block_y = outh_bordered / 4; - - // pad to 4n+2 - { - Option opt_pad = opt; - opt_pad.blob_vkallocator = opt.workspace_vkallocator; - - VkImageMat padding_param_blob(6, (size_t)4u, 1, opt.staging_vkallocator); - int* padding_params = padding_param_blob.mapped(); - - padding_params[0] = 0; - padding_params[1] = h_bordered - bottom_blob_bordered.h; - padding_params[2] = 0; - padding_params[3] = w_bordered - bottom_blob_bordered.w; - padding_params[4] = 0; - padding_params[5] = 0; - - std::vector padding_inputs(2); - padding_inputs[0] = bottom_blob_bordered; - padding_inputs[1] = padding_param_blob; - - std::vector padding_outputs(1); - winograd_padding->forward(padding_inputs, padding_outputs, cmd, opt_pad); - bottom_blob_bordered = padding_outputs[0]; - } - - // transform input - VkImageMat bottom_tm_blob; - { - bottom_tm_blob.create(block_x * block_y, channels, 36, elemsize, elempack, opt.workspace_vkallocator); - if (bottom_tm_blob.empty()) - return -100; - - std::vector bindings(2); - bindings[0] = bottom_blob_bordered; - bindings[1] = bottom_tm_blob; - - std::vector constants(7); - constants[0].i = bottom_blob_bordered.w; - constants[1].i = bottom_blob_bordered.h; - constants[2].i = bottom_blob_bordered.c; - constants[3].i = 0; //bottom_blob_bordered.cstep; - constants[4].i = 0; //bottom_tm_blob.cstep; - constants[5].i = block_x; - constants[6].i = block_y; - - VkImageMat dispatcher; - dispatcher.w = block_x; - dispatcher.h = block_y; - dispatcher.c = bottom_tm_blob.h; - - cmd.record_pipeline(pipeline_convolution_3x3s1d1_winograd43_transform_input, bindings, constants, dispatcher); - } - - // gemm - VkImageMat top_tm_blob; - { - top_tm_blob.create(block_x * block_y, num_output / out_elempack, 36, out_elemsize, out_elempack, opt.workspace_vkallocator); - if (top_tm_blob.empty()) - return -100; - - std::vector bindings(3); - bindings[0] = bottom_tm_blob; - bindings[1] = top_tm_blob; - bindings[2] = weight_data_gpu_tm_image; - - std::vector constants(5); - constants[0].i = bottom_tm_blob.h; - constants[1].i = 0; //bottom_tm_blob.cstep; - constants[2].i = top_tm_blob.w; - constants[3].i = top_tm_blob.h; - constants[4].i = 0; //top_tm_blob.cstep; - - VkImageMat dispatcher; - dispatcher.w = (top_tm_blob.w + 3) / 4; - dispatcher.h = top_tm_blob.h; - dispatcher.c = 36; - - cmd.record_pipeline(pipeline_convolution_3x3s1d1_winograd43_gemm, bindings, constants, dispatcher); - } - - // transform output - VkImageMat top_blob_bordered; - { - top_blob_bordered.create(outw_bordered, outh_bordered, num_output / out_elempack, out_elemsize, out_elempack, opt.blob_vkallocator); - if (top_blob_bordered.empty()) - return -100; - - std::vector bindings(3); - bindings[0] = top_tm_blob; - bindings[1] = top_blob_bordered; - bindings[2] = bias_data_gpu_image; - - std::vector constants(7); - constants[0].i = top_tm_blob.h; - constants[1].i = 0; //top_tm_blob.cstep; - constants[2].i = block_x; - constants[3].i = block_y; - constants[4].i = top_blob_bordered.w; - constants[5].i = top_blob_bordered.h; - constants[6].i = 0; //top_blob_bordered.cstep; - - VkImageMat dispatcher; - dispatcher.w = block_x; - dispatcher.h = block_y; - dispatcher.c = top_blob_bordered.c; - - cmd.record_pipeline(pipeline_convolution_3x3s1d1_winograd43_transform_output, bindings, constants, dispatcher); - } - - // crop top_blob - { - VkImageMat crop_param_blob(6, (size_t)4u, 1, opt.staging_vkallocator); - int* crop_params = crop_param_blob.mapped(); - - crop_params[0] = 0; - crop_params[1] = 0; - crop_params[2] = 0; - crop_params[3] = outw; - crop_params[4] = outh; - crop_params[5] = num_output; - - std::vector crop_inputs(2); - crop_inputs[0] = top_blob_bordered; - crop_inputs[1] = crop_param_blob; - - std::vector crop_outputs(1); - winograd_crop->forward(crop_inputs, crop_outputs, cmd, opt); - top_blob = crop_outputs[0]; - } - - return 0; - } if (opt.use_winograd_convolution && is_conv3x3s1d1 && channels * elempack >= 16 && num_output >= 16 && ((elempack == 4 && out_elempack == 4) || (elempack == 8 && out_elempack == 8))) { - // winograd23 - int outw_bordered = (outw + 1) / 2 * 2; - int outh_bordered = (outh + 1) / 2 * 2; + bool pre_winograd43 = true; + if (vkdev->info.type() == 0 && (w <= 24 && h <= 24 && (w != 22 && h != 22))) + pre_winograd43 = false; + if (vkdev->info.type() != 0 && (w <= 12 && h <= 12)) + pre_winograd43 = false; - int w_bordered = outw_bordered + 2; - int h_bordered = outh_bordered + 2; - - int block_x = outw_bordered / 2; - int block_y = outh_bordered / 2; - - // pad to 2n+2 + if (pre_winograd43) { - Option opt_pad = opt; - opt_pad.blob_vkallocator = opt.workspace_vkallocator; + // winograd43 + int outw_bordered = (outw + 3) / 4 * 4; + int outh_bordered = (outh + 3) / 4 * 4; - VkImageMat padding_param_blob(6, (size_t)4u, 1, opt.staging_vkallocator); - int* padding_params = padding_param_blob.mapped(); - - padding_params[0] = 0; - padding_params[1] = h_bordered - bottom_blob_bordered.h; - padding_params[2] = 0; - padding_params[3] = w_bordered - bottom_blob_bordered.w; - padding_params[4] = 0; - padding_params[5] = 0; - - std::vector padding_inputs(2); - padding_inputs[0] = bottom_blob_bordered; - padding_inputs[1] = padding_param_blob; + int w_bordered = outw_bordered + 2; + int h_bordered = outh_bordered + 2; - std::vector padding_outputs(1); - winograd_padding->forward(padding_inputs, padding_outputs, cmd, opt_pad); - bottom_blob_bordered = padding_outputs[0]; - } + int block_x = outw_bordered / 4; + int block_y = outh_bordered / 4; - // transform input - VkImageMat bottom_tm_blob; - { - bottom_tm_blob.create(block_x * block_y, channels, 16, elemsize, elempack, opt.workspace_vkallocator); - if (bottom_tm_blob.empty()) - return -100; + // pad to 4n+2 + { + Option opt_pad = opt; + opt_pad.blob_vkallocator = opt.workspace_vkallocator; + + VkImageMat padding_param_blob(6, (size_t)4u, 1, opt.staging_vkallocator); + int* padding_params = padding_param_blob.mapped(); + + padding_params[0] = 0; + padding_params[1] = h_bordered - bottom_blob_bordered.h; + padding_params[2] = 0; + padding_params[3] = w_bordered - bottom_blob_bordered.w; + padding_params[4] = 0; + padding_params[5] = 0; + + std::vector padding_inputs(2); + padding_inputs[0] = bottom_blob_bordered; + padding_inputs[1] = padding_param_blob; + + std::vector padding_outputs(1); + winograd_padding->forward(padding_inputs, padding_outputs, cmd, opt_pad); + bottom_blob_bordered = padding_outputs[0]; + } - std::vector bindings(2); - bindings[0] = bottom_blob_bordered; - bindings[1] = bottom_tm_blob; + // transform input + VkImageMat bottom_tm_blob; + { + bottom_tm_blob.create(block_x * block_y, channels, 36, elemsize, elempack, opt.workspace_vkallocator); + if (bottom_tm_blob.empty()) + return -100; + + std::vector bindings(2); + bindings[0] = bottom_blob_bordered; + bindings[1] = bottom_tm_blob; + + std::vector constants(7); + constants[0].i = bottom_blob_bordered.w; + constants[1].i = bottom_blob_bordered.h; + constants[2].i = bottom_blob_bordered.c; + constants[3].i = 0; //bottom_blob_bordered.cstep; + constants[4].i = 0; //bottom_tm_blob.cstep; + constants[5].i = block_x; + constants[6].i = block_y; + + VkImageMat dispatcher; + dispatcher.w = block_x; + dispatcher.h = block_y; + dispatcher.c = bottom_tm_blob.h; + + cmd.record_pipeline(pipeline_convolution_3x3s1d1_winograd43_transform_input, bindings, constants, dispatcher); + } - std::vector constants(7); - constants[0].i = bottom_blob_bordered.w; - constants[1].i = bottom_blob_bordered.h; - constants[2].i = bottom_blob_bordered.c; - constants[3].i = 0; //bottom_blob_bordered.cstep; - constants[4].i = 0; //bottom_tm_blob.cstep; - constants[5].i = block_x; - constants[6].i = block_y; + // gemm + VkImageMat top_tm_blob; + { + top_tm_blob.create(block_x * block_y, num_output / out_elempack, 36, out_elemsize, out_elempack, opt.workspace_vkallocator); + if (top_tm_blob.empty()) + return -100; + + std::vector bindings(3); + bindings[0] = bottom_tm_blob; + bindings[1] = top_tm_blob; + bindings[2] = weight_data_gpu_tm_winograd43_image; + + std::vector constants(5); + constants[0].i = bottom_tm_blob.h; + constants[1].i = 0; //bottom_tm_blob.cstep; + constants[2].i = top_tm_blob.w; + constants[3].i = top_tm_blob.h; + constants[4].i = 0; //top_tm_blob.cstep; + + VkImageMat dispatcher; + dispatcher.w = (top_tm_blob.w + 3) / 4; + dispatcher.h = top_tm_blob.h; + dispatcher.c = 36; + + cmd.record_pipeline(pipeline_convolution_3x3s1d1_winograd43_gemm, bindings, constants, dispatcher); + } - VkImageMat dispatcher; - dispatcher.w = block_x; - dispatcher.h = block_y; - dispatcher.c = bottom_tm_blob.h; + // transform output + VkImageMat top_blob_bordered; + { + top_blob_bordered.create(outw_bordered, outh_bordered, num_output / out_elempack, out_elemsize, out_elempack, opt.blob_vkallocator); + if (top_blob_bordered.empty()) + return -100; + + std::vector bindings(3); + bindings[0] = top_tm_blob; + bindings[1] = top_blob_bordered; + bindings[2] = bias_data_gpu_image; + + std::vector constants(7); + constants[0].i = top_tm_blob.h; + constants[1].i = 0; //top_tm_blob.cstep; + constants[2].i = block_x; + constants[3].i = block_y; + constants[4].i = top_blob_bordered.w; + constants[5].i = top_blob_bordered.h; + constants[6].i = 0; //top_blob_bordered.cstep; + + VkImageMat dispatcher; + dispatcher.w = block_x; + dispatcher.h = block_y; + dispatcher.c = top_blob_bordered.c; + + cmd.record_pipeline(pipeline_convolution_3x3s1d1_winograd43_transform_output, bindings, constants, dispatcher); + } - cmd.record_pipeline(pipeline_convolution_3x3s1d1_winograd23_transform_input, bindings, constants, dispatcher); + // crop top_blob + { + VkImageMat crop_param_blob(6, (size_t)4u, 1, opt.staging_vkallocator); + int* crop_params = crop_param_blob.mapped(); + + crop_params[0] = 0; + crop_params[1] = 0; + crop_params[2] = 0; + crop_params[3] = outw; + crop_params[4] = outh; + crop_params[5] = num_output; + + std::vector crop_inputs(2); + crop_inputs[0] = top_blob_bordered; + crop_inputs[1] = crop_param_blob; + + std::vector crop_outputs(1); + winograd_crop->forward(crop_inputs, crop_outputs, cmd, opt); + top_blob = crop_outputs[0]; + } } - - // gemm - VkImageMat top_tm_blob; + else { - top_tm_blob.create(block_x * block_y, num_output / out_elempack, 16, out_elemsize, out_elempack, opt.workspace_vkallocator); - if (top_tm_blob.empty()) - return -100; + // winograd23 + int outw_bordered = (outw + 1) / 2 * 2; + int outh_bordered = (outh + 1) / 2 * 2; - std::vector bindings(3); - bindings[0] = bottom_tm_blob; - bindings[1] = top_tm_blob; - bindings[2] = weight_data_gpu_tm_image; - - std::vector constants(5); - constants[0].i = bottom_tm_blob.c; - constants[1].i = 0; //bottom_tm_blob.cstep; - constants[2].i = top_tm_blob.w; - constants[3].i = top_tm_blob.h; - constants[4].i = 0; //top_tm_blob.cstep; - - VkImageMat dispatcher; - dispatcher.w = (top_tm_blob.w + 3) / 4; - dispatcher.h = top_tm_blob.h; - dispatcher.c = 16; - - cmd.record_pipeline(pipeline_convolution_3x3s1d1_winograd23_gemm, bindings, constants, dispatcher); - } + int w_bordered = outw_bordered + 2; + int h_bordered = outh_bordered + 2; - // transform output - VkImageMat top_blob_bordered; - { - top_blob_bordered.create(outw_bordered, outh_bordered, num_output / out_elempack, out_elemsize, out_elempack, opt.blob_vkallocator); - if (top_blob_bordered.empty()) - return -100; + int block_x = outw_bordered / 2; + int block_y = outh_bordered / 2; - std::vector bindings(3); - bindings[0] = top_tm_blob; - bindings[1] = top_blob_bordered; - bindings[2] = bias_data_gpu_image; + // pad to 2n+2 + { + Option opt_pad = opt; + opt_pad.blob_vkallocator = opt.workspace_vkallocator; + + VkImageMat padding_param_blob(6, (size_t)4u, 1, opt.staging_vkallocator); + int* padding_params = padding_param_blob.mapped(); + + padding_params[0] = 0; + padding_params[1] = h_bordered - bottom_blob_bordered.h; + padding_params[2] = 0; + padding_params[3] = w_bordered - bottom_blob_bordered.w; + padding_params[4] = 0; + padding_params[5] = 0; + + std::vector padding_inputs(2); + padding_inputs[0] = bottom_blob_bordered; + padding_inputs[1] = padding_param_blob; + + std::vector padding_outputs(1); + winograd_padding->forward(padding_inputs, padding_outputs, cmd, opt_pad); + bottom_blob_bordered = padding_outputs[0]; + } - std::vector constants(7); - constants[0].i = top_tm_blob.h; - constants[1].i = 0; //top_tm_blob.cstep; - constants[2].i = block_x; - constants[3].i = block_y; - constants[4].i = top_blob_bordered.w; - constants[5].i = top_blob_bordered.h; - constants[6].i = 0; //top_blob_bordered.cstep; + // transform input + VkImageMat bottom_tm_blob; + { + bottom_tm_blob.create(block_x * block_y, channels, 16, elemsize, elempack, opt.workspace_vkallocator); + if (bottom_tm_blob.empty()) + return -100; + + std::vector bindings(2); + bindings[0] = bottom_blob_bordered; + bindings[1] = bottom_tm_blob; + + std::vector constants(7); + constants[0].i = bottom_blob_bordered.w; + constants[1].i = bottom_blob_bordered.h; + constants[2].i = bottom_blob_bordered.c; + constants[3].i = 0; //bottom_blob_bordered.cstep; + constants[4].i = 0; //bottom_tm_blob.cstep; + constants[5].i = block_x; + constants[6].i = block_y; + + VkImageMat dispatcher; + dispatcher.w = block_x; + dispatcher.h = block_y; + dispatcher.c = bottom_tm_blob.h; + + cmd.record_pipeline(pipeline_convolution_3x3s1d1_winograd23_transform_input, bindings, constants, dispatcher); + } - VkImageMat dispatcher; - dispatcher.w = block_x; - dispatcher.h = block_y; - dispatcher.c = top_blob_bordered.c; + // gemm + VkImageMat top_tm_blob; + { + top_tm_blob.create(block_x * block_y, num_output / out_elempack, 16, out_elemsize, out_elempack, opt.workspace_vkallocator); + if (top_tm_blob.empty()) + return -100; + + std::vector bindings(3); + bindings[0] = bottom_tm_blob; + bindings[1] = top_tm_blob; + bindings[2] = weight_data_gpu_tm_winograd23_image; + + std::vector constants(5); + constants[0].i = bottom_tm_blob.c; + constants[1].i = 0; //bottom_tm_blob.cstep; + constants[2].i = top_tm_blob.w; + constants[3].i = top_tm_blob.h; + constants[4].i = 0; //top_tm_blob.cstep; + + VkImageMat dispatcher; + dispatcher.w = (top_tm_blob.w + 3) / 4; + dispatcher.h = top_tm_blob.h; + dispatcher.c = 16; + + cmd.record_pipeline(pipeline_convolution_3x3s1d1_winograd23_gemm, bindings, constants, dispatcher); + } - cmd.record_pipeline(pipeline_convolution_3x3s1d1_winograd23_transform_output, bindings, constants, dispatcher); - } + // transform output + VkImageMat top_blob_bordered; + { + top_blob_bordered.create(outw_bordered, outh_bordered, num_output / out_elempack, out_elemsize, out_elempack, opt.blob_vkallocator); + if (top_blob_bordered.empty()) + return -100; + + std::vector bindings(3); + bindings[0] = top_tm_blob; + bindings[1] = top_blob_bordered; + bindings[2] = bias_data_gpu_image; + + std::vector constants(7); + constants[0].i = top_tm_blob.h; + constants[1].i = 0; //top_tm_blob.cstep; + constants[2].i = block_x; + constants[3].i = block_y; + constants[4].i = top_blob_bordered.w; + constants[5].i = top_blob_bordered.h; + constants[6].i = 0; //top_blob_bordered.cstep; + + VkImageMat dispatcher; + dispatcher.w = block_x; + dispatcher.h = block_y; + dispatcher.c = top_blob_bordered.c; + + cmd.record_pipeline(pipeline_convolution_3x3s1d1_winograd23_transform_output, bindings, constants, dispatcher); + } - // crop top_blob - { - VkImageMat crop_param_blob(6, (size_t)4u, 1, opt.staging_vkallocator); - int* crop_params = crop_param_blob.mapped(); - - crop_params[0] = 0; - crop_params[1] = 0; - crop_params[2] = 0; - crop_params[3] = outw; - crop_params[4] = outh; - crop_params[5] = num_output; - - std::vector crop_inputs(2); - crop_inputs[0] = top_blob_bordered; - crop_inputs[1] = crop_param_blob; - - std::vector crop_outputs(1); - winograd_crop->forward(crop_inputs, crop_outputs, cmd, opt); - top_blob = crop_outputs[0]; + // crop top_blob + { + VkImageMat crop_param_blob(6, (size_t)4u, 1, opt.staging_vkallocator); + int* crop_params = crop_param_blob.mapped(); + + crop_params[0] = 0; + crop_params[1] = 0; + crop_params[2] = 0; + crop_params[3] = outw; + crop_params[4] = outh; + crop_params[5] = num_output; + + std::vector crop_inputs(2); + crop_inputs[0] = top_blob_bordered; + crop_inputs[1] = crop_param_blob; + + std::vector crop_outputs(1); + winograd_crop->forward(crop_inputs, crop_outputs, cmd, opt); + top_blob = crop_outputs[0]; + } } return 0; diff --git a/src/layer/vulkan/convolution_vulkan.h b/src/layer/vulkan/convolution_vulkan.h index 106df5d97..6efe6a8af 100644 --- a/src/layer/vulkan/convolution_vulkan.h +++ b/src/layer/vulkan/convolution_vulkan.h @@ -50,12 +50,14 @@ public: // winograd23 and winograd43 ncnn::Layer* winograd_padding; ncnn::Layer* winograd_crop; - VkMat weight_data_gpu_tm; - VkImageMat weight_data_gpu_tm_image; + VkMat weight_data_gpu_tm_winograd23; + VkImageMat weight_data_gpu_tm_winograd23_image; Pipeline* pipeline_convolution_3x3s1d1_winograd23_transform_input; Pipeline* pipeline_convolution_3x3s1d1_winograd23_gemm; Pipeline* pipeline_convolution_3x3s1d1_winograd23_transform_output; + VkMat weight_data_gpu_tm_winograd43; + VkImageMat weight_data_gpu_tm_winograd43_image; Pipeline* pipeline_convolution_3x3s1d1_winograd43_transform_input; Pipeline* pipeline_convolution_3x3s1d1_winograd43_gemm; Pipeline* pipeline_convolution_3x3s1d1_winograd43_transform_output; diff --git a/tests/test_convolution.cpp b/tests/test_convolution.cpp index 6094e827a..29711b5fb 100644 --- a/tests/test_convolution.cpp +++ b/tests/test_convolution.cpp @@ -40,7 +40,12 @@ static int test_convolution(int w, int h, int c, int outch, int kernel, int dila if (bias) weights[1] = RandomMat(outch); - int ret = test_layer("Convolution", pd, weights, a); + float epsilon = 0.001; + // larget epsilon for winograd optimization + if (kernel == 3 && dilation == 1 && stride == 1 && c >= 16 && outch >= 16) + epsilon = 0.002; + + int ret = test_layer("Convolution", pd, weights, a, epsilon); if (ret != 0) { fprintf(stderr, "test_convolution failed w=%d h=%d c=%d outch=%d kernel=%d dilation=%d stride=%d pad=%d bias=%d act=%d actparams=[%f,%f]\n", w, h, c, outch, kernel, dilation, stride, pad, bias, activation_type, activation_params[0], activation_params[1]);