1/* brig-control-handler.cc -- brig control directive handling
2   Copyright (C) 2016-2020 Free Software Foundation, Inc.
3   Contributed by Pekka Jaaskelainen <pekka.jaaskelainen@parmance.com>
4   for General Processor Tech.
5
6This file is part of GCC.
7
8GCC is free software; you can redistribute it and/or modify it under
9the terms of the GNU General Public License as published by the Free
10Software Foundation; either version 3, or (at your option) any later
11version.
12
13GCC is distributed in the hope that it will be useful, but WITHOUT ANY
14WARRANTY; without even the implied warranty of MERCHANTABILITY or
15FITNESS FOR A PARTICULAR PURPOSE.  See the GNU General Public License
16for more details.
17
18You should have received a copy of the GNU General Public License
19along 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
26size_t
27brig_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.", inst->control);
106    }
107  return base->byteCount;
108}
109