gcc/brig/brigfrontend/brig-control-handler.cc

   1 /* brig-control-handler.cc -- brig control directive handling
   2    Copyright (C) 2016-2018 Free Software Foundation, Inc.
   3    Contributed by Pekka Jaaskelainen <pekka.jaaskelainen@parmance.com>
   4    for General Processor Tech.
   5
   6 This file is part of GCC.
   7
   8 GCC is free software; you can redistribute it and/or modify it under
   9 the terms of the GNU General Public License as published by the Free
  10 Software Foundation; either version 3, or (at your option) any later
  11 version.
  12
  13 GCC is distributed in the hope that it will be useful, but WITHOUT ANY
  14 WARRANTY; without even the implied warranty of MERCHANTABILITY or
  15 FITNESS FOR A PARTICULAR PURPOSE.  See the GNU General Public License
  16 for more details.
  17
  18 You should have received a copy of the GNU General Public License
  19 along with GCC; see the file COPYING3.  If not see
  20 <http://www.gnu.org/licenses/>.  */
  21
  22 #include "brig-code-entry-handler.h"
  23 #include "diagnostic.h"
  24 #include "print-tree.h"
  25
  26 size_t
  27 brig_directive_control_handler::operator () (const BrigBase *base)
  28 {
  29   const BrigDirectiveControl *inst = (const BrigDirectiveControl *) base;
  30   const BrigData *operand_entries
  31     = m_parent.get_brig_data_entry (inst->operands);
  32
  33   /* Parse the constant integer operands.  */
  34   std::vector<tree> operands;
  35   for (size_t i = 0; i < operand_entries->byteCount / 4; ++i)
  36     {
  37       uint32_t operand_offset
  38         = ((const uint32_t *) &operand_entries->bytes)[i];
  39       const BrigBase *operand_data
  40         = m_parent.get_brig_operand_entry (operand_offset);
  41
  42       tree operand_type
  43         = (inst->control == BRIG_CONTROL_REQUIREDGRIDSIZE
  44            || inst->control == BRIG_CONTROL_MAXFLATGRIDSIZE) ?
  45         uint64_type_node : uint32_type_node;
  46       operands.push_back
  47         (build_tree_operand (*(const BrigInstBase*)inst, *operand_data,
  48                              operand_type));
  49     }
  50
  51   switch (inst->control)
  52     {
  53     case BRIG_CONTROL_MAXDYNAMICGROUPSIZE:
  54       {
  55         m_parent.m_cf->m_descriptor.max_dynamic_group_size
  56           = brig_function::int_constant_value (operands.at (0));
  57         break;
  58       }
  59     case BRIG_CONTROL_MAXFLATGRIDSIZE:
  60       {
  61         m_parent.m_cf->m_descriptor.max_flat_grid_size
  62           = brig_function::int_constant_value (operands.at (0));
  63         break;
  64       }
  65     case BRIG_CONTROL_MAXFLATWORKGROUPSIZE:
  66       {
  67         m_parent.m_cf->m_descriptor.max_flat_workgroup_size
  68           = brig_function::int_constant_value (operands.at (0));
  69         break;
  70       }
  71     case BRIG_CONTROL_REQUIREDDIM:
  72       {
  73         m_parent.m_cf->m_descriptor.required_dim
  74           = brig_function::int_constant_value (operands.at (0));
  75         break;
  76       }
  77     case BRIG_CONTROL_REQUIREDGRIDSIZE:
  78       {
  79         m_parent.m_cf->m_descriptor.required_grid_size[0]
  80           = brig_function::int_constant_value (operands.at (0));
  81         m_parent.m_cf->m_descriptor.required_grid_size[1]
  82           = brig_function::int_constant_value (operands.at (1));
  83         m_parent.m_cf->m_descriptor.required_grid_size[2]
  84           = brig_function::int_constant_value (operands.at (2));
  85         break;
  86       }
  87     case BRIG_CONTROL_REQUIREDWORKGROUPSIZE:
  88       {
  89         m_parent.m_cf->m_descriptor.required_workgroup_size[0]
  90           = brig_function::int_constant_value (operands.at (0));
  91         m_parent.m_cf->m_descriptor.required_workgroup_size[1]
  92           = brig_function::int_constant_value (operands.at (1));
  93         m_parent.m_cf->m_descriptor.required_workgroup_size[2]
  94           = brig_function::int_constant_value (operands.at (2));
  95         break;
  96       }
  97     case BRIG_CONTROL_REQUIRENOPARTIALWORKGROUPS:
  98       /* Performance hint only, ignored for now.  */
  99       break;
 100     case BRIG_CONTROL_ENABLEBREAKEXCEPTIONS:
 101     case BRIG_CONTROL_ENABLEDETECTEXCEPTIONS:
 102       /* Unimplemented.  */
 103       break;
 104     default:
 105       sorry ("Unsupported control directive %x.\n", inst->control);
 106     }
 107   return base->byteCount;
 108 }