inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reorder_weights.cl

   1 // Copyright (c) 2016-2017 Intel Corporation
   2 //
   3 // Licensed under the Apache License, Version 2.0 (the "License");
   4 // you may not use this file except in compliance with the License.
   5 // You may obtain a copy of the License at
   6 //
   7 //      http://www.apache.org/licenses/LICENSE-2.0
   8 //
   9 // Unless required by applicable law or agreed to in writing, software
  10 // distributed under the License is distributed on an "AS IS" BASIS,
  11 // WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
  12 // See the License for the specific language governing permissions and
  13 // limitations under the License.
  14
  15
  16 #include "include/fetch.cl"
  17 #include "include/reshape_dims.cl"
  18 #include "include/data_types.cl"
  19
  20
  21 ///////////////////////// Input Index /////////////////////////
  22 inline uint FUNC(get_input_index)(uint o, uint i, uint y, uint x)
  23 {
  24 #if   INPUT0_SIMPLE
  25     return GET_FILTER_INDEX(INPUT0, o, i, y, x);
  26 #elif defined INPUT0_LAYOUT_OS_IYX_OSV16    || \
  27       defined INPUT0_LAYOUT_OS_I_OSV16      || \
  28       defined INPUT0_LAYOUT_OS_I_OSV8__AI8  || \
  29       defined INPUT0_LAYOUT_OS_I_OSV16__AI8
  30     return GET_FILTER_OS_IYX_OSV8_INDEX(INPUT0, o, i, y, x, SUB_GROUP_SIZE);
  31 #elif defined INPUT0_LAYOUT_IYX_OSV32
  32     return GET_FILTER_OS_IYX_OSV8_INDEX(INPUT0, o, i, y, x, 32);
  33 #elif defined INPUT0_LAYOUT_IYX_OSV64
  34     return GET_FILTER_OS_IYX_OSV8_INDEX(INPUT0, o, i, y, x, 64);
  35 #elif defined INPUT0_LAYOUT_OS_IYX_OSV16_ROTATE_180
  36     return GET_FILTER_OS_IYX_OSV8_ROTATE_180_INDEX(INPUT0, o, i, y, x, SUB_GROUP_SIZE);
  37 #elif defined INPUT0_LAYOUT_I_YXS_OS_YXSV2_OSV16
  38     return GET_FILTER_I_YXS_OS_YXSV2_OSV_INDEX(INPUT0, o, i, y, x, SUB_GROUP_SIZE);
  39 #elif defined INPUT0_LAYOUT_IY_XS_OS_XSV2_OSV16__AO32 || defined OUTPUT_LAYOUT_IY_XS_OS_XSV2_OSV8__AO32
  40     return GET_FILTER_IY_XS_OS_XSV2_OSV_INDEX(INPUT0, o, i, y, x, SUB_GROUP_SIZE);
  41 #elif defined INPUT0_LAYOUT_IMAGE_2D_WEIGHTS_C1_B_FYX
  42     #error - not supported yet
  43 #elif defined INPUT0_LAYOUT_OS_IS_YX_ISA8_OSV8_ISV4
  44         return GET_FILTER_OS_IS_YX_ISA8_OSV8_ISV4(INPUT0, o, i, y, x);
  45 #elif defined INPUT0_LAYOUT_IS_O_YX_ISV32
  46     return GET_FILTER_IS_O_YX_ISV32(INPUT0, o, i, y, x);
  47 #elif defined INPUT0_LAYOUT_IS_O32_YX_ISV32_SWIZZLED_BY_4
  48     return GET_FILTER_IS_O32_YX_ISV32_SWIZZLED_BY_4(INPUT0, o, i, y, x);
  49 #elif defined INPUT0_LAYOUT_OS_IS_Y_X8_OSV8_ISV4
  50     return GET_FILTER_OS_IS_Y_X8_OSV8_ISV4(INPUT0, o, i, y, x);
  51 #else
  52 #error reorder_weights.cl: input format - not supported
  53 #endif
  54 }
  55
  56 ///////////////////////// Output Index /////////////////////////
  57
  58 inline uint FUNC(get_output_index)(uint o, uint i, uint y, uint x)
  59 {
  60 #if   OUTPUT_SIMPLE
  61     return GET_FILTER_INDEX(OUTPUT, o, i, y, x);
  62 #elif defined OUTPUT_LAYOUT_OS_IYX_OSV16    || \
  63       defined OUTPUT_LAYOUT_OS_I_OSV16      || \
  64       defined OUTPUT_LAYOUT_OS_I_OSV8__AI8  || \
  65       defined OUTPUT_LAYOUT_OS_I_OSV16__AI8
  66     return GET_FILTER_OS_IYX_OSV8_INDEX(OUTPUT, o, i, y, x, SUB_GROUP_SIZE);
  67 #elif defined OUTPUT_LAYOUT_OS_IYX_OSV32
  68     return GET_FILTER_OS_IYX_OSV8_INDEX(OUTPUT, o, i, y, x, 32);
  69 #elif defined OUTPUT_LAYOUT_OS_IYX_OSV64
  70     return GET_FILTER_OS_IYX_OSV8_INDEX(OUTPUT, o, i, y, x, 64);
  71 #elif defined OUTPUT_LAYOUT_OS_IYX_OSV16_ROTATE_180
  72     return GET_FILTER_OS_IYX_OSV8_ROTATE_180_INDEX(OUTPUT, o, i, y, x, SUB_GROUP_SIZE);
  73 #elif defined OUTPUT_LAYOUT_I_YXS_OS_YXSV2_OSV16
  74     return GET_FILTER_I_YXS_OS_YXSV2_OSV_INDEX(OUTPUT, o, i, y, x, SUB_GROUP_SIZE);
  75 #elif defined OUTPUT_LAYOUT_IY_XS_OS_XSV2_OSV16__AO32 || defined OUTPUT_LAYOUT_IY_XS_OS_XSV2_OSV8__AO32
  76     return GET_FILTER_IY_XS_OS_XSV2_OSV_INDEX(OUTPUT, o, i, y, x, SUB_GROUP_SIZE);
  77 #elif defined OUTPUT_LAYOUT_IMAGE_2D_WEIGHTS_C1_B_FYX
  78     return 0; //will not be used for images
  79 #elif defined OUTPUT_LAYOUT_OS_IS_YX_ISA8_OSV8_ISV4
  80         return GET_FILTER_OS_IS_YX_ISA8_OSV8_ISV4(OUTPUT, o, i, y, x);
  81 #elif defined OUTPUT_LAYOUT_IS_O_YX_ISV32
  82     return GET_FILTER_IS_O_YX_ISV32(OUTPUT, o, i, y, x);
  83 #elif defined OUTPUT_LAYOUT_IS_O32_YX_ISV32_SWIZZLED_BY_4
  84     return GET_FILTER_IS_O32_YX_ISV32_SWIZZLED_BY_4(OUTPUT, o, i, y, x);
  85 #elif defined OUTPUT_LAYOUT_OS_IS_Y_X8_OSV8_ISV4
  86     return GET_FILTER_OS_IS_Y_X8_OSV8_ISV4(OUTPUT, o, i, y, x);
  87 #elif defined OUTPUT_LAYOUT_OS_IS_YX_OSV16_ISV4
  88     return GET_FILTER_OS_IS_YX_OSV16_ISV4_INDEX(OUTPUT, o, i, y, x);
  89 #elif defined OUTPUT_LAYOUT_OS_IS_YX_ISA8_OSV8_ISV4_SWIZZLED_BY_4
  90     return GET_FILTER_OS_IS_YX_ISA8_OSV8_ISV4_SWIZZLED_BY_4_INDEX(OUTPUT, o, i, y, x);
  91 #else
  92 #error reorder_weights.cl: output format - not supported
  93 #endif
  94 }
  95
  96 #if OUTPUT_LAYOUT_IMAGE_2D_WEIGHTS_C1_B_FYX
  97 KERNEL (reorder_weights)(const __global INPUT0_TYPE* input, write_only image2d_t output)
  98 {
  99     const unsigned o = get_global_id(0);
 100     const unsigned iyx = get_global_id(1);
 101     const unsigned x = iyx % INPUT0_SIZE_X;
 102     const unsigned y = (iyx / INPUT0_SIZE_X) % INPUT0_SIZE_Y;
 103     const unsigned i = (iyx / INPUT0_SIZE_X) / INPUT0_SIZE_Y;
 104
 105     MAKE_VECTOR_TYPE(UNIT_TYPE, 4) input_val = (MAKE_VECTOR_TYPE(UNIT_TYPE, 4))(UNIT_VAL_ZERO, UNIT_VAL_ZERO, UNIT_VAL_ZERO, UNIT_VAL_ZERO);
 106     const int2 coord = (int2)(o, iyx);
 107     uint4 ir = FUNC_CALL(reshape_dims)(o,i,y,x, OUTPUT_SIZE_Y, OUTPUT_SIZE_X, INPUT0_SIZE_Y, INPUT0_SIZE_X, OUTPUT_DIMS, INPUT0_DIMS);
 108     input_val.s0 = TO_OUTPUT_TYPE(input[FUNC_CALL(get_input_index)(ir[0],ir[1],ir[2],ir[3])]);
 109     IMAGE_WRITE(output, coord, input_val);
 110 }
 111 #else
 112 KERNEL (reorder_weights)(const __global INPUT0_TYPE* input, __global OUTPUT_TYPE* output)
 113 {
 114     const unsigned o = get_global_id(0);
 115     const unsigned i = get_global_id(1);
 116 #if   OUTPUT_DIMS == 2
 117     const unsigned y = 0;
 118     const unsigned x = 0;
 119 #elif OUTPUT_DIMS == 4
 120     const unsigned y = get_global_id(2) / INPUT0_SIZE_X;
 121     const unsigned x = get_global_id(2) % INPUT0_SIZE_X;
 122 #endif
 123     uint4 ir = FUNC_CALL(reshape_dims)(o,i,y,x, OUTPUT_SIZE_Y, OUTPUT_SIZE_X, INPUT0_SIZE_Y, INPUT0_SIZE_X, OUTPUT_DIMS, INPUT0_DIMS);
 124     output[FUNC_CALL(get_output_index)(o, i, y, x)] = TO_OUTPUT_TYPE(input[FUNC_CALL(get_input_index)(ir[0],ir[1],ir[2],ir[3])]);
 125 }
 126 #endif