code clean up

Signed-off-by: Min, Byungil <[email protected]>

Author: byungilm

src/plugins/intel_gpu/src/kernel_selector/cl_kernels/col_to_im_opt.cl

@@ -18,25 +18,20 @@ KERNEL(col_to_im_opt)(const __global INPUT0_TYPE* input,
     const uint pads_begin[2] = {PAD_BEGIN_SIZE_X, PAD_BEGIN_SIZE_Y};
     const uint pads_end[2] = {PAD_END_SIZE_X, PAD_END_SIZE_Y};
 
-    const uint num_blocks = INPUT0_SIZE_Y;
-    const uint kernel_product = KERNEL_SIZE_X * KERNEL_SIZE_Y;
-    const uint channels_per_column = INPUT0_FEATURE_NUM; 
-    const uint channel_count = channels_per_column / kernel_product;
-
     const uint batch_count = INPUT0_BATCH_NUM;
     const uint batch = get_global_id(2);
 
-    // printf("batch(%d) num_blocks(%u) output(%u, %u), channel(%u, %u) original_height(%u) original_width(%u) \n",
+    // printf("batch(%d) num_blocks(%u) output(%u, %u), kernel(%u, %u) original_height(%u) original_width(%u) \n",
     //         batch, num_blocks, (uint)OUT_SIZE_X, (uint)OUT_SIZE_Y, (uint)KERNEL_SIZE_X, (uint)KERNEL_SIZE_Y, ORIG_HEIGHT, ORIG_WIDTH);
 
     // for (uint batch = 0; batch < batch_count; ++batch) {
-        for (uint column = 0; column < channels_per_column; ++column) {
+        for (uint column = 0; column < NUM_ELEMENTS_FOR_BLOCK; ++column) {
             const uint width_offset = column % kernel_size[1];
             const uint height_offset = (column / kernel_size[1]) % kernel_size[0];
-            const uint channel_idx = column / kernel_product;
+            const uint channel_idx = column / KERNEL_PRODUCT;
 
-            const uint out_idx = (batch * channel_count + channel_idx) * output_size[0];
-            const uint height_idx = (batch * channels_per_column + column) * ORIG_HEIGHT;
+            const uint out_idx = (batch * NUM_CHANNELS + channel_idx) * output_size[0];
+            const uint height_idx = (batch * NUM_ELEMENTS_FOR_BLOCK + column) * ORIG_HEIGHT;
 
             for (uint column_height_idx = 0; column_height_idx < ORIG_HEIGHT; ++column_height_idx) {
                 // get_image_dimension_index(column_height_idx, height_offset, 0);

src/plugins/intel_gpu/src/kernel_selector/kernels/col_to_im/col_to_im_kernel_base.cpp

@@ -28,18 +28,28 @@ bool ColToImKernelBase::Validate(const Params& p) const {
 
 JitConstants ColToImKernelBase::GetJitConstants(const col_to_im_params& params) const {
     JitConstants jit = MakeBaseParamsJitConstants(params);
+    auto input = params.inputs[0];
     const auto& output_size = params.output_size;
     const auto& stride = params.stride;
     const auto& dilation = params.dilation;
     const auto& pads_begin = params.padding_begin;
     const auto& pads_end = params.padding_begin;
 
-    // Get original dimension
-    const uint orig_height = (output_size.x + pads_begin.x + pads_end.x - (dilation.x * (params.kernel_size.x - 1) + 1)) / stride.x + 1;
-    const uint orig_width = (output_size.y + pads_begin.y + pads_end.y - (dilation.y * (params.kernel_size.y - 1) + 1)) / stride.y + 1;
+    const auto orig_height = (output_size.x + pads_begin.x + pads_end.x - (dilation.x * (params.kernel_size.x - 1) + 1)) / stride.x + 1;
+    const auto orig_width = (output_size.y + pads_begin.y + pads_end.y - (dilation.y * (params.kernel_size.y - 1) + 1)) / stride.y + 1;
     jit.AddConstant(MakeJitConstant("ORIG_HEIGHT", orig_height));
     jit.AddConstant(MakeJitConstant("ORIG_WIDTH", orig_width));
 
+    // Consider input tensor : (N, C * Product(kernel_size), L)
+    const auto num_elements_for_block = input.Feature().v;
+    const auto num_blocks = input.Y().v;
+    const auto kernel_product = params.kernel_size.x * params.kernel_size.y;
+    const auto num_channels = num_elements_for_block / kernel_product;
+    jit.AddConstant(MakeJitConstant("NUM_ELEMENTS_FOR_BLOCK", num_elements_for_block));
+    jit.AddConstant(MakeJitConstant("KERNEL_PRODUCT", kernel_product));
+    jit.AddConstant(MakeJitConstant("NUM_CHANNELS", num_channels));
+    jit.AddConstant(MakeJitConstant("NUM_BLOCKS", num_blocks));
+
     jit.AddConstant(MakeJitConstant("OUT", params.output_size));
     jit.AddConstant(MakeJitConstant("KERNEL", params.kernel_size));
     jit.AddConstant(MakeJitConstant("STRIDE", stride));

src/plugins/intel_gpu/src/kernel_selector/kernels/col_to_im/col_to_im_kernel_opt.cpp

@@ -56,8 +56,8 @@ CommonDispatchData ColToImKernelOpt::SetDefault(const col_to_im_params& params)
     // dispatchData.gws = {1, 1, 1};
     dispatchData.lws = {1, 1, 1};
 
-    std::cout << " Select ColToImKernelOpt : gws(" << dispatchData.gws[0] << ", " << dispatchData.gws[1] << ", " << dispatchData.gws[2] << ")" << std::endl;
-    std::cout << " Select ColToImKernelOpt : lws(" << dispatchData.lws[0] << ", " << dispatchData.lws[1] << ", " << dispatchData.lws[2] << ")" << std::endl;
+    // std::cout << " Select ColToImKernelOpt : gws(" << dispatchData.gws[0] << ", " << dispatchData.gws[1] << ", " << dispatchData.gws[2] << ")" << std::endl;
+    // std::cout << " Select ColToImKernelOpt : lws(" << dispatchData.lws[0] << ", " << dispatchData.lws[1] << ", " << dispatchData.lws[2] << ")" << std::endl;
 
     return dispatchData;
 }

src/plugins/intel_gpu/src/plugin/ops/col_to_im.cpp

@@ -33,6 +33,8 @@ static void CreateCol2ImOp(ProgramBuilder& p, const std::shared_ptr<ov::op::v15:
     for (auto p: op->get_pads_end())
         padding_end.push_back(p);
 
+    // std::cout << ">> col2im : " << op->get_friendly_name() << std::endl;
+
     auto output_shape_const = ov::as_type_ptr<ov::op::v0::Constant>(op->get_input_node_shared_ptr(1));
     auto vec_output_shape = output_shape_const->cast_vector<size_t>();
     ov::Shape output_shape(vec_output_shape);
@@ -41,7 +43,11 @@ static void CreateCol2ImOp(ProgramBuilder& p, const std::shared_ptr<ov::op::v15:
     auto kernel_size = kernel_size_const->cast_vector<size_t>();
     ov::Shape kernel_shape(kernel_size);
 
+    // std::cout << "  -- output shape : " << vec_output_shape[0] << ", " << vec_output_shape[1] << std::endl;
+    // std::cout << "  -- kernel size : " << kernel_shape.to_string() << std::endl;
+
     // Create col2im prim
+    // iputs : data, output size,  kernel_size(required)
     auto CallToImPrim = cldnn::col_to_im(layerName,
                                             inputPrimitives[0],
                                             inputPrimitives[1],