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 operator ()(const BrigBase * base)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 = 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 = 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 = 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 = 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 = int_constant_value (operands.at (0)); 81 m_parent.m_cf->m_descriptor.required_grid_size[1] 82 = int_constant_value (operands.at (1)); 83 m_parent.m_cf->m_descriptor.required_grid_size[2] 84 = 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 = int_constant_value (operands.at (0)); 91 m_parent.m_cf->m_descriptor.required_workgroup_size[1] 92 = int_constant_value (operands.at (1)); 93 m_parent.m_cf->m_descriptor.required_workgroup_size[2] 94 = 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 } 109