tensorflow/lite/delegates/gpu/cl/kernels/fully_connected.cc
// We split vec vec dot (every thread do vec vec dot product in basic
// vec mat mult) on 4 parts to create more threads
// tid.y thread process every 4-th element in vec vec dot
// Good results for ~1024 x 1024 sizes, for other can be written more
// optimized shaders
We split vec vec dot (every thread do vec vec dot product in basic
vec mat mult) on 4 parts to create more threads
tid.y thread process every 4-th element in vec vec dot
Good results for ~1024 x 1024 sizes, for other can be written more
optimized shaders
We split vec vec dot ( every thread do vec vec dot product in basic vec mat mult ) on 4 parts to create more threads tid . y thread process every 4 - th element in vec vec dot Good results for ~1024 x 1024 sizes for other can be written more optimized shaders
std::string FullyConnected::GetFullyConnectedKernelCode(
    const OperationDef& op_def, const int3& work_group_size) {
  AddSrcTensor("src_tensor", op_def.src_tensors[0]);
  AddDstTensor("dst_tensor", op_def.dst_tensors[0]);
  std::string c = GetCommonDefines(op_def.precision);
  switch (op_def.precision) {
    case CalculationsPrecision::F32:
      c += "#define FLT16 float16\n";
      break;
    case CalculationsPrecision::F32_F16:
    case CalculationsPrecision::F16:
      c += "#define FLT16 half16\n";
      break;
  }
  const std::string wg_x = std::to_string(work_group_size.x);
  const std::string wg_y = std::to_string(work_group_size.y);
  c += "__kernel void main_function(\n";
  c += "$0) {\n";
  c += "  int gid = get_global_id(0);\n";
  c += "  if (gid >= args.dst_tensor.Slices()) {\n";
  c += "    return;\n";
  c += "  }\n";
  c += "  int2 tid = (int2)(get_local_id(0), get_local_id(1));\n";
  c += "  ACCUM_FLT4 s = (ACCUM_FLT4)(0.0f);\n";
  c += "  for (uint c = tid.y; c < args.src_tensor.Slices(); c += " + wg_y +
       ") {\n";
  c += "    FLT4 v = args.src_tensor.Read(0, 0, c);\n";
  c += "    FLT16 w = args.weights.Read(c * args.dst_tensor.Slices() + gid);\n";
  c += "    s.x += dot(v, w.s0123);\n";
  c += "    s.y += dot(v, w.s4567);\n";
  c += "    s.z += dot(v, w.s89ab);\n";
  c += "    s.w += dot(v, w.scdef);\n";
  c += "  }\n";
  c += "  __local ACCUM_FLT4 temp[" + wg_x + "][" + wg_y + "];\n";
  c += "  temp[tid.x][tid.y] = s;\n";
  c += "  barrier(CLK_LOCAL_MEM_FENCE);\n";
  c += "  if (tid.y == 0) {\n";
  for (int i = 1; i < work_group_size.y; ++i) {
    c += "    s += temp[tid.x][" + std::to_string(i) + "];\n";
  }
  c += "    FLT4 r0 = TO_FLT4(s) + args.biases.Read(gid);\n";
  c += "    args.dst_tensor.Write(r0, 0, 0, gid);\n";
  c += "  }\n";
  c += "}\n";
  return c;
}
std :: string FullyConnected :: GetFullyConnectedKernelCode ( const OperationDef & op_def , const int3 & work_group_size ) { AddSrcTensor ( " " , op_def . src_tensors [ 0 ] ) ; AddDstTensor ( " " , op_def . dst_tensors [ 0 ] ) ; std :: string c = GetCommonDefines ( op_def . precision ) ; switch ( op_def . precision ) { case CalculationsPrecision :: F32 : c += " \n " ; break ; case CalculationsPrecision :: F32_F16 : case CalculationsPrecision :: F16 : c += " \n " ; break ; } const std :: string wg_x = std :: to_string ( work_group_size . x ) ; const std :: string wg_y = std :: to_string ( work_group_size . y ) ; c += " \n " ; c += " \n " ; c += " \n " ; c += " \n " ; c += " \n " ; c += " \n " ; c += " \n " ; c += " \n " ; c += " " + wg_y + " \n " ; c += " \n " ; c += " \n " ; c += " \n " ; c += " \n " ; c += " \n " ; c += " \n " ; c += " \n " ; c += " " + wg_x + " " + wg_y + " \n " ; c += " \n " ; c += " \n " ; c += " \n " ; for ( int i = 1 ; i < work_group_size . y ; ++ i ) { c += " " + std :: to_string ( i ) + " \n " ; } c += " \n " ; c += " \n " ; c += " \n " ; c += " \n " ; return c ; }
We split vec vec dot (every thread do vec vec dot product in basic
vec mat mult) on 4 parts to create more threads
tid.y thread process every 4-th element in vec vec dot
Good results for ~1024 x 1024 sizes, for other can be written more
optimized shaders
We split vec vec dot ( every thread do vec vec dot product in basic vec mat mult ) on 4 parts to create more threads tid . y thread process every 4 - th element in vec vec dot Good results for ~1024 x 1024 sizes for other can be written more optimized shaders