1 /* A pass for lowering gimple to HSAIL 2 Copyright (C) 2013-2018 Free Software Foundation, Inc. 3 Contributed by Martin Jambor <mjambor@suse.cz> and 4 Martin Liska <mliska@suse.cz>. 5 6 This file is part of GCC. 7 8 GCC is free software; you can redistribute it and/or modify 9 it under the terms of the GNU General Public License as published by 10 the Free Software Foundation; either version 3, or (at your option) 11 any later version. 12 13 GCC is distributed in the hope that it will be useful, 14 but WITHOUT ANY WARRANTY; without even the implied warranty of 15 MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the 16 GNU General Public License 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 "config.h" 23 #include "system.h" 24 #include "coretypes.h" 25 #include "memmodel.h" 26 #include "tm.h" 27 #include "is-a.h" 28 #include "hash-table.h" 29 #include "vec.h" 30 #include "tree.h" 31 #include "tree-pass.h" 32 #include "function.h" 33 #include "basic-block.h" 34 #include "cfg.h" 35 #include "fold-const.h" 36 #include "gimple.h" 37 #include "gimple-iterator.h" 38 #include "bitmap.h" 39 #include "dumpfile.h" 40 #include "gimple-pretty-print.h" 41 #include "diagnostic-core.h" 42 #include "gimple-ssa.h" 43 #include "tree-phinodes.h" 44 #include "stringpool.h" 45 #include "tree-vrp.h" 46 #include "tree-ssanames.h" 47 #include "tree-dfa.h" 48 #include "ssa-iterators.h" 49 #include "cgraph.h" 50 #include "print-tree.h" 51 #include "symbol-summary.h" 52 #include "hsa-common.h" 53 #include "cfghooks.h" 54 #include "tree-cfg.h" 55 #include "cfgloop.h" 56 #include "cfganal.h" 57 #include "builtins.h" 58 #include "params.h" 59 #include "gomp-constants.h" 60 #include "internal-fn.h" 61 #include "builtins.h" 62 #include "stor-layout.h" 63 #include "stringpool.h" 64 #include "attribs.h" 65 66 /* Print a warning message and set that we have seen an error. */ 67 68 #define HSA_SORRY_ATV(location, message, ...) \ 69 do \ 70 { \ 71 hsa_fail_cfun (); \ 72 if (warning_at (EXPR_LOCATION (hsa_cfun->m_decl), OPT_Whsa, \ 73 HSA_SORRY_MSG)) \ 74 inform (location, message, __VA_ARGS__); \ 75 } \ 76 while (false) 77 78 /* Same as previous, but highlight a location. */ 79 80 #define HSA_SORRY_AT(location, message) \ 81 do \ 82 { \ 83 hsa_fail_cfun (); \ 84 if (warning_at (EXPR_LOCATION (hsa_cfun->m_decl), OPT_Whsa, \ 85 HSA_SORRY_MSG)) \ 86 inform (location, message); \ 87 } \ 88 while (false) 89 90 /* Default number of threads used by kernel dispatch. */ 91 92 #define HSA_DEFAULT_NUM_THREADS 64 93 94 /* Following structures are defined in the final version 95 of HSA specification. */ 96 97 /* HSA queue packet is shadow structure, originally provided by AMD. */ 98 99 struct hsa_queue_packet 100 { 101 uint16_t header; 102 uint16_t setup; 103 uint16_t workgroup_size_x; 104 uint16_t workgroup_size_y; 105 uint16_t workgroup_size_z; 106 uint16_t reserved0; 107 uint32_t grid_size_x; 108 uint32_t grid_size_y; 109 uint32_t grid_size_z; 110 uint32_t private_segment_size; 111 uint32_t group_segment_size; 112 uint64_t kernel_object; 113 void *kernarg_address; 114 uint64_t reserved2; 115 uint64_t completion_signal; 116 }; 117 118 /* HSA queue is shadow structure, originally provided by AMD. */ 119 120 struct hsa_queue 121 { 122 int type; 123 uint32_t features; 124 void *base_address; 125 uint64_t doorbell_signal; 126 uint32_t size; 127 uint32_t reserved1; 128 uint64_t id; 129 }; 130 131 static struct obstack hsa_obstack; 132 133 /* List of pointers to all instructions that come from an object allocator. */ 134 static vec <hsa_insn_basic *> hsa_instructions; 135 136 /* List of pointers to all operands that come from an object allocator. */ 137 static vec <hsa_op_base *> hsa_operands; 138 139 hsa_symbol::hsa_symbol () 140 : m_decl (NULL_TREE), m_name (NULL), m_name_number (0), 141 m_directive_offset (0), m_type (BRIG_TYPE_NONE), 142 m_segment (BRIG_SEGMENT_NONE), m_linkage (BRIG_LINKAGE_NONE), m_dim (0), 143 m_cst_value (NULL), m_global_scope_p (false), m_seen_error (false), 144 m_allocation (BRIG_ALLOCATION_AUTOMATIC), m_emitted_to_brig (false) 145 { 146 } 147 148 149 hsa_symbol::hsa_symbol (BrigType16_t type, BrigSegment8_t segment, 150 BrigLinkage8_t linkage, bool global_scope_p, 151 BrigAllocation allocation, BrigAlignment8_t align) 152 : m_decl (NULL_TREE), m_name (NULL), m_name_number (0), 153 m_directive_offset (0), m_type (type), m_segment (segment), 154 m_linkage (linkage), m_dim (0), m_cst_value (NULL), 155 m_global_scope_p (global_scope_p), m_seen_error (false), 156 m_allocation (allocation), m_emitted_to_brig (false), m_align (align) 157 { 158 } 159 160 unsigned HOST_WIDE_INT 161 hsa_symbol::total_byte_size () 162 { 163 unsigned HOST_WIDE_INT s 164 = hsa_type_bit_size (~BRIG_TYPE_ARRAY_MASK & m_type); 165 gcc_assert (s % BITS_PER_UNIT == 0); 166 s /= BITS_PER_UNIT; 167 168 if (m_dim) 169 s *= m_dim; 170 171 return s; 172 } 173 174 /* Forward declaration. */ 175 176 static BrigType16_t 177 hsa_type_for_tree_type (const_tree type, unsigned HOST_WIDE_INT *dim_p, 178 bool min32int); 179 180 void 181 hsa_symbol::fillup_for_decl (tree decl) 182 { 183 m_decl = decl; 184 m_type = hsa_type_for_tree_type (TREE_TYPE (decl), &m_dim, false); 185 if (hsa_seen_error ()) 186 { 187 m_seen_error = true; 188 return; 189 } 190 191 m_align = MAX (m_align, hsa_natural_alignment (m_type)); 192 } 193 194 /* Constructor of class representing global HSA function/kernel information and 195 state. FNDECL is function declaration, KERNEL_P is true if the function 196 is going to become a HSA kernel. If the function has body, SSA_NAMES_COUNT 197 should be set to number of SSA names used in the function. 198 MODIFIED_CFG is set to true in case we modified control-flow graph 199 of the function. */ 200 201 hsa_function_representation::hsa_function_representation 202 (tree fdecl, bool kernel_p, unsigned ssa_names_count, bool modified_cfg) 203 : m_name (NULL), 204 m_reg_count (0), m_input_args (vNULL), 205 m_output_arg (NULL), m_spill_symbols (vNULL), m_global_symbols (vNULL), 206 m_private_variables (vNULL), m_called_functions (vNULL), 207 m_called_internal_fns (vNULL), m_hbb_count (0), 208 m_in_ssa (true), m_kern_p (kernel_p), m_declaration_p (false), 209 m_decl (fdecl), m_internal_fn (NULL), m_shadow_reg (NULL), 210 m_kernel_dispatch_count (0), m_maximum_omp_data_size (0), 211 m_seen_error (false), m_temp_symbol_count (0), m_ssa_map (), 212 m_modified_cfg (modified_cfg) 213 { 214 int sym_init_len = (vec_safe_length (cfun->local_decls) / 2) + 1; 215 m_local_symbols = new hash_table <hsa_noop_symbol_hasher> (sym_init_len); 216 m_ssa_map.safe_grow_cleared (ssa_names_count); 217 } 218 219 /* Constructor of class representing HSA function information that 220 is derived for an internal function. */ 221 hsa_function_representation::hsa_function_representation (hsa_internal_fn *fn) 222 : m_reg_count (0), m_input_args (vNULL), 223 m_output_arg (NULL), m_local_symbols (NULL), 224 m_spill_symbols (vNULL), m_global_symbols (vNULL), 225 m_private_variables (vNULL), m_called_functions (vNULL), 226 m_called_internal_fns (vNULL), m_hbb_count (0), 227 m_in_ssa (true), m_kern_p (false), m_declaration_p (true), m_decl (NULL), 228 m_internal_fn (fn), m_shadow_reg (NULL), m_kernel_dispatch_count (0), 229 m_maximum_omp_data_size (0), m_seen_error (false), m_temp_symbol_count (0), 230 m_ssa_map () {} 231 232 /* Destructor of class holding function/kernel-wide information and state. */ 233 234 hsa_function_representation::~hsa_function_representation () 235 { 236 /* Kernel names are deallocated at the end of BRIG output when deallocating 237 hsa_decl_kernel_mapping. */ 238 if (!m_kern_p || m_seen_error) 239 free (m_name); 240 241 for (unsigned i = 0; i < m_input_args.length (); i++) 242 delete m_input_args[i]; 243 m_input_args.release (); 244 245 delete m_output_arg; 246 delete m_local_symbols; 247 248 for (unsigned i = 0; i < m_spill_symbols.length (); i++) 249 delete m_spill_symbols[i]; 250 m_spill_symbols.release (); 251 252 hsa_symbol *sym; 253 for (unsigned i = 0; i < m_global_symbols.iterate (i, &sym); i++) 254 if (sym->m_linkage != BRIG_ALLOCATION_PROGRAM) 255 delete sym; 256 m_global_symbols.release (); 257 258 for (unsigned i = 0; i < m_private_variables.length (); i++) 259 delete m_private_variables[i]; 260 m_private_variables.release (); 261 m_called_functions.release (); 262 m_ssa_map.release (); 263 264 for (unsigned i = 0; i < m_called_internal_fns.length (); i++) 265 delete m_called_internal_fns[i]; 266 } 267 268 hsa_op_reg * 269 hsa_function_representation::get_shadow_reg () 270 { 271 /* If we compile a function with kernel dispatch and does not set 272 an optimization level, the function won't be inlined and 273 we return NULL. */ 274 if (!m_kern_p) 275 return NULL; 276 277 if (m_shadow_reg) 278 return m_shadow_reg; 279 280 /* Append the shadow argument. */ 281 hsa_symbol *shadow = new hsa_symbol (BRIG_TYPE_U64, BRIG_SEGMENT_KERNARG, 282 BRIG_LINKAGE_FUNCTION); 283 m_input_args.safe_push (shadow); 284 shadow->m_name = "hsa_runtime_shadow"; 285 286 hsa_op_reg *r = new hsa_op_reg (BRIG_TYPE_U64); 287 hsa_op_address *addr = new hsa_op_address (shadow); 288 289 hsa_insn_mem *mem = new hsa_insn_mem (BRIG_OPCODE_LD, BRIG_TYPE_U64, r, addr); 290 hsa_bb_for_bb (ENTRY_BLOCK_PTR_FOR_FN (cfun))->append_insn (mem); 291 m_shadow_reg = r; 292 293 return r; 294 } 295 296 bool hsa_function_representation::has_shadow_reg_p () 297 { 298 return m_shadow_reg != NULL; 299 } 300 301 void 302 hsa_function_representation::init_extra_bbs () 303 { 304 hsa_init_new_bb (ENTRY_BLOCK_PTR_FOR_FN (cfun)); 305 hsa_init_new_bb (EXIT_BLOCK_PTR_FOR_FN (cfun)); 306 } 307 308 void 309 hsa_function_representation::update_dominance () 310 { 311 if (m_modified_cfg) 312 { 313 free_dominance_info (CDI_DOMINATORS); 314 calculate_dominance_info (CDI_DOMINATORS); 315 } 316 } 317 318 hsa_symbol * 319 hsa_function_representation::create_hsa_temporary (BrigType16_t type) 320 { 321 hsa_symbol *s = new hsa_symbol (type, BRIG_SEGMENT_PRIVATE, 322 BRIG_LINKAGE_FUNCTION); 323 s->m_name_number = m_temp_symbol_count++; 324 325 hsa_cfun->m_private_variables.safe_push (s); 326 return s; 327 } 328 329 BrigLinkage8_t 330 hsa_function_representation::get_linkage () 331 { 332 if (m_internal_fn) 333 return BRIG_LINKAGE_PROGRAM; 334 335 return m_kern_p || TREE_PUBLIC (m_decl) ? 336 BRIG_LINKAGE_PROGRAM : BRIG_LINKAGE_MODULE; 337 } 338 339 /* Hash map of simple OMP builtins. */ 340 static hash_map <nofree_string_hash, omp_simple_builtin> *omp_simple_builtins 341 = NULL; 342 343 /* Warning messages for OMP builtins. */ 344 345 #define HSA_WARN_LOCK_ROUTINE "support for HSA does not implement OpenMP " \ 346 "lock routines" 347 #define HSA_WARN_TIMING_ROUTINE "support for HSA does not implement OpenMP " \ 348 "timing routines" 349 #define HSA_WARN_MEMORY_ROUTINE "OpenMP device memory library routines have " \ 350 "undefined semantics within target regions, support for HSA ignores them" 351 #define HSA_WARN_AFFINITY "Support for HSA does not implement OpenMP " \ 352 "affinity feateres" 353 354 /* Initialize hash map with simple OMP builtins. */ 355 356 static void 357 hsa_init_simple_builtins () 358 { 359 if (omp_simple_builtins != NULL) 360 return; 361 362 omp_simple_builtins 363 = new hash_map <nofree_string_hash, omp_simple_builtin> (); 364 365 omp_simple_builtin omp_builtins[] = 366 { 367 omp_simple_builtin ("omp_get_initial_device", NULL, false, 368 new hsa_op_immed (GOMP_DEVICE_HOST, 369 (BrigType16_t) BRIG_TYPE_S32)), 370 omp_simple_builtin ("omp_is_initial_device", NULL, false, 371 new hsa_op_immed (0, (BrigType16_t) BRIG_TYPE_S32)), 372 omp_simple_builtin ("omp_get_dynamic", NULL, false, 373 new hsa_op_immed (0, (BrigType16_t) BRIG_TYPE_S32)), 374 omp_simple_builtin ("omp_set_dynamic", NULL, false, NULL), 375 omp_simple_builtin ("omp_init_lock", HSA_WARN_LOCK_ROUTINE, true), 376 omp_simple_builtin ("omp_init_lock_with_hint", HSA_WARN_LOCK_ROUTINE, 377 true), 378 omp_simple_builtin ("omp_init_nest_lock_with_hint", HSA_WARN_LOCK_ROUTINE, 379 true), 380 omp_simple_builtin ("omp_destroy_lock", HSA_WARN_LOCK_ROUTINE, true), 381 omp_simple_builtin ("omp_set_lock", HSA_WARN_LOCK_ROUTINE, true), 382 omp_simple_builtin ("omp_unset_lock", HSA_WARN_LOCK_ROUTINE, true), 383 omp_simple_builtin ("omp_test_lock", HSA_WARN_LOCK_ROUTINE, true), 384 omp_simple_builtin ("omp_get_wtime", HSA_WARN_TIMING_ROUTINE, true), 385 omp_simple_builtin ("omp_get_wtick", HSA_WARN_TIMING_ROUTINE, true), 386 omp_simple_builtin ("omp_target_alloc", HSA_WARN_MEMORY_ROUTINE, false, 387 new hsa_op_immed (0, (BrigType16_t) BRIG_TYPE_U64)), 388 omp_simple_builtin ("omp_target_free", HSA_WARN_MEMORY_ROUTINE, false), 389 omp_simple_builtin ("omp_target_is_present", HSA_WARN_MEMORY_ROUTINE, 390 false, 391 new hsa_op_immed (-1, (BrigType16_t) BRIG_TYPE_S32)), 392 omp_simple_builtin ("omp_target_memcpy", HSA_WARN_MEMORY_ROUTINE, false, 393 new hsa_op_immed (-1, (BrigType16_t) BRIG_TYPE_S32)), 394 omp_simple_builtin ("omp_target_memcpy_rect", HSA_WARN_MEMORY_ROUTINE, 395 false, 396 new hsa_op_immed (-1, (BrigType16_t) BRIG_TYPE_S32)), 397 omp_simple_builtin ("omp_target_associate_ptr", HSA_WARN_MEMORY_ROUTINE, 398 false, 399 new hsa_op_immed (-1, (BrigType16_t) BRIG_TYPE_S32)), 400 omp_simple_builtin ("omp_target_disassociate_ptr", 401 HSA_WARN_MEMORY_ROUTINE, 402 false, 403 new hsa_op_immed (-1, (BrigType16_t) BRIG_TYPE_S32)), 404 omp_simple_builtin ("omp_set_max_active_levels", 405 "Support for HSA only allows only one active level, " 406 "call to omp_set_max_active_levels will be ignored " 407 "in the generated HSAIL", 408 false, NULL), 409 omp_simple_builtin ("omp_get_max_active_levels", NULL, false, 410 new hsa_op_immed (1, (BrigType16_t) BRIG_TYPE_S32)), 411 omp_simple_builtin ("omp_in_final", NULL, false, 412 new hsa_op_immed (0, (BrigType16_t) BRIG_TYPE_S32)), 413 omp_simple_builtin ("omp_get_proc_bind", HSA_WARN_AFFINITY, false, 414 new hsa_op_immed (0, (BrigType16_t) BRIG_TYPE_S32)), 415 omp_simple_builtin ("omp_get_num_places", HSA_WARN_AFFINITY, false, 416 new hsa_op_immed (0, (BrigType16_t) BRIG_TYPE_S32)), 417 omp_simple_builtin ("omp_get_place_num_procs", HSA_WARN_AFFINITY, false, 418 new hsa_op_immed (0, (BrigType16_t) BRIG_TYPE_S32)), 419 omp_simple_builtin ("omp_get_place_proc_ids", HSA_WARN_AFFINITY, false, 420 NULL), 421 omp_simple_builtin ("omp_get_place_num", HSA_WARN_AFFINITY, false, 422 new hsa_op_immed (-1, (BrigType16_t) BRIG_TYPE_S32)), 423 omp_simple_builtin ("omp_get_partition_num_places", HSA_WARN_AFFINITY, 424 false, 425 new hsa_op_immed (0, (BrigType16_t) BRIG_TYPE_S32)), 426 omp_simple_builtin ("omp_get_partition_place_nums", HSA_WARN_AFFINITY, 427 false, NULL), 428 omp_simple_builtin ("omp_set_default_device", 429 "omp_set_default_device has undefined semantics " 430 "within target regions, support for HSA ignores it", 431 false, NULL), 432 omp_simple_builtin ("omp_get_default_device", 433 "omp_get_default_device has undefined semantics " 434 "within target regions, support for HSA ignores it", 435 false, 436 new hsa_op_immed (0, (BrigType16_t) BRIG_TYPE_S32)), 437 omp_simple_builtin ("omp_get_num_devices", 438 "omp_get_num_devices has undefined semantics " 439 "within target regions, support for HSA ignores it", 440 false, 441 new hsa_op_immed (0, (BrigType16_t) BRIG_TYPE_S32)), 442 omp_simple_builtin ("omp_get_num_procs", NULL, true, NULL), 443 omp_simple_builtin ("omp_get_cancellation", NULL, true, NULL), 444 omp_simple_builtin ("omp_set_nested", NULL, true, NULL), 445 omp_simple_builtin ("omp_get_nested", NULL, true, NULL), 446 omp_simple_builtin ("omp_set_schedule", NULL, true, NULL), 447 omp_simple_builtin ("omp_get_schedule", NULL, true, NULL), 448 omp_simple_builtin ("omp_get_thread_limit", NULL, true, NULL), 449 omp_simple_builtin ("omp_get_team_size", NULL, true, NULL), 450 omp_simple_builtin ("omp_get_ancestor_thread_num", NULL, true, NULL), 451 omp_simple_builtin ("omp_get_max_task_priority", NULL, true, NULL) 452 }; 453 454 unsigned count = sizeof (omp_builtins) / sizeof (omp_simple_builtin); 455 456 for (unsigned i = 0; i < count; i++) 457 omp_simple_builtins->put (omp_builtins[i].m_name, omp_builtins[i]); 458 } 459 460 /* Allocate HSA structures that we need only while generating with this. */ 461 462 static void 463 hsa_init_data_for_cfun () 464 { 465 hsa_init_compilation_unit_data (); 466 gcc_obstack_init (&hsa_obstack); 467 } 468 469 /* Deinitialize HSA subsystem and free all allocated memory. */ 470 471 static void 472 hsa_deinit_data_for_cfun (void) 473 { 474 basic_block bb; 475 476 FOR_ALL_BB_FN (bb, cfun) 477 if (bb->aux) 478 { 479 hsa_bb *hbb = hsa_bb_for_bb (bb); 480 hbb->~hsa_bb (); 481 bb->aux = NULL; 482 } 483 484 for (unsigned int i = 0; i < hsa_operands.length (); i++) 485 hsa_destroy_operand (hsa_operands[i]); 486 487 hsa_operands.release (); 488 489 for (unsigned i = 0; i < hsa_instructions.length (); i++) 490 hsa_destroy_insn (hsa_instructions[i]); 491 492 hsa_instructions.release (); 493 494 if (omp_simple_builtins != NULL) 495 { 496 delete omp_simple_builtins; 497 omp_simple_builtins = NULL; 498 } 499 500 obstack_free (&hsa_obstack, NULL); 501 delete hsa_cfun; 502 } 503 504 /* Return the type which holds addresses in the given SEGMENT. */ 505 506 static BrigType16_t 507 hsa_get_segment_addr_type (BrigSegment8_t segment) 508 { 509 switch (segment) 510 { 511 case BRIG_SEGMENT_NONE: 512 gcc_unreachable (); 513 514 case BRIG_SEGMENT_FLAT: 515 case BRIG_SEGMENT_GLOBAL: 516 case BRIG_SEGMENT_READONLY: 517 case BRIG_SEGMENT_KERNARG: 518 return hsa_machine_large_p () ? BRIG_TYPE_U64 : BRIG_TYPE_U32; 519 520 case BRIG_SEGMENT_GROUP: 521 case BRIG_SEGMENT_PRIVATE: 522 case BRIG_SEGMENT_SPILL: 523 case BRIG_SEGMENT_ARG: 524 return BRIG_TYPE_U32; 525 } 526 gcc_unreachable (); 527 } 528 529 /* Return integer brig type according to provided SIZE in bytes. If SIGN 530 is set to true, return signed integer type. */ 531 532 static BrigType16_t 533 get_integer_type_by_bytes (unsigned size, bool sign) 534 { 535 if (sign) 536 switch (size) 537 { 538 case 1: 539 return BRIG_TYPE_S8; 540 case 2: 541 return BRIG_TYPE_S16; 542 case 4: 543 return BRIG_TYPE_S32; 544 case 8: 545 return BRIG_TYPE_S64; 546 default: 547 break; 548 } 549 else 550 switch (size) 551 { 552 case 1: 553 return BRIG_TYPE_U8; 554 case 2: 555 return BRIG_TYPE_U16; 556 case 4: 557 return BRIG_TYPE_U32; 558 case 8: 559 return BRIG_TYPE_U64; 560 default: 561 break; 562 } 563 564 return 0; 565 } 566 567 /* If T points to an integral type smaller than 32 bits, change it to a 32bit 568 equivalent and return the result. Otherwise just return the result. */ 569 570 static BrigType16_t 571 hsa_extend_inttype_to_32bit (BrigType16_t t) 572 { 573 if (t == BRIG_TYPE_U8 || t == BRIG_TYPE_U16) 574 return BRIG_TYPE_U32; 575 else if (t == BRIG_TYPE_S8 || t == BRIG_TYPE_S16) 576 return BRIG_TYPE_S32; 577 return t; 578 } 579 580 /* Return HSA type for tree TYPE, which has to fit into BrigType16_t. Pointers 581 are assumed to use flat addressing. If min32int is true, always expand 582 integer types to one that has at least 32 bits. */ 583 584 static BrigType16_t 585 hsa_type_for_scalar_tree_type (const_tree type, bool min32int) 586 { 587 HOST_WIDE_INT bsize; 588 const_tree base; 589 BrigType16_t res = BRIG_TYPE_NONE; 590 591 gcc_checking_assert (TYPE_P (type)); 592 gcc_checking_assert (!AGGREGATE_TYPE_P (type)); 593 if (POINTER_TYPE_P (type)) 594 return hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT); 595 596 if (TREE_CODE (type) == VECTOR_TYPE) 597 base = TREE_TYPE (type); 598 else if (TREE_CODE (type) == COMPLEX_TYPE) 599 { 600 base = TREE_TYPE (type); 601 min32int = true; 602 } 603 else 604 base = type; 605 606 if (!tree_fits_uhwi_p (TYPE_SIZE (base))) 607 { 608 HSA_SORRY_ATV (EXPR_LOCATION (type), 609 "support for HSA does not implement huge or " 610 "variable-sized type %qT", type); 611 return res; 612 } 613 614 bsize = tree_to_uhwi (TYPE_SIZE (base)); 615 unsigned byte_size = bsize / BITS_PER_UNIT; 616 if (INTEGRAL_TYPE_P (base)) 617 res = get_integer_type_by_bytes (byte_size, !TYPE_UNSIGNED (base)); 618 else if (SCALAR_FLOAT_TYPE_P (base)) 619 { 620 switch (bsize) 621 { 622 case 16: 623 res = BRIG_TYPE_F16; 624 break; 625 case 32: 626 res = BRIG_TYPE_F32; 627 break; 628 case 64: 629 res = BRIG_TYPE_F64; 630 break; 631 default: 632 break; 633 } 634 } 635 636 if (res == BRIG_TYPE_NONE) 637 { 638 HSA_SORRY_ATV (EXPR_LOCATION (type), 639 "support for HSA does not implement type %qT", type); 640 return res; 641 } 642 643 if (TREE_CODE (type) == VECTOR_TYPE) 644 { 645 HOST_WIDE_INT tsize = tree_to_uhwi (TYPE_SIZE (type)); 646 647 if (bsize == tsize) 648 { 649 HSA_SORRY_ATV (EXPR_LOCATION (type), 650 "support for HSA does not implement a vector type " 651 "where a type and unit size are equal: %qT", type); 652 return res; 653 } 654 655 switch (tsize) 656 { 657 case 32: 658 res |= BRIG_TYPE_PACK_32; 659 break; 660 case 64: 661 res |= BRIG_TYPE_PACK_64; 662 break; 663 case 128: 664 res |= BRIG_TYPE_PACK_128; 665 break; 666 default: 667 HSA_SORRY_ATV (EXPR_LOCATION (type), 668 "support for HSA does not implement type %qT", type); 669 } 670 } 671 672 if (min32int) 673 /* Registers/immediate operands can only be 32bit or more except for 674 f16. */ 675 res = hsa_extend_inttype_to_32bit (res); 676 677 if (TREE_CODE (type) == COMPLEX_TYPE) 678 { 679 unsigned bsize = 2 * hsa_type_bit_size (res); 680 res = hsa_bittype_for_bitsize (bsize); 681 } 682 683 return res; 684 } 685 686 /* Returns the BRIG type we need to load/store entities of TYPE. */ 687 688 static BrigType16_t 689 mem_type_for_type (BrigType16_t type) 690 { 691 /* HSA has non-intuitive constraints on load/store types. If it's 692 a bit-type it _must_ be B128, if it's not a bit-type it must be 693 64bit max. So for loading entities of 128 bits (e.g. vectors) 694 we have to use B128, while for loading the rest we have to use the 695 input type (??? or maybe also flattened to a equally sized non-vector 696 unsigned type?). */ 697 if ((type & BRIG_TYPE_PACK_MASK) == BRIG_TYPE_PACK_128) 698 return BRIG_TYPE_B128; 699 else if (hsa_btype_p (type) || hsa_type_packed_p (type)) 700 { 701 unsigned bitsize = hsa_type_bit_size (type); 702 if (bitsize < 128) 703 return hsa_uint_for_bitsize (bitsize); 704 else 705 return hsa_bittype_for_bitsize (bitsize); 706 } 707 return type; 708 } 709 710 /* Return HSA type for tree TYPE. If it cannot fit into BrigType16_t, some 711 kind of array will be generated, setting DIM appropriately. Otherwise, it 712 will be set to zero. */ 713 714 static BrigType16_t 715 hsa_type_for_tree_type (const_tree type, unsigned HOST_WIDE_INT *dim_p = NULL, 716 bool min32int = false) 717 { 718 gcc_checking_assert (TYPE_P (type)); 719 if (!tree_fits_uhwi_p (TYPE_SIZE_UNIT (type))) 720 { 721 HSA_SORRY_ATV (EXPR_LOCATION (type), "support for HSA does not " 722 "implement huge or variable-sized type %qT", type); 723 return BRIG_TYPE_NONE; 724 } 725 726 if (RECORD_OR_UNION_TYPE_P (type)) 727 { 728 if (dim_p) 729 *dim_p = tree_to_uhwi (TYPE_SIZE_UNIT (type)); 730 return BRIG_TYPE_U8 | BRIG_TYPE_ARRAY; 731 } 732 733 if (TREE_CODE (type) == ARRAY_TYPE) 734 { 735 /* We try to be nice and use the real base-type when this is an array of 736 scalars and only resort to an array of bytes if the type is more 737 complex. */ 738 739 unsigned HOST_WIDE_INT dim = 1; 740 741 while (TREE_CODE (type) == ARRAY_TYPE) 742 { 743 tree domain = TYPE_DOMAIN (type); 744 if (!TYPE_MIN_VALUE (domain) 745 || !TYPE_MAX_VALUE (domain) 746 || !tree_fits_shwi_p (TYPE_MIN_VALUE (domain)) 747 || !tree_fits_shwi_p (TYPE_MAX_VALUE (domain))) 748 { 749 HSA_SORRY_ATV (EXPR_LOCATION (type), 750 "support for HSA does not implement array " 751 "%qT with unknown bounds", type); 752 return BRIG_TYPE_NONE; 753 } 754 HOST_WIDE_INT min = tree_to_shwi (TYPE_MIN_VALUE (domain)); 755 HOST_WIDE_INT max = tree_to_shwi (TYPE_MAX_VALUE (domain)); 756 dim = dim * (unsigned HOST_WIDE_INT) (max - min + 1); 757 type = TREE_TYPE (type); 758 } 759 760 BrigType16_t res; 761 if (RECORD_OR_UNION_TYPE_P (type)) 762 { 763 dim = dim * tree_to_uhwi (TYPE_SIZE_UNIT (type)); 764 res = BRIG_TYPE_U8; 765 } 766 else 767 res = hsa_type_for_scalar_tree_type (type, false); 768 769 if (dim_p) 770 *dim_p = dim; 771 return res | BRIG_TYPE_ARRAY; 772 } 773 774 /* Scalar case: */ 775 if (dim_p) 776 *dim_p = 0; 777 778 return hsa_type_for_scalar_tree_type (type, min32int); 779 } 780 781 /* Returns true if converting from STYPE into DTYPE needs the _CVT 782 opcode. If false a normal _MOV is enough. */ 783 784 static bool 785 hsa_needs_cvt (BrigType16_t dtype, BrigType16_t stype) 786 { 787 if (hsa_btype_p (dtype)) 788 return false; 789 790 /* float <-> int conversions are real converts. */ 791 if (hsa_type_float_p (dtype) != hsa_type_float_p (stype)) 792 return true; 793 /* When both types have different size, then we need CVT as well. */ 794 if (hsa_type_bit_size (dtype) != hsa_type_bit_size (stype)) 795 return true; 796 return false; 797 } 798 799 /* Return declaration name if it exists or create one from UID if it does not. 800 If DECL is a local variable, make UID part of its name. */ 801 802 const char * 803 hsa_get_declaration_name (tree decl) 804 { 805 if (!DECL_NAME (decl)) 806 { 807 char buf[64]; 808 snprintf (buf, 64, "__hsa_anon_%u", DECL_UID (decl)); 809 size_t len = strlen (buf); 810 char *copy = (char *) obstack_alloc (&hsa_obstack, len + 1); 811 memcpy (copy, buf, len + 1); 812 return copy; 813 } 814 815 tree name_tree; 816 if (TREE_CODE (decl) == FUNCTION_DECL 817 || (TREE_CODE (decl) == VAR_DECL && is_global_var (decl))) 818 name_tree = DECL_ASSEMBLER_NAME (decl); 819 else 820 name_tree = DECL_NAME (decl); 821 822 const char *name = IDENTIFIER_POINTER (name_tree); 823 /* User-defined assembly names have prepended asterisk symbol. */ 824 if (name[0] == '*') 825 name++; 826 827 if ((TREE_CODE (decl) == VAR_DECL) 828 && decl_function_context (decl)) 829 { 830 size_t len = strlen (name); 831 char *buf = (char *) alloca (len + 32); 832 snprintf (buf, len + 32, "%s_%u", name, DECL_UID (decl)); 833 len = strlen (buf); 834 char *copy = (char *) obstack_alloc (&hsa_obstack, len + 1); 835 memcpy (copy, buf, len + 1); 836 return copy; 837 } 838 else 839 return name; 840 } 841 842 /* Lookup or create the associated hsa_symbol structure with a given VAR_DECL 843 or lookup the hsa_structure corresponding to a PARM_DECL. */ 844 845 static hsa_symbol * 846 get_symbol_for_decl (tree decl) 847 { 848 hsa_symbol **slot; 849 hsa_symbol dummy (BRIG_TYPE_NONE, BRIG_SEGMENT_NONE, BRIG_LINKAGE_NONE); 850 851 gcc_assert (TREE_CODE (decl) == PARM_DECL 852 || TREE_CODE (decl) == RESULT_DECL 853 || TREE_CODE (decl) == VAR_DECL 854 || TREE_CODE (decl) == CONST_DECL); 855 856 dummy.m_decl = decl; 857 858 bool is_in_global_vars = ((TREE_CODE (decl) == VAR_DECL) 859 && !decl_function_context (decl)); 860 861 if (is_in_global_vars) 862 slot = hsa_global_variable_symbols->find_slot (&dummy, INSERT); 863 else 864 slot = hsa_cfun->m_local_symbols->find_slot (&dummy, INSERT); 865 866 gcc_checking_assert (slot); 867 if (*slot) 868 { 869 hsa_symbol *sym = (*slot); 870 871 /* If the symbol is problematic, mark current function also as 872 problematic. */ 873 if (sym->m_seen_error) 874 hsa_fail_cfun (); 875 876 /* PR hsa/70234: If a global variable was marked to be emitted, 877 but HSAIL generation of a function using the variable fails, 878 we should retry to emit the variable in context of a different 879 function. 880 881 Iterate elements whether a symbol is already in m_global_symbols 882 of not. */ 883 if (is_in_global_vars && !sym->m_emitted_to_brig) 884 { 885 for (unsigned i = 0; i < hsa_cfun->m_global_symbols.length (); i++) 886 if (hsa_cfun->m_global_symbols[i] == sym) 887 return *slot; 888 hsa_cfun->m_global_symbols.safe_push (sym); 889 } 890 891 return *slot; 892 } 893 else 894 { 895 hsa_symbol *sym; 896 /* PARM_DECLs and RESULT_DECL should be already in m_local_symbols. */ 897 gcc_assert (TREE_CODE (decl) == VAR_DECL 898 || TREE_CODE (decl) == CONST_DECL); 899 BrigAlignment8_t align = hsa_object_alignment (decl); 900 901 if (is_in_global_vars) 902 { 903 gcc_checking_assert (TREE_CODE (decl) != CONST_DECL); 904 sym = new hsa_symbol (BRIG_TYPE_NONE, BRIG_SEGMENT_GLOBAL, 905 BRIG_LINKAGE_PROGRAM, true, 906 BRIG_ALLOCATION_PROGRAM, align); 907 hsa_cfun->m_global_symbols.safe_push (sym); 908 sym->fillup_for_decl (decl); 909 if (sym->m_align > align) 910 { 911 sym->m_seen_error = true; 912 HSA_SORRY_ATV (EXPR_LOCATION (decl), 913 "HSA specification requires that %E is at least " 914 "naturally aligned", decl); 915 } 916 } 917 else 918 { 919 /* As generation of efficient memory copy instructions relies 920 on alignment greater or equal to 8 bytes, 921 we need to increase alignment of all aggregate types.. */ 922 if (AGGREGATE_TYPE_P (TREE_TYPE (decl))) 923 align = MAX ((BrigAlignment8_t) BRIG_ALIGNMENT_8, align); 924 925 BrigAllocation allocation = BRIG_ALLOCATION_AUTOMATIC; 926 BrigSegment8_t segment; 927 if (TREE_CODE (decl) == CONST_DECL) 928 { 929 segment = BRIG_SEGMENT_READONLY; 930 allocation = BRIG_ALLOCATION_AGENT; 931 } 932 else if (lookup_attribute ("hsa_group_segment", 933 DECL_ATTRIBUTES (decl))) 934 segment = BRIG_SEGMENT_GROUP; 935 else if (TREE_STATIC (decl)) 936 { 937 segment = BRIG_SEGMENT_GLOBAL; 938 allocation = BRIG_ALLOCATION_PROGRAM; 939 } 940 else if (lookup_attribute ("hsa_global_segment", 941 DECL_ATTRIBUTES (decl))) 942 segment = BRIG_SEGMENT_GLOBAL; 943 else 944 segment = BRIG_SEGMENT_PRIVATE; 945 946 sym = new hsa_symbol (BRIG_TYPE_NONE, segment, BRIG_LINKAGE_FUNCTION, 947 false, allocation, align); 948 sym->fillup_for_decl (decl); 949 hsa_cfun->m_private_variables.safe_push (sym); 950 } 951 952 sym->m_name = hsa_get_declaration_name (decl); 953 *slot = sym; 954 return sym; 955 } 956 } 957 958 /* For a given HSA function declaration, return a host 959 function declaration. */ 960 961 tree 962 hsa_get_host_function (tree decl) 963 { 964 hsa_function_summary *s 965 = hsa_summaries->get (cgraph_node::get_create (decl)); 966 gcc_assert (s->m_kind != HSA_NONE); 967 gcc_assert (s->m_gpu_implementation_p); 968 969 return s->m_bound_function ? s->m_bound_function->decl : NULL; 970 } 971 972 /* Return true if function DECL has a host equivalent function. */ 973 974 static char * 975 get_brig_function_name (tree decl) 976 { 977 tree d = decl; 978 979 hsa_function_summary *s = hsa_summaries->get (cgraph_node::get_create (d)); 980 if (s->m_kind != HSA_NONE 981 && s->m_gpu_implementation_p 982 && s->m_bound_function) 983 d = s->m_bound_function->decl; 984 985 /* IPA split can create a function that has no host equivalent. */ 986 if (d == NULL) 987 d = decl; 988 989 char *name = xstrdup (hsa_get_declaration_name (d)); 990 hsa_sanitize_name (name); 991 992 return name; 993 } 994 995 /* Create a spill symbol of type TYPE. */ 996 997 hsa_symbol * 998 hsa_get_spill_symbol (BrigType16_t type) 999 { 1000 hsa_symbol *sym = new hsa_symbol (type, BRIG_SEGMENT_SPILL, 1001 BRIG_LINKAGE_FUNCTION); 1002 hsa_cfun->m_spill_symbols.safe_push (sym); 1003 return sym; 1004 } 1005 1006 /* Create a symbol for a read-only string constant. */ 1007 hsa_symbol * 1008 hsa_get_string_cst_symbol (tree string_cst) 1009 { 1010 gcc_checking_assert (TREE_CODE (string_cst) == STRING_CST); 1011 1012 hsa_symbol **slot = hsa_cfun->m_string_constants_map.get (string_cst); 1013 if (slot) 1014 return *slot; 1015 1016 hsa_op_immed *cst = new hsa_op_immed (string_cst); 1017 hsa_symbol *sym = new hsa_symbol (cst->m_type, BRIG_SEGMENT_GLOBAL, 1018 BRIG_LINKAGE_MODULE, true, 1019 BRIG_ALLOCATION_AGENT); 1020 sym->m_cst_value = cst; 1021 sym->m_dim = TREE_STRING_LENGTH (string_cst); 1022 sym->m_name_number = hsa_cfun->m_global_symbols.length (); 1023 1024 hsa_cfun->m_global_symbols.safe_push (sym); 1025 hsa_cfun->m_string_constants_map.put (string_cst, sym); 1026 return sym; 1027 } 1028 1029 /* Make the type of a MOV instruction larger if mandated by HSAIL rules. */ 1030 1031 static void 1032 hsa_fixup_mov_insn_type (hsa_insn_basic *insn) 1033 { 1034 insn->m_type = hsa_extend_inttype_to_32bit (insn->m_type); 1035 if (insn->m_type == BRIG_TYPE_B8 || insn->m_type == BRIG_TYPE_B16) 1036 insn->m_type = BRIG_TYPE_B32; 1037 } 1038 1039 /* Constructor of the ancestor of all operands. K is BRIG kind that identified 1040 what the operator is. */ 1041 1042 hsa_op_base::hsa_op_base (BrigKind16_t k) 1043 : m_next (NULL), m_brig_op_offset (0), m_kind (k) 1044 { 1045 hsa_operands.safe_push (this); 1046 } 1047 1048 /* Constructor of ancestor of all operands which have a type. K is BRIG kind 1049 that identified what the operator is. T is the type of the operator. */ 1050 1051 hsa_op_with_type::hsa_op_with_type (BrigKind16_t k, BrigType16_t t) 1052 : hsa_op_base (k), m_type (t) 1053 { 1054 } 1055 1056 hsa_op_with_type * 1057 hsa_op_with_type::get_in_type (BrigType16_t dtype, hsa_bb *hbb) 1058 { 1059 if (m_type == dtype) 1060 return this; 1061 1062 hsa_op_reg *dest; 1063 1064 if (hsa_needs_cvt (dtype, m_type)) 1065 { 1066 dest = new hsa_op_reg (dtype); 1067 hbb->append_insn (new hsa_insn_cvt (dest, this)); 1068 } 1069 else if (is_a <hsa_op_reg *> (this)) 1070 { 1071 /* In the end, HSA registers do not really have types, only sizes, so if 1072 the sizes match, we can use the register directly. */ 1073 gcc_checking_assert (hsa_type_bit_size (dtype) 1074 == hsa_type_bit_size (m_type)); 1075 return this; 1076 } 1077 else 1078 { 1079 dest = new hsa_op_reg (m_type); 1080 1081 hsa_insn_basic *mov = new hsa_insn_basic (2, BRIG_OPCODE_MOV, 1082 dest->m_type, dest, this); 1083 hsa_fixup_mov_insn_type (mov); 1084 hbb->append_insn (mov); 1085 /* We cannot simply for instance: 'mov_u32 $_3, 48 (s32)' because 1086 type of the operand must be same as type of the instruction. */ 1087 dest->m_type = dtype; 1088 } 1089 1090 return dest; 1091 } 1092 1093 /* If this operand has integer type smaller than 32 bits, extend it to 32 bits, 1094 adding instructions to HBB if needed. */ 1095 1096 hsa_op_with_type * 1097 hsa_op_with_type::extend_int_to_32bit (hsa_bb *hbb) 1098 { 1099 if (m_type == BRIG_TYPE_U8 || m_type == BRIG_TYPE_U16) 1100 return get_in_type (BRIG_TYPE_U32, hbb); 1101 else if (m_type == BRIG_TYPE_S8 || m_type == BRIG_TYPE_S16) 1102 return get_in_type (BRIG_TYPE_S32, hbb); 1103 else 1104 return this; 1105 } 1106 1107 /* Constructor of class representing HSA immediate values. TREE_VAL is the 1108 tree representation of the immediate value. If min32int is true, 1109 always expand integer types to one that has at least 32 bits. */ 1110 1111 hsa_op_immed::hsa_op_immed (tree tree_val, bool min32int) 1112 : hsa_op_with_type (BRIG_KIND_OPERAND_CONSTANT_BYTES, 1113 hsa_type_for_tree_type (TREE_TYPE (tree_val), NULL, 1114 min32int)) 1115 { 1116 if (hsa_seen_error ()) 1117 return; 1118 1119 gcc_checking_assert ((is_gimple_min_invariant (tree_val) 1120 && (!POINTER_TYPE_P (TREE_TYPE (tree_val)) 1121 || TREE_CODE (tree_val) == INTEGER_CST)) 1122 || TREE_CODE (tree_val) == CONSTRUCTOR); 1123 m_tree_value = tree_val; 1124 1125 /* Verify that all elements of a constructor are constants. */ 1126 if (TREE_CODE (m_tree_value) == CONSTRUCTOR) 1127 for (unsigned i = 0; i < CONSTRUCTOR_NELTS (m_tree_value); i++) 1128 { 1129 tree v = CONSTRUCTOR_ELT (m_tree_value, i)->value; 1130 if (!CONSTANT_CLASS_P (v)) 1131 { 1132 HSA_SORRY_AT (EXPR_LOCATION (tree_val), 1133 "HSA ctor should have only constants"); 1134 return; 1135 } 1136 } 1137 } 1138 1139 /* Constructor of class representing HSA immediate values. INTEGER_VALUE is the 1140 integer representation of the immediate value. TYPE is BRIG type. */ 1141 1142 hsa_op_immed::hsa_op_immed (HOST_WIDE_INT integer_value, BrigType16_t type) 1143 : hsa_op_with_type (BRIG_KIND_OPERAND_CONSTANT_BYTES, type), 1144 m_tree_value (NULL) 1145 { 1146 gcc_assert (hsa_type_integer_p (type)); 1147 m_int_value = integer_value; 1148 } 1149 1150 hsa_op_immed::hsa_op_immed () 1151 : hsa_op_with_type (BRIG_KIND_NONE, BRIG_TYPE_NONE) 1152 { 1153 } 1154 1155 /* New operator to allocate immediate operands from obstack. */ 1156 1157 void * 1158 hsa_op_immed::operator new (size_t size) 1159 { 1160 return obstack_alloc (&hsa_obstack, size); 1161 } 1162 1163 /* Destructor. */ 1164 1165 hsa_op_immed::~hsa_op_immed () 1166 { 1167 } 1168 1169 /* Change type of the immediate value to T. */ 1170 1171 void 1172 hsa_op_immed::set_type (BrigType16_t t) 1173 { 1174 m_type = t; 1175 } 1176 1177 /* Constructor of class representing HSA registers and pseudo-registers. T is 1178 the BRIG type of the new register. */ 1179 1180 hsa_op_reg::hsa_op_reg (BrigType16_t t) 1181 : hsa_op_with_type (BRIG_KIND_OPERAND_REGISTER, t), m_gimple_ssa (NULL_TREE), 1182 m_def_insn (NULL), m_spill_sym (NULL), m_order (hsa_cfun->m_reg_count++), 1183 m_lr_begin (0), m_lr_end (0), m_reg_class (0), m_hard_num (0) 1184 { 1185 } 1186 1187 /* New operator to allocate a register from obstack. */ 1188 1189 void * 1190 hsa_op_reg::operator new (size_t size) 1191 { 1192 return obstack_alloc (&hsa_obstack, size); 1193 } 1194 1195 /* Verify register operand. */ 1196 1197 void 1198 hsa_op_reg::verify_ssa () 1199 { 1200 /* Verify that each HSA register has a definition assigned. 1201 Exceptions are VAR_DECL and PARM_DECL that are a default 1202 definition. */ 1203 gcc_checking_assert (m_def_insn 1204 || (m_gimple_ssa != NULL 1205 && (!SSA_NAME_VAR (m_gimple_ssa) 1206 || (TREE_CODE (SSA_NAME_VAR (m_gimple_ssa)) 1207 != PARM_DECL)) 1208 && SSA_NAME_IS_DEFAULT_DEF (m_gimple_ssa))); 1209 1210 /* Verify that every use of the register is really present 1211 in an instruction. */ 1212 for (unsigned i = 0; i < m_uses.length (); i++) 1213 { 1214 hsa_insn_basic *use = m_uses[i]; 1215 1216 bool is_visited = false; 1217 for (unsigned j = 0; j < use->operand_count (); j++) 1218 { 1219 hsa_op_base *u = use->get_op (j); 1220 hsa_op_address *addr; addr = dyn_cast <hsa_op_address *> (u); 1221 if (addr && addr->m_reg) 1222 u = addr->m_reg; 1223 1224 if (u == this) 1225 { 1226 bool r = !addr && use->op_output_p (j); 1227 1228 if (r) 1229 { 1230 error ("HSA SSA name defined by instruction that is supposed " 1231 "to be using it"); 1232 debug_hsa_operand (this); 1233 debug_hsa_insn (use); 1234 internal_error ("HSA SSA verification failed"); 1235 } 1236 1237 is_visited = true; 1238 } 1239 } 1240 1241 if (!is_visited) 1242 { 1243 error ("HSA SSA name not among operands of instruction that is " 1244 "supposed to use it"); 1245 debug_hsa_operand (this); 1246 debug_hsa_insn (use); 1247 internal_error ("HSA SSA verification failed"); 1248 } 1249 } 1250 } 1251 1252 hsa_op_address::hsa_op_address (hsa_symbol *sym, hsa_op_reg *r, 1253 HOST_WIDE_INT offset) 1254 : hsa_op_base (BRIG_KIND_OPERAND_ADDRESS), m_symbol (sym), m_reg (r), 1255 m_imm_offset (offset) 1256 { 1257 } 1258 1259 hsa_op_address::hsa_op_address (hsa_symbol *sym, HOST_WIDE_INT offset) 1260 : hsa_op_base (BRIG_KIND_OPERAND_ADDRESS), m_symbol (sym), m_reg (NULL), 1261 m_imm_offset (offset) 1262 { 1263 } 1264 1265 hsa_op_address::hsa_op_address (hsa_op_reg *r, HOST_WIDE_INT offset) 1266 : hsa_op_base (BRIG_KIND_OPERAND_ADDRESS), m_symbol (NULL), m_reg (r), 1267 m_imm_offset (offset) 1268 { 1269 } 1270 1271 /* New operator to allocate address operands from obstack. */ 1272 1273 void * 1274 hsa_op_address::operator new (size_t size) 1275 { 1276 return obstack_alloc (&hsa_obstack, size); 1277 } 1278 1279 /* Constructor of an operand referring to HSAIL code. */ 1280 1281 hsa_op_code_ref::hsa_op_code_ref () : hsa_op_base (BRIG_KIND_OPERAND_CODE_REF), 1282 m_directive_offset (0) 1283 { 1284 } 1285 1286 /* Constructor of an operand representing a code list. Set it up so that it 1287 can contain ELEMENTS number of elements. */ 1288 1289 hsa_op_code_list::hsa_op_code_list (unsigned elements) 1290 : hsa_op_base (BRIG_KIND_OPERAND_CODE_LIST) 1291 { 1292 m_offsets.create (1); 1293 m_offsets.safe_grow_cleared (elements); 1294 } 1295 1296 /* New operator to allocate code list operands from obstack. */ 1297 1298 void * 1299 hsa_op_code_list::operator new (size_t size) 1300 { 1301 return obstack_alloc (&hsa_obstack, size); 1302 } 1303 1304 /* Constructor of an operand representing an operand list. 1305 Set it up so that it can contain ELEMENTS number of elements. */ 1306 1307 hsa_op_operand_list::hsa_op_operand_list (unsigned elements) 1308 : hsa_op_base (BRIG_KIND_OPERAND_OPERAND_LIST) 1309 { 1310 m_offsets.create (elements); 1311 m_offsets.safe_grow (elements); 1312 } 1313 1314 /* New operator to allocate operand list operands from obstack. */ 1315 1316 void * 1317 hsa_op_operand_list::operator new (size_t size) 1318 { 1319 return obstack_alloc (&hsa_obstack, size); 1320 } 1321 1322 hsa_op_operand_list::~hsa_op_operand_list () 1323 { 1324 m_offsets.release (); 1325 } 1326 1327 1328 hsa_op_reg * 1329 hsa_function_representation::reg_for_gimple_ssa (tree ssa) 1330 { 1331 hsa_op_reg *hreg; 1332 1333 gcc_checking_assert (TREE_CODE (ssa) == SSA_NAME); 1334 if (m_ssa_map[SSA_NAME_VERSION (ssa)]) 1335 return m_ssa_map[SSA_NAME_VERSION (ssa)]; 1336 1337 hreg = new hsa_op_reg (hsa_type_for_scalar_tree_type (TREE_TYPE (ssa), 1338 false)); 1339 hreg->m_gimple_ssa = ssa; 1340 m_ssa_map[SSA_NAME_VERSION (ssa)] = hreg; 1341 1342 return hreg; 1343 } 1344 1345 void 1346 hsa_op_reg::set_definition (hsa_insn_basic *insn) 1347 { 1348 if (hsa_cfun->m_in_ssa) 1349 { 1350 gcc_checking_assert (!m_def_insn); 1351 m_def_insn = insn; 1352 } 1353 else 1354 m_def_insn = NULL; 1355 } 1356 1357 /* Constructor of the class which is the bases of all instructions and directly 1358 represents the most basic ones. NOPS is the number of operands that the 1359 operand vector will contain (and which will be cleared). OP is the opcode 1360 of the instruction. This constructor does not set type. */ 1361 1362 hsa_insn_basic::hsa_insn_basic (unsigned nops, int opc) 1363 : m_prev (NULL), 1364 m_next (NULL), m_bb (NULL), m_opcode (opc), m_number (0), 1365 m_type (BRIG_TYPE_NONE), m_brig_offset (0) 1366 { 1367 if (nops > 0) 1368 m_operands.safe_grow_cleared (nops); 1369 1370 hsa_instructions.safe_push (this); 1371 } 1372 1373 /* Make OP the operand number INDEX of operands of this instruction. If OP is a 1374 register or an address containing a register, then either set the definition 1375 of the register to this instruction if it an output operand or add this 1376 instruction to the uses if it is an input one. */ 1377 1378 void 1379 hsa_insn_basic::set_op (int index, hsa_op_base *op) 1380 { 1381 /* Each address operand is always use. */ 1382 hsa_op_address *addr = dyn_cast <hsa_op_address *> (op); 1383 if (addr && addr->m_reg) 1384 addr->m_reg->m_uses.safe_push (this); 1385 else 1386 { 1387 hsa_op_reg *reg = dyn_cast <hsa_op_reg *> (op); 1388 if (reg) 1389 { 1390 if (op_output_p (index)) 1391 reg->set_definition (this); 1392 else 1393 reg->m_uses.safe_push (this); 1394 } 1395 } 1396 1397 m_operands[index] = op; 1398 } 1399 1400 /* Get INDEX-th operand of the instruction. */ 1401 1402 hsa_op_base * 1403 hsa_insn_basic::get_op (int index) 1404 { 1405 return m_operands[index]; 1406 } 1407 1408 /* Get address of INDEX-th operand of the instruction. */ 1409 1410 hsa_op_base ** 1411 hsa_insn_basic::get_op_addr (int index) 1412 { 1413 return &m_operands[index]; 1414 } 1415 1416 /* Get number of operands of the instruction. */ 1417 unsigned int 1418 hsa_insn_basic::operand_count () 1419 { 1420 return m_operands.length (); 1421 } 1422 1423 /* Constructor of the class which is the bases of all instructions and directly 1424 represents the most basic ones. NOPS is the number of operands that the 1425 operand vector will contain (and which will be cleared). OPC is the opcode 1426 of the instruction, T is the type of the instruction. */ 1427 1428 hsa_insn_basic::hsa_insn_basic (unsigned nops, int opc, BrigType16_t t, 1429 hsa_op_base *arg0, hsa_op_base *arg1, 1430 hsa_op_base *arg2, hsa_op_base *arg3) 1431 : m_prev (NULL), m_next (NULL), m_bb (NULL), m_opcode (opc),m_number (0), 1432 m_type (t), m_brig_offset (0) 1433 { 1434 if (nops > 0) 1435 m_operands.safe_grow_cleared (nops); 1436 1437 if (arg0 != NULL) 1438 { 1439 gcc_checking_assert (nops >= 1); 1440 set_op (0, arg0); 1441 } 1442 1443 if (arg1 != NULL) 1444 { 1445 gcc_checking_assert (nops >= 2); 1446 set_op (1, arg1); 1447 } 1448 1449 if (arg2 != NULL) 1450 { 1451 gcc_checking_assert (nops >= 3); 1452 set_op (2, arg2); 1453 } 1454 1455 if (arg3 != NULL) 1456 { 1457 gcc_checking_assert (nops >= 4); 1458 set_op (3, arg3); 1459 } 1460 1461 hsa_instructions.safe_push (this); 1462 } 1463 1464 /* New operator to allocate basic instruction from obstack. */ 1465 1466 void * 1467 hsa_insn_basic::operator new (size_t size) 1468 { 1469 return obstack_alloc (&hsa_obstack, size); 1470 } 1471 1472 /* Verify the instruction. */ 1473 1474 void 1475 hsa_insn_basic::verify () 1476 { 1477 hsa_op_address *addr; 1478 hsa_op_reg *reg; 1479 1480 /* Iterate all register operands and verify that the instruction 1481 is set in uses of the register. */ 1482 for (unsigned i = 0; i < operand_count (); i++) 1483 { 1484 hsa_op_base *use = get_op (i); 1485 1486 if ((addr = dyn_cast <hsa_op_address *> (use)) && addr->m_reg) 1487 { 1488 gcc_assert (addr->m_reg->m_def_insn != this); 1489 use = addr->m_reg; 1490 } 1491 1492 if ((reg = dyn_cast <hsa_op_reg *> (use)) && !op_output_p (i)) 1493 { 1494 unsigned j; 1495 for (j = 0; j < reg->m_uses.length (); j++) 1496 { 1497 if (reg->m_uses[j] == this) 1498 break; 1499 } 1500 1501 if (j == reg->m_uses.length ()) 1502 { 1503 error ("HSA instruction uses a register but is not among " 1504 "recorded register uses"); 1505 debug_hsa_operand (reg); 1506 debug_hsa_insn (this); 1507 internal_error ("HSA instruction verification failed"); 1508 } 1509 } 1510 } 1511 } 1512 1513 /* Constructor of an instruction representing a PHI node. NOPS is the number 1514 of operands (equal to the number of predecessors). */ 1515 1516 hsa_insn_phi::hsa_insn_phi (unsigned nops, hsa_op_reg *dst) 1517 : hsa_insn_basic (nops, HSA_OPCODE_PHI), m_dest (dst) 1518 { 1519 dst->set_definition (this); 1520 } 1521 1522 /* Constructor of class representing instructions for control flow and 1523 sychronization, */ 1524 1525 hsa_insn_br::hsa_insn_br (unsigned nops, int opc, BrigType16_t t, 1526 BrigWidth8_t width, hsa_op_base *arg0, 1527 hsa_op_base *arg1, hsa_op_base *arg2, 1528 hsa_op_base *arg3) 1529 : hsa_insn_basic (nops, opc, t, arg0, arg1, arg2, arg3), 1530 m_width (width) 1531 { 1532 } 1533 1534 /* Constructor of class representing instruction for conditional jump, CTRL is 1535 the control register determining whether the jump will be carried out, the 1536 new instruction is automatically added to its uses list. */ 1537 1538 hsa_insn_cbr::hsa_insn_cbr (hsa_op_reg *ctrl) 1539 : hsa_insn_br (1, BRIG_OPCODE_CBR, BRIG_TYPE_B1, BRIG_WIDTH_1, ctrl) 1540 { 1541 } 1542 1543 /* Constructor of class representing instruction for switch jump, CTRL is 1544 the index register. */ 1545 1546 hsa_insn_sbr::hsa_insn_sbr (hsa_op_reg *index, unsigned jump_count) 1547 : hsa_insn_basic (1, BRIG_OPCODE_SBR, BRIG_TYPE_B1, index), 1548 m_width (BRIG_WIDTH_1), m_jump_table (vNULL), 1549 m_label_code_list (new hsa_op_code_list (jump_count)) 1550 { 1551 } 1552 1553 /* Replace all occurrences of OLD_BB with NEW_BB in the statements 1554 jump table. */ 1555 1556 void 1557 hsa_insn_sbr::replace_all_labels (basic_block old_bb, basic_block new_bb) 1558 { 1559 for (unsigned i = 0; i < m_jump_table.length (); i++) 1560 if (m_jump_table[i] == old_bb) 1561 m_jump_table[i] = new_bb; 1562 } 1563 1564 hsa_insn_sbr::~hsa_insn_sbr () 1565 { 1566 m_jump_table.release (); 1567 } 1568 1569 /* Constructor of comparison instruction. CMP is the comparison operation and T 1570 is the result type. */ 1571 1572 hsa_insn_cmp::hsa_insn_cmp (BrigCompareOperation8_t cmp, BrigType16_t t, 1573 hsa_op_base *arg0, hsa_op_base *arg1, 1574 hsa_op_base *arg2) 1575 : hsa_insn_basic (3 , BRIG_OPCODE_CMP, t, arg0, arg1, arg2), m_compare (cmp) 1576 { 1577 } 1578 1579 /* Constructor of classes representing memory accesses. OPC is the opcode (must 1580 be BRIG_OPCODE_ST or BRIG_OPCODE_LD) and T is the type. The instruction 1581 operands are provided as ARG0 and ARG1. */ 1582 1583 hsa_insn_mem::hsa_insn_mem (int opc, BrigType16_t t, hsa_op_base *arg0, 1584 hsa_op_base *arg1) 1585 : hsa_insn_basic (2, opc, t, arg0, arg1), 1586 m_align (hsa_natural_alignment (t)), m_equiv_class (0) 1587 { 1588 gcc_checking_assert (opc == BRIG_OPCODE_LD || opc == BRIG_OPCODE_ST); 1589 } 1590 1591 /* Constructor for descendants allowing different opcodes and number of 1592 operands, it passes its arguments directly to hsa_insn_basic 1593 constructor. The instruction operands are provided as ARG[0-3]. */ 1594 1595 1596 hsa_insn_mem::hsa_insn_mem (unsigned nops, int opc, BrigType16_t t, 1597 hsa_op_base *arg0, hsa_op_base *arg1, 1598 hsa_op_base *arg2, hsa_op_base *arg3) 1599 : hsa_insn_basic (nops, opc, t, arg0, arg1, arg2, arg3), 1600 m_align (hsa_natural_alignment (t)), m_equiv_class (0) 1601 { 1602 } 1603 1604 /* Constructor of class representing atomic instructions. OPC is the principal 1605 opcode, AOP is the specific atomic operation opcode. T is the type of the 1606 instruction. The instruction operands are provided as ARG[0-3]. */ 1607 1608 hsa_insn_atomic::hsa_insn_atomic (int nops, int opc, 1609 enum BrigAtomicOperation aop, 1610 BrigType16_t t, BrigMemoryOrder memorder, 1611 hsa_op_base *arg0, 1612 hsa_op_base *arg1, hsa_op_base *arg2, 1613 hsa_op_base *arg3) 1614 : hsa_insn_mem (nops, opc, t, arg0, arg1, arg2, arg3), m_atomicop (aop), 1615 m_memoryorder (memorder), 1616 m_memoryscope (BRIG_MEMORY_SCOPE_SYSTEM) 1617 { 1618 gcc_checking_assert (opc == BRIG_OPCODE_ATOMICNORET || 1619 opc == BRIG_OPCODE_ATOMIC || 1620 opc == BRIG_OPCODE_SIGNAL || 1621 opc == BRIG_OPCODE_SIGNALNORET); 1622 } 1623 1624 /* Constructor of class representing signal instructions. OPC is the prinicpal 1625 opcode, SOP is the specific signal operation opcode. T is the type of the 1626 instruction. The instruction operands are provided as ARG[0-3]. */ 1627 1628 hsa_insn_signal::hsa_insn_signal (int nops, int opc, 1629 enum BrigAtomicOperation sop, 1630 BrigType16_t t, BrigMemoryOrder memorder, 1631 hsa_op_base *arg0, hsa_op_base *arg1, 1632 hsa_op_base *arg2, hsa_op_base *arg3) 1633 : hsa_insn_basic (nops, opc, t, arg0, arg1, arg2, arg3), 1634 m_memory_order (memorder), m_signalop (sop) 1635 { 1636 } 1637 1638 /* Constructor of class representing segment conversion instructions. OPC is 1639 the opcode which must be either BRIG_OPCODE_STOF or BRIG_OPCODE_FTOS. DEST 1640 and SRCT are destination and source types respectively, SEG is the segment 1641 we are converting to or from. The instruction operands are 1642 provided as ARG0 and ARG1. */ 1643 1644 hsa_insn_seg::hsa_insn_seg (int opc, BrigType16_t dest, BrigType16_t srct, 1645 BrigSegment8_t seg, hsa_op_base *arg0, 1646 hsa_op_base *arg1) 1647 : hsa_insn_basic (2, opc, dest, arg0, arg1), m_src_type (srct), 1648 m_segment (seg) 1649 { 1650 gcc_checking_assert (opc == BRIG_OPCODE_STOF || opc == BRIG_OPCODE_FTOS); 1651 } 1652 1653 /* Constructor of class representing a call instruction. CALLEE is the tree 1654 representation of the function being called. */ 1655 1656 hsa_insn_call::hsa_insn_call (tree callee) 1657 : hsa_insn_basic (0, BRIG_OPCODE_CALL), m_called_function (callee), 1658 m_output_arg (NULL), m_args_code_list (NULL), m_result_code_list (NULL) 1659 { 1660 } 1661 1662 hsa_insn_call::hsa_insn_call (hsa_internal_fn *fn) 1663 : hsa_insn_basic (0, BRIG_OPCODE_CALL), m_called_function (NULL), 1664 m_called_internal_fn (fn), m_output_arg (NULL), m_args_code_list (NULL), 1665 m_result_code_list (NULL) 1666 { 1667 } 1668 1669 hsa_insn_call::~hsa_insn_call () 1670 { 1671 for (unsigned i = 0; i < m_input_args.length (); i++) 1672 delete m_input_args[i]; 1673 1674 delete m_output_arg; 1675 1676 m_input_args.release (); 1677 m_input_arg_insns.release (); 1678 } 1679 1680 /* Constructor of class representing the argument block required to invoke 1681 a call in HSAIL. */ 1682 hsa_insn_arg_block::hsa_insn_arg_block (BrigKind brig_kind, 1683 hsa_insn_call * call) 1684 : hsa_insn_basic (0, HSA_OPCODE_ARG_BLOCK), m_kind (brig_kind), 1685 m_call_insn (call) 1686 { 1687 } 1688 1689 hsa_insn_comment::hsa_insn_comment (const char *s) 1690 : hsa_insn_basic (0, BRIG_KIND_DIRECTIVE_COMMENT) 1691 { 1692 unsigned l = strlen (s); 1693 1694 /* Append '// ' to the string. */ 1695 char *buf = XNEWVEC (char, l + 4); 1696 sprintf (buf, "// %s", s); 1697 m_comment = buf; 1698 } 1699 1700 hsa_insn_comment::~hsa_insn_comment () 1701 { 1702 gcc_checking_assert (m_comment); 1703 free (m_comment); 1704 m_comment = NULL; 1705 } 1706 1707 /* Constructor of class representing the queue instruction in HSAIL. */ 1708 1709 hsa_insn_queue::hsa_insn_queue (int nops, int opcode, BrigSegment segment, 1710 BrigMemoryOrder memory_order, 1711 hsa_op_base *arg0, hsa_op_base *arg1, 1712 hsa_op_base *arg2, hsa_op_base *arg3) 1713 : hsa_insn_basic (nops, opcode, BRIG_TYPE_U64, arg0, arg1, arg2, arg3), 1714 m_segment (segment), m_memory_order (memory_order) 1715 { 1716 } 1717 1718 /* Constructor of class representing the source type instruction in HSAIL. */ 1719 1720 hsa_insn_srctype::hsa_insn_srctype (int nops, BrigOpcode opcode, 1721 BrigType16_t destt, BrigType16_t srct, 1722 hsa_op_base *arg0, hsa_op_base *arg1, 1723 hsa_op_base *arg2 = NULL) 1724 : hsa_insn_basic (nops, opcode, destt, arg0, arg1, arg2), 1725 m_source_type (srct) 1726 {} 1727 1728 /* Constructor of class representing the packed instruction in HSAIL. */ 1729 1730 hsa_insn_packed::hsa_insn_packed (int nops, BrigOpcode opcode, 1731 BrigType16_t destt, BrigType16_t srct, 1732 hsa_op_base *arg0, hsa_op_base *arg1, 1733 hsa_op_base *arg2) 1734 : hsa_insn_srctype (nops, opcode, destt, srct, arg0, arg1, arg2) 1735 { 1736 m_operand_list = new hsa_op_operand_list (nops - 1); 1737 } 1738 1739 /* Constructor of class representing the convert instruction in HSAIL. */ 1740 1741 hsa_insn_cvt::hsa_insn_cvt (hsa_op_with_type *dest, hsa_op_with_type *src) 1742 : hsa_insn_basic (2, BRIG_OPCODE_CVT, dest->m_type, dest, src) 1743 { 1744 } 1745 1746 /* Constructor of class representing the alloca in HSAIL. */ 1747 1748 hsa_insn_alloca::hsa_insn_alloca (hsa_op_with_type *dest, 1749 hsa_op_with_type *size, unsigned alignment) 1750 : hsa_insn_basic (2, BRIG_OPCODE_ALLOCA, dest->m_type, dest, size), 1751 m_align (BRIG_ALIGNMENT_8) 1752 { 1753 gcc_assert (dest->m_type == BRIG_TYPE_U32); 1754 if (alignment) 1755 m_align = hsa_alignment_encoding (alignment); 1756 } 1757 1758 /* Append an instruction INSN into the basic block. */ 1759 1760 void 1761 hsa_bb::append_insn (hsa_insn_basic *insn) 1762 { 1763 gcc_assert (insn->m_opcode != 0 || insn->operand_count () == 0); 1764 gcc_assert (!insn->m_bb); 1765 1766 insn->m_bb = m_bb; 1767 insn->m_prev = m_last_insn; 1768 insn->m_next = NULL; 1769 if (m_last_insn) 1770 m_last_insn->m_next = insn; 1771 m_last_insn = insn; 1772 if (!m_first_insn) 1773 m_first_insn = insn; 1774 } 1775 1776 void 1777 hsa_bb::append_phi (hsa_insn_phi *hphi) 1778 { 1779 hphi->m_bb = m_bb; 1780 1781 hphi->m_prev = m_last_phi; 1782 hphi->m_next = NULL; 1783 if (m_last_phi) 1784 m_last_phi->m_next = hphi; 1785 m_last_phi = hphi; 1786 if (!m_first_phi) 1787 m_first_phi = hphi; 1788 } 1789 1790 /* Insert HSA instruction NEW_INSN immediately before an existing instruction 1791 OLD_INSN. */ 1792 1793 static void 1794 hsa_insert_insn_before (hsa_insn_basic *new_insn, hsa_insn_basic *old_insn) 1795 { 1796 hsa_bb *hbb = hsa_bb_for_bb (old_insn->m_bb); 1797 1798 if (hbb->m_first_insn == old_insn) 1799 hbb->m_first_insn = new_insn; 1800 new_insn->m_prev = old_insn->m_prev; 1801 new_insn->m_next = old_insn; 1802 if (old_insn->m_prev) 1803 old_insn->m_prev->m_next = new_insn; 1804 old_insn->m_prev = new_insn; 1805 } 1806 1807 /* Append HSA instruction NEW_INSN immediately after an existing instruction 1808 OLD_INSN. */ 1809 1810 static void 1811 hsa_append_insn_after (hsa_insn_basic *new_insn, hsa_insn_basic *old_insn) 1812 { 1813 hsa_bb *hbb = hsa_bb_for_bb (old_insn->m_bb); 1814 1815 if (hbb->m_last_insn == old_insn) 1816 hbb->m_last_insn = new_insn; 1817 new_insn->m_prev = old_insn; 1818 new_insn->m_next = old_insn->m_next; 1819 if (old_insn->m_next) 1820 old_insn->m_next->m_prev = new_insn; 1821 old_insn->m_next = new_insn; 1822 } 1823 1824 /* Return a register containing the calculated value of EXP which must be an 1825 expression consisting of PLUS_EXPRs, MULT_EXPRs, NOP_EXPRs, SSA_NAMEs and 1826 integer constants as returned by get_inner_reference. 1827 Newly generated HSA instructions will be appended to HBB. 1828 Perform all calculations in ADDRTYPE. */ 1829 1830 static hsa_op_with_type * 1831 gen_address_calculation (tree exp, hsa_bb *hbb, BrigType16_t addrtype) 1832 { 1833 int opcode; 1834 1835 if (TREE_CODE (exp) == NOP_EXPR) 1836 exp = TREE_OPERAND (exp, 0); 1837 1838 switch (TREE_CODE (exp)) 1839 { 1840 case SSA_NAME: 1841 return hsa_cfun->reg_for_gimple_ssa (exp)->get_in_type (addrtype, hbb); 1842 1843 case INTEGER_CST: 1844 { 1845 hsa_op_immed *imm = new hsa_op_immed (exp); 1846 if (addrtype != imm->m_type) 1847 imm->m_type = addrtype; 1848 return imm; 1849 } 1850 1851 case PLUS_EXPR: 1852 opcode = BRIG_OPCODE_ADD; 1853 break; 1854 1855 case MULT_EXPR: 1856 opcode = BRIG_OPCODE_MUL; 1857 break; 1858 1859 default: 1860 gcc_unreachable (); 1861 } 1862 1863 hsa_op_reg *res = new hsa_op_reg (addrtype); 1864 hsa_insn_basic *insn = new hsa_insn_basic (3, opcode, addrtype); 1865 insn->set_op (0, res); 1866 1867 hsa_op_with_type *op1 = gen_address_calculation (TREE_OPERAND (exp, 0), hbb, 1868 addrtype); 1869 hsa_op_with_type *op2 = gen_address_calculation (TREE_OPERAND (exp, 1), hbb, 1870 addrtype); 1871 insn->set_op (1, op1); 1872 insn->set_op (2, op2); 1873 1874 hbb->append_insn (insn); 1875 return res; 1876 } 1877 1878 /* If R1 is NULL, just return R2, otherwise append an instruction adding them 1879 to HBB and return the register holding the result. */ 1880 1881 static hsa_op_reg * 1882 add_addr_regs_if_needed (hsa_op_reg *r1, hsa_op_reg *r2, hsa_bb *hbb) 1883 { 1884 gcc_checking_assert (r2); 1885 if (!r1) 1886 return r2; 1887 1888 hsa_op_reg *res = new hsa_op_reg (r1->m_type); 1889 gcc_assert (!hsa_needs_cvt (r1->m_type, r2->m_type)); 1890 hsa_insn_basic *insn = new hsa_insn_basic (3, BRIG_OPCODE_ADD, res->m_type); 1891 insn->set_op (0, res); 1892 insn->set_op (1, r1); 1893 insn->set_op (2, r2); 1894 hbb->append_insn (insn); 1895 return res; 1896 } 1897 1898 /* Helper of gen_hsa_addr. Update *SYMBOL, *ADDRTYPE, *REG and *OFFSET to 1899 reflect BASE which is the first operand of a MEM_REF or a TARGET_MEM_REF. */ 1900 1901 static void 1902 process_mem_base (tree base, hsa_symbol **symbol, BrigType16_t *addrtype, 1903 hsa_op_reg **reg, offset_int *offset, hsa_bb *hbb) 1904 { 1905 if (TREE_CODE (base) == SSA_NAME) 1906 { 1907 gcc_assert (!*reg); 1908 hsa_op_with_type *ssa 1909 = hsa_cfun->reg_for_gimple_ssa (base)->get_in_type (*addrtype, hbb); 1910 *reg = dyn_cast <hsa_op_reg *> (ssa); 1911 } 1912 else if (TREE_CODE (base) == ADDR_EXPR) 1913 { 1914 tree decl = TREE_OPERAND (base, 0); 1915 1916 if (!DECL_P (decl) || TREE_CODE (decl) == FUNCTION_DECL) 1917 { 1918 HSA_SORRY_AT (EXPR_LOCATION (base), 1919 "support for HSA does not implement a memory reference " 1920 "to a non-declaration type"); 1921 return; 1922 } 1923 1924 gcc_assert (!*symbol); 1925 1926 *symbol = get_symbol_for_decl (decl); 1927 *addrtype = hsa_get_segment_addr_type ((*symbol)->m_segment); 1928 } 1929 else if (TREE_CODE (base) == INTEGER_CST) 1930 *offset += wi::to_offset (base); 1931 else 1932 gcc_unreachable (); 1933 } 1934 1935 /* Forward declaration of a function. */ 1936 1937 static void 1938 gen_hsa_addr_insns (tree val, hsa_op_reg *dest, hsa_bb *hbb); 1939 1940 /* Generate HSA address operand for a given tree memory reference REF. If 1941 instructions need to be created to calculate the address, they will be added 1942 to the end of HBB. If a caller provider OUTPUT_BITSIZE and OUTPUT_BITPOS, 1943 the function assumes that the caller will handle possible 1944 bit-field references. Otherwise if we reference a bit-field, sorry message 1945 is displayed. */ 1946 1947 static hsa_op_address * 1948 gen_hsa_addr (tree ref, hsa_bb *hbb, HOST_WIDE_INT *output_bitsize = NULL, 1949 HOST_WIDE_INT *output_bitpos = NULL) 1950 { 1951 hsa_symbol *symbol = NULL; 1952 hsa_op_reg *reg = NULL; 1953 offset_int offset = 0; 1954 tree origref = ref; 1955 tree varoffset = NULL_TREE; 1956 BrigType16_t addrtype = hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT); 1957 HOST_WIDE_INT bitsize = 0, bitpos = 0; 1958 BrigType16_t flat_addrtype = hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT); 1959 1960 if (TREE_CODE (ref) == STRING_CST) 1961 { 1962 symbol = hsa_get_string_cst_symbol (ref); 1963 goto out; 1964 } 1965 else if (TREE_CODE (ref) == BIT_FIELD_REF 1966 && (!multiple_p (bit_field_size (ref), BITS_PER_UNIT) 1967 || !multiple_p (bit_field_offset (ref), BITS_PER_UNIT))) 1968 { 1969 HSA_SORRY_ATV (EXPR_LOCATION (origref), 1970 "support for HSA does not implement " 1971 "bit field references such as %E", ref); 1972 goto out; 1973 } 1974 1975 if (handled_component_p (ref)) 1976 { 1977 machine_mode mode; 1978 int unsignedp, volatilep, preversep; 1979 poly_int64 pbitsize, pbitpos; 1980 tree new_ref; 1981 1982 new_ref = get_inner_reference (ref, &pbitsize, &pbitpos, &varoffset, 1983 &mode, &unsignedp, &preversep, 1984 &volatilep); 1985 /* When this isn't true, the switch below will report an 1986 appropriate error. */ 1987 if (pbitsize.is_constant () && pbitpos.is_constant ()) 1988 { 1989 bitsize = pbitsize.to_constant (); 1990 bitpos = pbitpos.to_constant (); 1991 ref = new_ref; 1992 offset = bitpos; 1993 offset = wi::rshift (offset, LOG2_BITS_PER_UNIT, SIGNED); 1994 } 1995 } 1996 1997 switch (TREE_CODE (ref)) 1998 { 1999 case ADDR_EXPR: 2000 { 2001 addrtype = hsa_get_segment_addr_type (BRIG_SEGMENT_PRIVATE); 2002 symbol = hsa_cfun->create_hsa_temporary (flat_addrtype); 2003 hsa_op_reg *r = new hsa_op_reg (flat_addrtype); 2004 gen_hsa_addr_insns (ref, r, hbb); 2005 hbb->append_insn (new hsa_insn_mem (BRIG_OPCODE_ST, r->m_type, 2006 r, new hsa_op_address (symbol))); 2007 2008 break; 2009 } 2010 case SSA_NAME: 2011 { 2012 addrtype = hsa_get_segment_addr_type (BRIG_SEGMENT_PRIVATE); 2013 hsa_op_with_type *r = hsa_cfun->reg_for_gimple_ssa (ref); 2014 if (r->m_type == BRIG_TYPE_B1) 2015 r = r->get_in_type (BRIG_TYPE_U32, hbb); 2016 symbol = hsa_cfun->create_hsa_temporary (r->m_type); 2017 2018 hbb->append_insn (new hsa_insn_mem (BRIG_OPCODE_ST, r->m_type, 2019 r, new hsa_op_address (symbol))); 2020 2021 break; 2022 } 2023 case PARM_DECL: 2024 case VAR_DECL: 2025 case RESULT_DECL: 2026 case CONST_DECL: 2027 gcc_assert (!symbol); 2028 symbol = get_symbol_for_decl (ref); 2029 addrtype = hsa_get_segment_addr_type (symbol->m_segment); 2030 break; 2031 2032 case MEM_REF: 2033 process_mem_base (TREE_OPERAND (ref, 0), &symbol, &addrtype, ®, 2034 &offset, hbb); 2035 2036 if (!integer_zerop (TREE_OPERAND (ref, 1))) 2037 offset += wi::to_offset (TREE_OPERAND (ref, 1)); 2038 break; 2039 2040 case TARGET_MEM_REF: 2041 process_mem_base (TMR_BASE (ref), &symbol, &addrtype, ®, &offset, hbb); 2042 if (TMR_INDEX (ref)) 2043 { 2044 hsa_op_reg *disp1; 2045 hsa_op_base *idx = hsa_cfun->reg_for_gimple_ssa 2046 (TMR_INDEX (ref))->get_in_type (addrtype, hbb); 2047 if (TMR_STEP (ref) && !integer_onep (TMR_STEP (ref))) 2048 { 2049 disp1 = new hsa_op_reg (addrtype); 2050 hsa_insn_basic *insn = new hsa_insn_basic (3, BRIG_OPCODE_MUL, 2051 addrtype); 2052 2053 /* As step must respect addrtype, we overwrite the type 2054 of an immediate value. */ 2055 hsa_op_immed *step = new hsa_op_immed (TMR_STEP (ref)); 2056 step->m_type = addrtype; 2057 2058 insn->set_op (0, disp1); 2059 insn->set_op (1, idx); 2060 insn->set_op (2, step); 2061 hbb->append_insn (insn); 2062 } 2063 else 2064 disp1 = as_a <hsa_op_reg *> (idx); 2065 reg = add_addr_regs_if_needed (reg, disp1, hbb); 2066 } 2067 if (TMR_INDEX2 (ref)) 2068 { 2069 if (TREE_CODE (TMR_INDEX2 (ref)) == SSA_NAME) 2070 { 2071 hsa_op_base *disp2 = hsa_cfun->reg_for_gimple_ssa 2072 (TMR_INDEX2 (ref))->get_in_type (addrtype, hbb); 2073 reg = add_addr_regs_if_needed (reg, as_a <hsa_op_reg *> (disp2), 2074 hbb); 2075 } 2076 else if (TREE_CODE (TMR_INDEX2 (ref)) == INTEGER_CST) 2077 offset += wi::to_offset (TMR_INDEX2 (ref)); 2078 else 2079 gcc_unreachable (); 2080 } 2081 offset += wi::to_offset (TMR_OFFSET (ref)); 2082 break; 2083 case FUNCTION_DECL: 2084 HSA_SORRY_AT (EXPR_LOCATION (origref), 2085 "support for HSA does not implement function pointers"); 2086 goto out; 2087 default: 2088 HSA_SORRY_ATV (EXPR_LOCATION (origref), "support for HSA does " 2089 "not implement memory access to %E", origref); 2090 goto out; 2091 } 2092 2093 if (varoffset) 2094 { 2095 if (TREE_CODE (varoffset) == INTEGER_CST) 2096 offset += wi::to_offset (varoffset); 2097 else 2098 { 2099 hsa_op_base *off_op = gen_address_calculation (varoffset, hbb, 2100 addrtype); 2101 reg = add_addr_regs_if_needed (reg, as_a <hsa_op_reg *> (off_op), 2102 hbb); 2103 } 2104 } 2105 2106 gcc_checking_assert ((symbol 2107 && addrtype 2108 == hsa_get_segment_addr_type (symbol->m_segment)) 2109 || (!symbol 2110 && addrtype 2111 == hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT))); 2112 out: 2113 HOST_WIDE_INT hwi_offset = offset.to_shwi (); 2114 2115 /* Calculate remaining bitsize offset (if presented). */ 2116 bitpos %= BITS_PER_UNIT; 2117 /* If bitsize is a power of two that is greater or equal to BITS_PER_UNIT, it 2118 is not a reason to think this is a bit-field access. */ 2119 if (bitpos == 0 2120 && (bitsize >= BITS_PER_UNIT) 2121 && !(bitsize & (bitsize - 1))) 2122 bitsize = 0; 2123 2124 if ((bitpos || bitsize) && (output_bitpos == NULL || output_bitsize == NULL)) 2125 HSA_SORRY_ATV (EXPR_LOCATION (origref), "support for HSA does not " 2126 "implement unhandled bit field reference such as %E", ref); 2127 2128 if (output_bitsize != NULL && output_bitpos != NULL) 2129 { 2130 *output_bitsize = bitsize; 2131 *output_bitpos = bitpos; 2132 } 2133 2134 return new hsa_op_address (symbol, reg, hwi_offset); 2135 } 2136 2137 /* Generate HSA address operand for a given tree memory reference REF. If 2138 instructions need to be created to calculate the address, they will be added 2139 to the end of HBB. OUTPUT_ALIGN is alignment of the created address. */ 2140 2141 static hsa_op_address * 2142 gen_hsa_addr_with_align (tree ref, hsa_bb *hbb, BrigAlignment8_t *output_align) 2143 { 2144 hsa_op_address *addr = gen_hsa_addr (ref, hbb); 2145 if (addr->m_reg || !addr->m_symbol) 2146 *output_align = hsa_object_alignment (ref); 2147 else 2148 { 2149 /* If the address consists only of a symbol and an offset, we 2150 compute the alignment ourselves to take into account any alignment 2151 promotions we might have done for the HSA symbol representation. */ 2152 unsigned align = hsa_byte_alignment (addr->m_symbol->m_align); 2153 unsigned misalign = addr->m_imm_offset & (align - 1); 2154 if (misalign) 2155 align = least_bit_hwi (misalign); 2156 *output_align = hsa_alignment_encoding (BITS_PER_UNIT * align); 2157 } 2158 return addr; 2159 } 2160 2161 /* Generate HSA address for a function call argument of given TYPE. 2162 INDEX is used to generate corresponding name of the arguments. 2163 Special value -1 represents fact that result value is created. */ 2164 2165 static hsa_op_address * 2166 gen_hsa_addr_for_arg (tree tree_type, int index) 2167 { 2168 hsa_symbol *sym = new hsa_symbol (BRIG_TYPE_NONE, BRIG_SEGMENT_ARG, 2169 BRIG_LINKAGE_ARG); 2170 sym->m_type = hsa_type_for_tree_type (tree_type, &sym->m_dim); 2171 2172 if (index == -1) /* Function result. */ 2173 sym->m_name = "res"; 2174 else /* Function call arguments. */ 2175 { 2176 sym->m_name = NULL; 2177 sym->m_name_number = index; 2178 } 2179 2180 return new hsa_op_address (sym); 2181 } 2182 2183 /* Generate HSA instructions that process all necessary conversions 2184 of an ADDR to flat addressing and place the result into DEST. 2185 Instructions are appended to HBB. */ 2186 2187 static void 2188 convert_addr_to_flat_segment (hsa_op_address *addr, hsa_op_reg *dest, 2189 hsa_bb *hbb) 2190 { 2191 hsa_insn_basic *insn = new hsa_insn_basic (2, BRIG_OPCODE_LDA); 2192 insn->set_op (1, addr); 2193 if (addr->m_symbol && addr->m_symbol->m_segment != BRIG_SEGMENT_GLOBAL) 2194 { 2195 /* LDA produces segment-relative address, we need to convert 2196 it to the flat one. */ 2197 hsa_op_reg *tmp; 2198 tmp = new hsa_op_reg (hsa_get_segment_addr_type 2199 (addr->m_symbol->m_segment)); 2200 hsa_insn_seg *seg; 2201 seg = new hsa_insn_seg (BRIG_OPCODE_STOF, 2202 hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT), 2203 tmp->m_type, addr->m_symbol->m_segment, dest, 2204 tmp); 2205 2206 insn->set_op (0, tmp); 2207 insn->m_type = tmp->m_type; 2208 hbb->append_insn (insn); 2209 hbb->append_insn (seg); 2210 } 2211 else 2212 { 2213 insn->set_op (0, dest); 2214 insn->m_type = hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT); 2215 hbb->append_insn (insn); 2216 } 2217 } 2218 2219 /* Generate HSA instructions that calculate address of VAL including all 2220 necessary conversions to flat addressing and place the result into DEST. 2221 Instructions are appended to HBB. */ 2222 2223 static void 2224 gen_hsa_addr_insns (tree val, hsa_op_reg *dest, hsa_bb *hbb) 2225 { 2226 /* Handle cases like tmp = NULL, where we just emit a move instruction 2227 to a register. */ 2228 if (TREE_CODE (val) == INTEGER_CST) 2229 { 2230 hsa_op_immed *c = new hsa_op_immed (val); 2231 hsa_insn_basic *insn = new hsa_insn_basic (2, BRIG_OPCODE_MOV, 2232 dest->m_type, dest, c); 2233 hbb->append_insn (insn); 2234 return; 2235 } 2236 2237 hsa_op_address *addr; 2238 2239 gcc_assert (dest->m_type == hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT)); 2240 if (TREE_CODE (val) == ADDR_EXPR) 2241 val = TREE_OPERAND (val, 0); 2242 addr = gen_hsa_addr (val, hbb); 2243 2244 if (TREE_CODE (val) == CONST_DECL 2245 && is_gimple_reg_type (TREE_TYPE (val))) 2246 { 2247 gcc_assert (addr->m_symbol 2248 && addr->m_symbol->m_segment == BRIG_SEGMENT_READONLY); 2249 /* CONST_DECLs are in readonly segment which however does not have 2250 addresses convertible to flat segments. So copy it to a private one 2251 and take address of that. */ 2252 BrigType16_t csttype 2253 = mem_type_for_type (hsa_type_for_scalar_tree_type (TREE_TYPE (val), 2254 false)); 2255 hsa_op_reg *r = new hsa_op_reg (csttype); 2256 hbb->append_insn (new hsa_insn_mem (BRIG_OPCODE_LD, csttype, r, 2257 new hsa_op_address (addr->m_symbol))); 2258 hsa_symbol *copysym = hsa_cfun->create_hsa_temporary (csttype); 2259 hbb->append_insn (new hsa_insn_mem (BRIG_OPCODE_ST, csttype, r, 2260 new hsa_op_address (copysym))); 2261 addr->m_symbol = copysym; 2262 } 2263 else if (addr->m_symbol && addr->m_symbol->m_segment == BRIG_SEGMENT_READONLY) 2264 { 2265 HSA_SORRY_ATV (EXPR_LOCATION (val), "support for HSA does " 2266 "not implement taking addresses of complex " 2267 "CONST_DECLs such as %E", val); 2268 return; 2269 } 2270 2271 2272 convert_addr_to_flat_segment (addr, dest, hbb); 2273 } 2274 2275 /* Return an HSA register or HSA immediate value operand corresponding to 2276 gimple operand OP. */ 2277 2278 static hsa_op_with_type * 2279 hsa_reg_or_immed_for_gimple_op (tree op, hsa_bb *hbb) 2280 { 2281 hsa_op_reg *tmp; 2282 2283 if (TREE_CODE (op) == SSA_NAME) 2284 tmp = hsa_cfun->reg_for_gimple_ssa (op); 2285 else if (!POINTER_TYPE_P (TREE_TYPE (op))) 2286 return new hsa_op_immed (op); 2287 else 2288 { 2289 tmp = new hsa_op_reg (hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT)); 2290 gen_hsa_addr_insns (op, tmp, hbb); 2291 } 2292 return tmp; 2293 } 2294 2295 /* Create a simple movement instruction with register destination DEST and 2296 register or immediate source SRC and append it to the end of HBB. */ 2297 2298 void 2299 hsa_build_append_simple_mov (hsa_op_reg *dest, hsa_op_base *src, hsa_bb *hbb) 2300 { 2301 /* Moves of packed data between registers need to adhere to the same type 2302 rules like when dealing with memory. */ 2303 BrigType16_t tp = mem_type_for_type (dest->m_type); 2304 hsa_insn_basic *insn = new hsa_insn_basic (2, BRIG_OPCODE_MOV, tp, dest, src); 2305 hsa_fixup_mov_insn_type (insn); 2306 unsigned dest_size = hsa_type_bit_size (dest->m_type); 2307 if (hsa_op_reg *sreg = dyn_cast <hsa_op_reg *> (src)) 2308 gcc_assert (dest_size == hsa_type_bit_size (sreg->m_type)); 2309 else 2310 { 2311 unsigned imm_size 2312 = hsa_type_bit_size (as_a <hsa_op_immed *> (src)->m_type); 2313 gcc_assert ((dest_size == imm_size) 2314 /* Eventually < 32bit registers will be promoted to 32bit. */ 2315 || (dest_size < 32 && imm_size == 32)); 2316 } 2317 hbb->append_insn (insn); 2318 } 2319 2320 /* Generate HSAIL instructions loading a bit field into register DEST. 2321 VALUE_REG is a register of a SSA name that is used in the bit field 2322 reference. To identify a bit field BITPOS is offset to the loaded memory 2323 and BITSIZE is number of bits of the bit field. 2324 Add instructions to HBB. */ 2325 2326 static void 2327 gen_hsa_insns_for_bitfield (hsa_op_reg *dest, hsa_op_reg *value_reg, 2328 HOST_WIDE_INT bitsize, HOST_WIDE_INT bitpos, 2329 hsa_bb *hbb) 2330 { 2331 unsigned type_bitsize 2332 = hsa_type_bit_size (hsa_extend_inttype_to_32bit (dest->m_type)); 2333 unsigned left_shift = type_bitsize - (bitsize + bitpos); 2334 unsigned right_shift = left_shift + bitpos; 2335 2336 if (left_shift) 2337 { 2338 hsa_op_reg *value_reg_2 2339 = new hsa_op_reg (hsa_extend_inttype_to_32bit (dest->m_type)); 2340 hsa_op_immed *c = new hsa_op_immed (left_shift, BRIG_TYPE_U32); 2341 2342 hsa_insn_basic *lshift 2343 = new hsa_insn_basic (3, BRIG_OPCODE_SHL, value_reg_2->m_type, 2344 value_reg_2, value_reg, c); 2345 2346 hbb->append_insn (lshift); 2347 2348 value_reg = value_reg_2; 2349 } 2350 2351 if (right_shift) 2352 { 2353 hsa_op_reg *value_reg_2 2354 = new hsa_op_reg (hsa_extend_inttype_to_32bit (dest->m_type)); 2355 hsa_op_immed *c = new hsa_op_immed (right_shift, BRIG_TYPE_U32); 2356 2357 hsa_insn_basic *rshift 2358 = new hsa_insn_basic (3, BRIG_OPCODE_SHR, value_reg_2->m_type, 2359 value_reg_2, value_reg, c); 2360 2361 hbb->append_insn (rshift); 2362 2363 value_reg = value_reg_2; 2364 } 2365 2366 hsa_insn_basic *assignment 2367 = new hsa_insn_basic (2, BRIG_OPCODE_MOV, dest->m_type, NULL, value_reg); 2368 hsa_fixup_mov_insn_type (assignment); 2369 hbb->append_insn (assignment); 2370 assignment->set_output_in_type (dest, 0, hbb); 2371 } 2372 2373 2374 /* Generate HSAIL instructions loading a bit field into register DEST. ADDR is 2375 prepared memory address which is used to load the bit field. To identify a 2376 bit field BITPOS is offset to the loaded memory and BITSIZE is number of 2377 bits of the bit field. Add instructions to HBB. Load must be performed in 2378 alignment ALIGN. */ 2379 2380 static void 2381 gen_hsa_insns_for_bitfield_load (hsa_op_reg *dest, hsa_op_address *addr, 2382 HOST_WIDE_INT bitsize, HOST_WIDE_INT bitpos, 2383 hsa_bb *hbb, BrigAlignment8_t align) 2384 { 2385 hsa_op_reg *value_reg = new hsa_op_reg (dest->m_type); 2386 hsa_insn_mem *mem 2387 = new hsa_insn_mem (BRIG_OPCODE_LD, 2388 hsa_extend_inttype_to_32bit (dest->m_type), 2389 value_reg, addr); 2390 mem->set_align (align); 2391 hbb->append_insn (mem); 2392 gen_hsa_insns_for_bitfield (dest, value_reg, bitsize, bitpos, hbb); 2393 } 2394 2395 /* Return the alignment of base memory accesses we issue to perform bit-field 2396 memory access REF. */ 2397 2398 static BrigAlignment8_t 2399 hsa_bitmemref_alignment (tree ref) 2400 { 2401 unsigned HOST_WIDE_INT bit_offset = 0; 2402 2403 while (true) 2404 { 2405 if (TREE_CODE (ref) == BIT_FIELD_REF) 2406 { 2407 if (!tree_fits_uhwi_p (TREE_OPERAND (ref, 2))) 2408 return BRIG_ALIGNMENT_1; 2409 bit_offset += tree_to_uhwi (TREE_OPERAND (ref, 2)); 2410 } 2411 else if (TREE_CODE (ref) == COMPONENT_REF 2412 && DECL_BIT_FIELD (TREE_OPERAND (ref, 1))) 2413 bit_offset += int_bit_position (TREE_OPERAND (ref, 1)); 2414 else 2415 break; 2416 ref = TREE_OPERAND (ref, 0); 2417 } 2418 2419 unsigned HOST_WIDE_INT bits = bit_offset % BITS_PER_UNIT; 2420 unsigned HOST_WIDE_INT byte_bits = bit_offset - bits; 2421 BrigAlignment8_t base = hsa_object_alignment (ref); 2422 if (byte_bits == 0) 2423 return base; 2424 return MIN (base, hsa_alignment_encoding (least_bit_hwi (byte_bits))); 2425 } 2426 2427 /* Generate HSAIL instructions loading something into register DEST. RHS is 2428 tree representation of the loaded data, which are loaded as type TYPE. Add 2429 instructions to HBB. */ 2430 2431 static void 2432 gen_hsa_insns_for_load (hsa_op_reg *dest, tree rhs, tree type, hsa_bb *hbb) 2433 { 2434 /* The destination SSA name will give us the type. */ 2435 if (TREE_CODE (rhs) == VIEW_CONVERT_EXPR) 2436 rhs = TREE_OPERAND (rhs, 0); 2437 2438 if (TREE_CODE (rhs) == SSA_NAME) 2439 { 2440 hsa_op_reg *src = hsa_cfun->reg_for_gimple_ssa (rhs); 2441 hsa_build_append_simple_mov (dest, src, hbb); 2442 } 2443 else if (is_gimple_min_invariant (rhs) 2444 || TREE_CODE (rhs) == ADDR_EXPR) 2445 { 2446 if (POINTER_TYPE_P (TREE_TYPE (rhs))) 2447 { 2448 if (dest->m_type != hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT)) 2449 { 2450 HSA_SORRY_ATV (EXPR_LOCATION (rhs), 2451 "support for HSA does not implement conversion " 2452 "of %E to the requested non-pointer type.", rhs); 2453 return; 2454 } 2455 2456 gen_hsa_addr_insns (rhs, dest, hbb); 2457 } 2458 else if (TREE_CODE (rhs) == COMPLEX_CST) 2459 { 2460 hsa_op_immed *real_part = new hsa_op_immed (TREE_REALPART (rhs)); 2461 hsa_op_immed *imag_part = new hsa_op_immed (TREE_IMAGPART (rhs)); 2462 2463 hsa_op_reg *real_part_reg 2464 = new hsa_op_reg (hsa_type_for_scalar_tree_type (TREE_TYPE (type), 2465 true)); 2466 hsa_op_reg *imag_part_reg 2467 = new hsa_op_reg (hsa_type_for_scalar_tree_type (TREE_TYPE (type), 2468 true)); 2469 2470 hsa_build_append_simple_mov (real_part_reg, real_part, hbb); 2471 hsa_build_append_simple_mov (imag_part_reg, imag_part, hbb); 2472 2473 BrigType16_t src_type = hsa_bittype_for_type (real_part_reg->m_type); 2474 2475 hsa_insn_packed *insn 2476 = new hsa_insn_packed (3, BRIG_OPCODE_COMBINE, dest->m_type, 2477 src_type, dest, real_part_reg, 2478 imag_part_reg); 2479 hbb->append_insn (insn); 2480 } 2481 else 2482 { 2483 hsa_op_immed *imm = new hsa_op_immed (rhs); 2484 hsa_build_append_simple_mov (dest, imm, hbb); 2485 } 2486 } 2487 else if (TREE_CODE (rhs) == REALPART_EXPR || TREE_CODE (rhs) == IMAGPART_EXPR) 2488 { 2489 tree pack_type = TREE_TYPE (TREE_OPERAND (rhs, 0)); 2490 2491 hsa_op_reg *packed_reg 2492 = new hsa_op_reg (hsa_type_for_scalar_tree_type (pack_type, true)); 2493 2494 tree complex_rhs = TREE_OPERAND (rhs, 0); 2495 gen_hsa_insns_for_load (packed_reg, complex_rhs, TREE_TYPE (complex_rhs), 2496 hbb); 2497 2498 hsa_op_reg *real_reg 2499 = new hsa_op_reg (hsa_type_for_scalar_tree_type (type, true)); 2500 2501 hsa_op_reg *imag_reg 2502 = new hsa_op_reg (hsa_type_for_scalar_tree_type (type, true)); 2503 2504 BrigKind16_t brig_type = packed_reg->m_type; 2505 hsa_insn_packed *packed 2506 = new hsa_insn_packed (3, BRIG_OPCODE_EXPAND, 2507 hsa_bittype_for_type (real_reg->m_type), 2508 brig_type, real_reg, imag_reg, packed_reg); 2509 2510 hbb->append_insn (packed); 2511 2512 hsa_op_reg *source = TREE_CODE (rhs) == REALPART_EXPR ? 2513 real_reg : imag_reg; 2514 2515 hsa_insn_basic *insn = new hsa_insn_basic (2, BRIG_OPCODE_MOV, 2516 dest->m_type, NULL, source); 2517 hsa_fixup_mov_insn_type (insn); 2518 hbb->append_insn (insn); 2519 insn->set_output_in_type (dest, 0, hbb); 2520 } 2521 else if (TREE_CODE (rhs) == BIT_FIELD_REF 2522 && TREE_CODE (TREE_OPERAND (rhs, 0)) == SSA_NAME) 2523 { 2524 tree ssa_name = TREE_OPERAND (rhs, 0); 2525 HOST_WIDE_INT bitsize = tree_to_uhwi (TREE_OPERAND (rhs, 1)); 2526 HOST_WIDE_INT bitpos = tree_to_uhwi (TREE_OPERAND (rhs, 2)); 2527 2528 hsa_op_reg *imm_value = hsa_cfun->reg_for_gimple_ssa (ssa_name); 2529 gen_hsa_insns_for_bitfield (dest, imm_value, bitsize, bitpos, hbb); 2530 } 2531 else if (DECL_P (rhs) || TREE_CODE (rhs) == MEM_REF 2532 || TREE_CODE (rhs) == TARGET_MEM_REF 2533 || handled_component_p (rhs)) 2534 { 2535 HOST_WIDE_INT bitsize, bitpos; 2536 2537 /* Load from memory. */ 2538 hsa_op_address *addr; 2539 addr = gen_hsa_addr (rhs, hbb, &bitsize, &bitpos); 2540 2541 /* Handle load of a bit field. */ 2542 if (bitsize > 64) 2543 { 2544 HSA_SORRY_AT (EXPR_LOCATION (rhs), 2545 "support for HSA does not implement load from a bit " 2546 "field bigger than 64 bits"); 2547 return; 2548 } 2549 2550 if (bitsize || bitpos) 2551 gen_hsa_insns_for_bitfield_load (dest, addr, bitsize, bitpos, hbb, 2552 hsa_bitmemref_alignment (rhs)); 2553 else 2554 { 2555 BrigType16_t mtype; 2556 /* Not dest->m_type, that's possibly extended. */ 2557 mtype = mem_type_for_type (hsa_type_for_scalar_tree_type (type, 2558 false)); 2559 hsa_insn_mem *mem = new hsa_insn_mem (BRIG_OPCODE_LD, mtype, dest, 2560 addr); 2561 mem->set_align (hsa_object_alignment (rhs)); 2562 hbb->append_insn (mem); 2563 } 2564 } 2565 else 2566 HSA_SORRY_ATV (EXPR_LOCATION (rhs), 2567 "support for HSA does not implement loading " 2568 "of expression %E", 2569 rhs); 2570 } 2571 2572 /* Return number of bits necessary for representation of a bit field, 2573 starting at BITPOS with size of BITSIZE. */ 2574 2575 static unsigned 2576 get_bitfield_size (unsigned bitpos, unsigned bitsize) 2577 { 2578 unsigned s = bitpos + bitsize; 2579 unsigned sizes[] = {8, 16, 32, 64}; 2580 2581 for (unsigned i = 0; i < 4; i++) 2582 if (s <= sizes[i]) 2583 return sizes[i]; 2584 2585 gcc_unreachable (); 2586 return 0; 2587 } 2588 2589 /* Generate HSAIL instructions storing into memory. LHS is the destination of 2590 the store, SRC is the source operand. Add instructions to HBB. */ 2591 2592 static void 2593 gen_hsa_insns_for_store (tree lhs, hsa_op_base *src, hsa_bb *hbb) 2594 { 2595 HOST_WIDE_INT bitsize = 0, bitpos = 0; 2596 BrigAlignment8_t req_align; 2597 BrigType16_t mtype; 2598 mtype = mem_type_for_type (hsa_type_for_scalar_tree_type (TREE_TYPE (lhs), 2599 false)); 2600 hsa_op_address *addr; 2601 addr = gen_hsa_addr (lhs, hbb, &bitsize, &bitpos); 2602 2603 /* Handle store to a bit field. */ 2604 if (bitsize > 64) 2605 { 2606 HSA_SORRY_AT (EXPR_LOCATION (lhs), 2607 "support for HSA does not implement store to a bit field " 2608 "bigger than 64 bits"); 2609 return; 2610 } 2611 2612 unsigned type_bitsize = get_bitfield_size (bitpos, bitsize); 2613 2614 /* HSAIL does not support MOV insn with 16-bits integers. */ 2615 if (type_bitsize < 32) 2616 type_bitsize = 32; 2617 2618 if (bitpos || (bitsize && type_bitsize != bitsize)) 2619 { 2620 unsigned HOST_WIDE_INT mask = 0; 2621 BrigType16_t mem_type 2622 = get_integer_type_by_bytes (type_bitsize / BITS_PER_UNIT, 2623 !TYPE_UNSIGNED (TREE_TYPE (lhs))); 2624 2625 for (unsigned i = 0; i < type_bitsize; i++) 2626 if (i < bitpos || i >= bitpos + bitsize) 2627 mask |= ((unsigned HOST_WIDE_INT)1 << i); 2628 2629 hsa_op_reg *value_reg = new hsa_op_reg (mem_type); 2630 2631 req_align = hsa_bitmemref_alignment (lhs); 2632 /* Load value from memory. */ 2633 hsa_insn_mem *mem = new hsa_insn_mem (BRIG_OPCODE_LD, mem_type, 2634 value_reg, addr); 2635 mem->set_align (req_align); 2636 hbb->append_insn (mem); 2637 2638 /* AND the loaded value with prepared mask. */ 2639 hsa_op_reg *cleared_reg = new hsa_op_reg (mem_type); 2640 2641 BrigType16_t t 2642 = get_integer_type_by_bytes (type_bitsize / BITS_PER_UNIT, false); 2643 hsa_op_immed *c = new hsa_op_immed (mask, t); 2644 2645 hsa_insn_basic *clearing 2646 = new hsa_insn_basic (3, BRIG_OPCODE_AND, mem_type, cleared_reg, 2647 value_reg, c); 2648 hbb->append_insn (clearing); 2649 2650 /* Shift to left a value that is going to be stored. */ 2651 hsa_op_reg *new_value_reg = new hsa_op_reg (mem_type); 2652 2653 hsa_insn_basic *basic = new hsa_insn_basic (2, BRIG_OPCODE_MOV, mem_type, 2654 new_value_reg, src); 2655 hsa_fixup_mov_insn_type (basic); 2656 hbb->append_insn (basic); 2657 2658 if (bitpos) 2659 { 2660 hsa_op_reg *shifted_value_reg = new hsa_op_reg (mem_type); 2661 c = new hsa_op_immed (bitpos, BRIG_TYPE_U32); 2662 2663 hsa_insn_basic *basic 2664 = new hsa_insn_basic (3, BRIG_OPCODE_SHL, mem_type, 2665 shifted_value_reg, new_value_reg, c); 2666 hbb->append_insn (basic); 2667 2668 new_value_reg = shifted_value_reg; 2669 } 2670 2671 /* OR the prepared value with prepared chunk loaded from memory. */ 2672 hsa_op_reg *prepared_reg= new hsa_op_reg (mem_type); 2673 basic = new hsa_insn_basic (3, BRIG_OPCODE_OR, mem_type, prepared_reg, 2674 new_value_reg, cleared_reg); 2675 hbb->append_insn (basic); 2676 2677 src = prepared_reg; 2678 mtype = mem_type; 2679 } 2680 else 2681 req_align = hsa_object_alignment (lhs); 2682 2683 hsa_insn_mem *mem = new hsa_insn_mem (BRIG_OPCODE_ST, mtype, src, addr); 2684 mem->set_align (req_align); 2685 2686 /* The HSAIL verifier has another constraint: if the source is an immediate 2687 then it must match the destination type. If it's a register the low bits 2688 will be used for sub-word stores. We're always allocating new operands so 2689 we can modify the above in place. */ 2690 if (hsa_op_immed *imm = dyn_cast <hsa_op_immed *> (src)) 2691 { 2692 if (!hsa_type_packed_p (imm->m_type)) 2693 imm->m_type = mem->m_type; 2694 else 2695 { 2696 /* ...and all vector immediates apparently need to be vectors of 2697 unsigned bytes. */ 2698 unsigned bs = hsa_type_bit_size (imm->m_type); 2699 gcc_assert (bs == hsa_type_bit_size (mem->m_type)); 2700 switch (bs) 2701 { 2702 case 32: 2703 imm->m_type = BRIG_TYPE_U8X4; 2704 break; 2705 case 64: 2706 imm->m_type = BRIG_TYPE_U8X8; 2707 break; 2708 case 128: 2709 imm->m_type = BRIG_TYPE_U8X16; 2710 break; 2711 default: 2712 gcc_unreachable (); 2713 } 2714 } 2715 } 2716 2717 hbb->append_insn (mem); 2718 } 2719 2720 /* Generate memory copy instructions that are going to be used 2721 for copying a SRC memory to TARGET memory, 2722 represented by pointer in a register. MIN_ALIGN is minimal alignment 2723 of provided HSA addresses. */ 2724 2725 static void 2726 gen_hsa_memory_copy (hsa_bb *hbb, hsa_op_address *target, hsa_op_address *src, 2727 unsigned size, BrigAlignment8_t min_align) 2728 { 2729 hsa_op_address *addr; 2730 hsa_insn_mem *mem; 2731 2732 unsigned offset = 0; 2733 unsigned min_byte_align = hsa_byte_alignment (min_align); 2734 2735 while (size) 2736 { 2737 unsigned s; 2738 if (size >= 8) 2739 s = 8; 2740 else if (size >= 4) 2741 s = 4; 2742 else if (size >= 2) 2743 s = 2; 2744 else 2745 s = 1; 2746 2747 if (s > min_byte_align) 2748 s = min_byte_align; 2749 2750 BrigType16_t t = get_integer_type_by_bytes (s, false); 2751 2752 hsa_op_reg *tmp = new hsa_op_reg (t); 2753 addr = new hsa_op_address (src->m_symbol, src->m_reg, 2754 src->m_imm_offset + offset); 2755 mem = new hsa_insn_mem (BRIG_OPCODE_LD, t, tmp, addr); 2756 hbb->append_insn (mem); 2757 2758 addr = new hsa_op_address (target->m_symbol, target->m_reg, 2759 target->m_imm_offset + offset); 2760 mem = new hsa_insn_mem (BRIG_OPCODE_ST, t, tmp, addr); 2761 hbb->append_insn (mem); 2762 offset += s; 2763 size -= s; 2764 } 2765 } 2766 2767 /* Create a memset mask that is created by copying a CONSTANT byte value 2768 to an integer of BYTE_SIZE bytes. */ 2769 2770 static unsigned HOST_WIDE_INT 2771 build_memset_value (unsigned HOST_WIDE_INT constant, unsigned byte_size) 2772 { 2773 if (constant == 0) 2774 return 0; 2775 2776 HOST_WIDE_INT v = constant; 2777 2778 for (unsigned i = 1; i < byte_size; i++) 2779 v |= constant << (8 * i); 2780 2781 return v; 2782 } 2783 2784 /* Generate memory set instructions that are going to be used 2785 for setting a CONSTANT byte value to TARGET memory of SIZE bytes. 2786 MIN_ALIGN is minimal alignment of provided HSA addresses. */ 2787 2788 static void 2789 gen_hsa_memory_set (hsa_bb *hbb, hsa_op_address *target, 2790 unsigned HOST_WIDE_INT constant, 2791 unsigned size, BrigAlignment8_t min_align) 2792 { 2793 hsa_op_address *addr; 2794 hsa_insn_mem *mem; 2795 2796 unsigned offset = 0; 2797 unsigned min_byte_align = hsa_byte_alignment (min_align); 2798 2799 while (size) 2800 { 2801 unsigned s; 2802 if (size >= 8) 2803 s = 8; 2804 else if (size >= 4) 2805 s = 4; 2806 else if (size >= 2) 2807 s = 2; 2808 else 2809 s = 1; 2810 2811 if (s > min_byte_align) 2812 s = min_byte_align; 2813 2814 addr = new hsa_op_address (target->m_symbol, target->m_reg, 2815 target->m_imm_offset + offset); 2816 2817 BrigType16_t t = get_integer_type_by_bytes (s, false); 2818 HOST_WIDE_INT c = build_memset_value (constant, s); 2819 2820 mem = new hsa_insn_mem (BRIG_OPCODE_ST, t, new hsa_op_immed (c, t), 2821 addr); 2822 hbb->append_insn (mem); 2823 offset += s; 2824 size -= s; 2825 } 2826 } 2827 2828 /* Generate HSAIL instructions for a single assignment 2829 of an empty constructor to an ADDR_LHS. Constructor is passed as a 2830 tree RHS and all instructions are appended to HBB. ALIGN is 2831 alignment of the address. */ 2832 2833 void 2834 gen_hsa_ctor_assignment (hsa_op_address *addr_lhs, tree rhs, hsa_bb *hbb, 2835 BrigAlignment8_t align) 2836 { 2837 if (CONSTRUCTOR_NELTS (rhs)) 2838 { 2839 HSA_SORRY_AT (EXPR_LOCATION (rhs), 2840 "support for HSA does not implement load from constructor"); 2841 return; 2842 } 2843 2844 unsigned size = tree_to_uhwi (TYPE_SIZE_UNIT (TREE_TYPE (rhs))); 2845 gen_hsa_memory_set (hbb, addr_lhs, 0, size, align); 2846 } 2847 2848 /* Generate HSA instructions for a single assignment of RHS to LHS. 2849 HBB is the basic block they will be appended to. */ 2850 2851 static void 2852 gen_hsa_insns_for_single_assignment (tree lhs, tree rhs, hsa_bb *hbb) 2853 { 2854 if (TREE_CODE (lhs) == SSA_NAME) 2855 { 2856 hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (lhs); 2857 if (hsa_seen_error ()) 2858 return; 2859 2860 gen_hsa_insns_for_load (dest, rhs, TREE_TYPE (lhs), hbb); 2861 } 2862 else if (TREE_CODE (rhs) == SSA_NAME 2863 || (is_gimple_min_invariant (rhs) && TREE_CODE (rhs) != STRING_CST)) 2864 { 2865 /* Store to memory. */ 2866 hsa_op_base *src = hsa_reg_or_immed_for_gimple_op (rhs, hbb); 2867 if (hsa_seen_error ()) 2868 return; 2869 2870 gen_hsa_insns_for_store (lhs, src, hbb); 2871 } 2872 else 2873 { 2874 BrigAlignment8_t lhs_align; 2875 hsa_op_address *addr_lhs = gen_hsa_addr_with_align (lhs, hbb, 2876 &lhs_align); 2877 2878 if (TREE_CODE (rhs) == CONSTRUCTOR) 2879 gen_hsa_ctor_assignment (addr_lhs, rhs, hbb, lhs_align); 2880 else 2881 { 2882 BrigAlignment8_t rhs_align; 2883 hsa_op_address *addr_rhs = gen_hsa_addr_with_align (rhs, hbb, 2884 &rhs_align); 2885 2886 unsigned size = tree_to_uhwi (TYPE_SIZE_UNIT (TREE_TYPE (rhs))); 2887 gen_hsa_memory_copy (hbb, addr_lhs, addr_rhs, size, 2888 MIN (lhs_align, rhs_align)); 2889 } 2890 } 2891 } 2892 2893 /* Prepend before INSN a load from spill symbol of SPILL_REG. Return the 2894 register into which we loaded. If this required another register to convert 2895 from a B1 type, return it in *PTMP2, otherwise store NULL into it. We 2896 assume we are out of SSA so the returned register does not have its 2897 definition set. */ 2898 2899 hsa_op_reg * 2900 hsa_spill_in (hsa_insn_basic *insn, hsa_op_reg *spill_reg, hsa_op_reg **ptmp2) 2901 { 2902 hsa_symbol *spill_sym = spill_reg->m_spill_sym; 2903 hsa_op_reg *reg = new hsa_op_reg (spill_sym->m_type); 2904 hsa_op_address *addr = new hsa_op_address (spill_sym); 2905 2906 hsa_insn_mem *mem = new hsa_insn_mem (BRIG_OPCODE_LD, spill_sym->m_type, 2907 reg, addr); 2908 hsa_insert_insn_before (mem, insn); 2909 2910 *ptmp2 = NULL; 2911 if (spill_reg->m_type == BRIG_TYPE_B1) 2912 { 2913 hsa_insn_basic *cvtinsn; 2914 *ptmp2 = reg; 2915 reg = new hsa_op_reg (spill_reg->m_type); 2916 2917 cvtinsn = new hsa_insn_cvt (reg, *ptmp2); 2918 hsa_insert_insn_before (cvtinsn, insn); 2919 } 2920 return reg; 2921 } 2922 2923 /* Append after INSN a store to spill symbol of SPILL_REG. Return the register 2924 from which we stored. If this required another register to convert to a B1 2925 type, return it in *PTMP2, otherwise store NULL into it. We assume we are 2926 out of SSA so the returned register does not have its use updated. */ 2927 2928 hsa_op_reg * 2929 hsa_spill_out (hsa_insn_basic *insn, hsa_op_reg *spill_reg, hsa_op_reg **ptmp2) 2930 { 2931 hsa_symbol *spill_sym = spill_reg->m_spill_sym; 2932 hsa_op_reg *reg = new hsa_op_reg (spill_sym->m_type); 2933 hsa_op_address *addr = new hsa_op_address (spill_sym); 2934 hsa_op_reg *returnreg; 2935 2936 *ptmp2 = NULL; 2937 returnreg = reg; 2938 if (spill_reg->m_type == BRIG_TYPE_B1) 2939 { 2940 hsa_insn_basic *cvtinsn; 2941 *ptmp2 = new hsa_op_reg (spill_sym->m_type); 2942 reg->m_type = spill_reg->m_type; 2943 2944 cvtinsn = new hsa_insn_cvt (*ptmp2, returnreg); 2945 hsa_append_insn_after (cvtinsn, insn); 2946 insn = cvtinsn; 2947 reg = *ptmp2; 2948 } 2949 2950 hsa_insn_mem *mem = new hsa_insn_mem (BRIG_OPCODE_ST, spill_sym->m_type, reg, 2951 addr); 2952 hsa_append_insn_after (mem, insn); 2953 return returnreg; 2954 } 2955 2956 /* Generate a comparison instruction that will compare LHS and RHS with 2957 comparison specified by CODE and put result into register DEST. DEST has to 2958 have its type set already but must not have its definition set yet. 2959 Generated instructions will be added to HBB. */ 2960 2961 static void 2962 gen_hsa_cmp_insn_from_gimple (enum tree_code code, tree lhs, tree rhs, 2963 hsa_op_reg *dest, hsa_bb *hbb) 2964 { 2965 BrigCompareOperation8_t compare; 2966 2967 switch (code) 2968 { 2969 case LT_EXPR: 2970 compare = BRIG_COMPARE_LT; 2971 break; 2972 case LE_EXPR: 2973 compare = BRIG_COMPARE_LE; 2974 break; 2975 case GT_EXPR: 2976 compare = BRIG_COMPARE_GT; 2977 break; 2978 case GE_EXPR: 2979 compare = BRIG_COMPARE_GE; 2980 break; 2981 case EQ_EXPR: 2982 compare = BRIG_COMPARE_EQ; 2983 break; 2984 case NE_EXPR: 2985 compare = BRIG_COMPARE_NE; 2986 break; 2987 case UNORDERED_EXPR: 2988 compare = BRIG_COMPARE_NAN; 2989 break; 2990 case ORDERED_EXPR: 2991 compare = BRIG_COMPARE_NUM; 2992 break; 2993 case UNLT_EXPR: 2994 compare = BRIG_COMPARE_LTU; 2995 break; 2996 case UNLE_EXPR: 2997 compare = BRIG_COMPARE_LEU; 2998 break; 2999 case UNGT_EXPR: 3000 compare = BRIG_COMPARE_GTU; 3001 break; 3002 case UNGE_EXPR: 3003 compare = BRIG_COMPARE_GEU; 3004 break; 3005 case UNEQ_EXPR: 3006 compare = BRIG_COMPARE_EQU; 3007 break; 3008 case LTGT_EXPR: 3009 compare = BRIG_COMPARE_NEU; 3010 break; 3011 3012 default: 3013 HSA_SORRY_ATV (EXPR_LOCATION (lhs), 3014 "support for HSA does not implement comparison tree " 3015 "code %s\n", get_tree_code_name (code)); 3016 return; 3017 } 3018 3019 /* CMP instruction returns e.g. 0xffffffff (for a 32-bit with integer) 3020 as a result of comparison. */ 3021 3022 BrigType16_t dest_type = hsa_type_integer_p (dest->m_type) 3023 ? (BrigType16_t) BRIG_TYPE_B1 : dest->m_type; 3024 3025 hsa_insn_cmp *cmp = new hsa_insn_cmp (compare, dest_type); 3026 hsa_op_with_type *op1 = hsa_reg_or_immed_for_gimple_op (lhs, hbb); 3027 cmp->set_op (1, op1->extend_int_to_32bit (hbb)); 3028 hsa_op_with_type *op2 = hsa_reg_or_immed_for_gimple_op (rhs, hbb); 3029 cmp->set_op (2, op2->extend_int_to_32bit (hbb)); 3030 3031 hbb->append_insn (cmp); 3032 cmp->set_output_in_type (dest, 0, hbb); 3033 } 3034 3035 /* Generate an unary instruction with OPCODE and append it to a basic block 3036 HBB. The instruction uses DEST as a destination and OP1 3037 as a single operand. */ 3038 3039 static void 3040 gen_hsa_unary_operation (BrigOpcode opcode, hsa_op_reg *dest, 3041 hsa_op_with_type *op1, hsa_bb *hbb) 3042 { 3043 gcc_checking_assert (dest); 3044 hsa_insn_basic *insn; 3045 3046 if (opcode == BRIG_OPCODE_MOV && hsa_needs_cvt (dest->m_type, op1->m_type)) 3047 { 3048 insn = new hsa_insn_cvt (dest, op1); 3049 hbb->append_insn (insn); 3050 return; 3051 } 3052 3053 op1 = op1->extend_int_to_32bit (hbb); 3054 if (opcode == BRIG_OPCODE_FIRSTBIT || opcode == BRIG_OPCODE_LASTBIT) 3055 { 3056 BrigType16_t srctype = hsa_type_integer_p (op1->m_type) ? op1->m_type 3057 : hsa_unsigned_type_for_type (op1->m_type); 3058 insn = new hsa_insn_srctype (2, opcode, BRIG_TYPE_U32, srctype, NULL, 3059 op1); 3060 } 3061 else 3062 { 3063 BrigType16_t optype = hsa_extend_inttype_to_32bit (dest->m_type); 3064 insn = new hsa_insn_basic (2, opcode, optype, NULL, op1); 3065 3066 if (opcode == BRIG_OPCODE_MOV) 3067 hsa_fixup_mov_insn_type (insn); 3068 else if (opcode == BRIG_OPCODE_ABS || opcode == BRIG_OPCODE_NEG) 3069 { 3070 /* ABS and NEG only exist in _s form :-/ */ 3071 if (insn->m_type == BRIG_TYPE_U32) 3072 insn->m_type = BRIG_TYPE_S32; 3073 else if (insn->m_type == BRIG_TYPE_U64) 3074 insn->m_type = BRIG_TYPE_S64; 3075 } 3076 } 3077 3078 hbb->append_insn (insn); 3079 insn->set_output_in_type (dest, 0, hbb); 3080 } 3081 3082 /* Generate a binary instruction with OPCODE and append it to a basic block 3083 HBB. The instruction uses DEST as a destination and operands OP1 3084 and OP2. */ 3085 3086 static void 3087 gen_hsa_binary_operation (int opcode, hsa_op_reg *dest, 3088 hsa_op_with_type *op1, hsa_op_with_type *op2, 3089 hsa_bb *hbb) 3090 { 3091 gcc_checking_assert (dest); 3092 3093 BrigType16_t optype = hsa_extend_inttype_to_32bit (dest->m_type); 3094 op1 = op1->extend_int_to_32bit (hbb); 3095 op2 = op2->extend_int_to_32bit (hbb); 3096 3097 if ((opcode == BRIG_OPCODE_SHL || opcode == BRIG_OPCODE_SHR) 3098 && is_a <hsa_op_immed *> (op2)) 3099 { 3100 hsa_op_immed *i = dyn_cast <hsa_op_immed *> (op2); 3101 i->set_type (BRIG_TYPE_U32); 3102 } 3103 if ((opcode == BRIG_OPCODE_OR 3104 || opcode == BRIG_OPCODE_XOR 3105 || opcode == BRIG_OPCODE_AND) 3106 && is_a <hsa_op_immed *> (op2)) 3107 { 3108 hsa_op_immed *i = dyn_cast <hsa_op_immed *> (op2); 3109 i->set_type (hsa_unsigned_type_for_type (i->m_type)); 3110 } 3111 3112 hsa_insn_basic *insn = new hsa_insn_basic (3, opcode, optype, NULL, 3113 op1, op2); 3114 hbb->append_insn (insn); 3115 insn->set_output_in_type (dest, 0, hbb); 3116 } 3117 3118 /* Generate HSA instructions for a single assignment. HBB is the basic block 3119 they will be appended to. */ 3120 3121 static void 3122 gen_hsa_insns_for_operation_assignment (gimple *assign, hsa_bb *hbb) 3123 { 3124 tree_code code = gimple_assign_rhs_code (assign); 3125 gimple_rhs_class rhs_class = get_gimple_rhs_class (gimple_expr_code (assign)); 3126 3127 tree lhs = gimple_assign_lhs (assign); 3128 tree rhs1 = gimple_assign_rhs1 (assign); 3129 tree rhs2 = gimple_assign_rhs2 (assign); 3130 tree rhs3 = gimple_assign_rhs3 (assign); 3131 3132 BrigOpcode opcode; 3133 3134 switch (code) 3135 { 3136 CASE_CONVERT: 3137 case FLOAT_EXPR: 3138 /* The opcode is changed to BRIG_OPCODE_CVT if BRIG types 3139 needs a conversion. */ 3140 opcode = BRIG_OPCODE_MOV; 3141 break; 3142 3143 case PLUS_EXPR: 3144 case POINTER_PLUS_EXPR: 3145 opcode = BRIG_OPCODE_ADD; 3146 break; 3147 case MINUS_EXPR: 3148 opcode = BRIG_OPCODE_SUB; 3149 break; 3150 case MULT_EXPR: 3151 opcode = BRIG_OPCODE_MUL; 3152 break; 3153 case MULT_HIGHPART_EXPR: 3154 opcode = BRIG_OPCODE_MULHI; 3155 break; 3156 case RDIV_EXPR: 3157 case TRUNC_DIV_EXPR: 3158 case EXACT_DIV_EXPR: 3159 opcode = BRIG_OPCODE_DIV; 3160 break; 3161 case CEIL_DIV_EXPR: 3162 case FLOOR_DIV_EXPR: 3163 case ROUND_DIV_EXPR: 3164 HSA_SORRY_AT (gimple_location (assign), 3165 "support for HSA does not implement CEIL_DIV_EXPR, " 3166 "FLOOR_DIV_EXPR or ROUND_DIV_EXPR"); 3167 return; 3168 case TRUNC_MOD_EXPR: 3169 opcode = BRIG_OPCODE_REM; 3170 break; 3171 case CEIL_MOD_EXPR: 3172 case FLOOR_MOD_EXPR: 3173 case ROUND_MOD_EXPR: 3174 HSA_SORRY_AT (gimple_location (assign), 3175 "support for HSA does not implement CEIL_MOD_EXPR, " 3176 "FLOOR_MOD_EXPR or ROUND_MOD_EXPR"); 3177 return; 3178 case NEGATE_EXPR: 3179 opcode = BRIG_OPCODE_NEG; 3180 break; 3181 case FMA_EXPR: 3182 /* There is a native HSA instruction for scalar FMAs but not for vector 3183 ones. */ 3184 if (TREE_CODE (TREE_TYPE (lhs)) == VECTOR_TYPE) 3185 { 3186 hsa_op_reg *dest 3187 = hsa_cfun->reg_for_gimple_ssa (gimple_assign_lhs (assign)); 3188 hsa_op_with_type *op1 = hsa_reg_or_immed_for_gimple_op (rhs1, hbb); 3189 hsa_op_with_type *op2 = hsa_reg_or_immed_for_gimple_op (rhs2, hbb); 3190 hsa_op_with_type *op3 = hsa_reg_or_immed_for_gimple_op (rhs3, hbb); 3191 hsa_op_reg *tmp = new hsa_op_reg (dest->m_type); 3192 gen_hsa_binary_operation (BRIG_OPCODE_MUL, tmp, op1, op2, hbb); 3193 gen_hsa_binary_operation (BRIG_OPCODE_ADD, dest, tmp, op3, hbb); 3194 return; 3195 } 3196 opcode = BRIG_OPCODE_MAD; 3197 break; 3198 case MIN_EXPR: 3199 opcode = BRIG_OPCODE_MIN; 3200 break; 3201 case MAX_EXPR: 3202 opcode = BRIG_OPCODE_MAX; 3203 break; 3204 case ABS_EXPR: 3205 opcode = BRIG_OPCODE_ABS; 3206 break; 3207 case LSHIFT_EXPR: 3208 opcode = BRIG_OPCODE_SHL; 3209 break; 3210 case RSHIFT_EXPR: 3211 opcode = BRIG_OPCODE_SHR; 3212 break; 3213 case LROTATE_EXPR: 3214 case RROTATE_EXPR: 3215 { 3216 hsa_insn_basic *insn = NULL; 3217 int code1 = code == LROTATE_EXPR ? BRIG_OPCODE_SHL : BRIG_OPCODE_SHR; 3218 int code2 = code != LROTATE_EXPR ? BRIG_OPCODE_SHL : BRIG_OPCODE_SHR; 3219 BrigType16_t btype = hsa_type_for_scalar_tree_type (TREE_TYPE (lhs), 3220 true); 3221 3222 hsa_op_with_type *src = hsa_reg_or_immed_for_gimple_op (rhs1, hbb); 3223 hsa_op_reg *op1 = new hsa_op_reg (btype); 3224 hsa_op_reg *op2 = new hsa_op_reg (btype); 3225 hsa_op_with_type *shift1 = hsa_reg_or_immed_for_gimple_op (rhs2, hbb); 3226 3227 tree type = TREE_TYPE (rhs2); 3228 unsigned HOST_WIDE_INT bitsize = TREE_INT_CST_LOW (TYPE_SIZE (type)); 3229 3230 hsa_op_with_type *shift2 = NULL; 3231 if (TREE_CODE (rhs2) == INTEGER_CST) 3232 shift2 = new hsa_op_immed (bitsize - tree_to_uhwi (rhs2), 3233 BRIG_TYPE_U32); 3234 else if (TREE_CODE (rhs2) == SSA_NAME) 3235 { 3236 hsa_op_reg *s = hsa_cfun->reg_for_gimple_ssa (rhs2); 3237 s = as_a <hsa_op_reg *> (s->extend_int_to_32bit (hbb)); 3238 hsa_op_reg *d = new hsa_op_reg (s->m_type); 3239 hsa_op_immed *size_imm = new hsa_op_immed (bitsize, BRIG_TYPE_U32); 3240 3241 insn = new hsa_insn_basic (3, BRIG_OPCODE_SUB, d->m_type, 3242 d, s, size_imm); 3243 hbb->append_insn (insn); 3244 3245 shift2 = d; 3246 } 3247 else 3248 gcc_unreachable (); 3249 3250 hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (lhs); 3251 gen_hsa_binary_operation (code1, op1, src, shift1, hbb); 3252 gen_hsa_binary_operation (code2, op2, src, shift2, hbb); 3253 gen_hsa_binary_operation (BRIG_OPCODE_OR, dest, op1, op2, hbb); 3254 3255 return; 3256 } 3257 case BIT_IOR_EXPR: 3258 opcode = BRIG_OPCODE_OR; 3259 break; 3260 case BIT_XOR_EXPR: 3261 opcode = BRIG_OPCODE_XOR; 3262 break; 3263 case BIT_AND_EXPR: 3264 opcode = BRIG_OPCODE_AND; 3265 break; 3266 case BIT_NOT_EXPR: 3267 opcode = BRIG_OPCODE_NOT; 3268 break; 3269 case FIX_TRUNC_EXPR: 3270 { 3271 hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (lhs); 3272 hsa_op_with_type *v = hsa_reg_or_immed_for_gimple_op (rhs1, hbb); 3273 3274 if (hsa_needs_cvt (dest->m_type, v->m_type)) 3275 { 3276 hsa_op_reg *tmp = new hsa_op_reg (v->m_type); 3277 3278 hsa_insn_basic *insn = new hsa_insn_basic (2, BRIG_OPCODE_TRUNC, 3279 tmp->m_type, tmp, v); 3280 hbb->append_insn (insn); 3281 3282 hsa_insn_basic *cvtinsn = new hsa_insn_cvt (dest, tmp); 3283 hbb->append_insn (cvtinsn); 3284 } 3285 else 3286 { 3287 hsa_insn_basic *insn = new hsa_insn_basic (2, BRIG_OPCODE_TRUNC, 3288 dest->m_type, dest, v); 3289 hbb->append_insn (insn); 3290 } 3291 3292 return; 3293 } 3294 opcode = BRIG_OPCODE_TRUNC; 3295 break; 3296 3297 case LT_EXPR: 3298 case LE_EXPR: 3299 case GT_EXPR: 3300 case GE_EXPR: 3301 case EQ_EXPR: 3302 case NE_EXPR: 3303 case UNORDERED_EXPR: 3304 case ORDERED_EXPR: 3305 case UNLT_EXPR: 3306 case UNLE_EXPR: 3307 case UNGT_EXPR: 3308 case UNGE_EXPR: 3309 case UNEQ_EXPR: 3310 case LTGT_EXPR: 3311 { 3312 hsa_op_reg *dest 3313 = hsa_cfun->reg_for_gimple_ssa (gimple_assign_lhs (assign)); 3314 3315 gen_hsa_cmp_insn_from_gimple (code, rhs1, rhs2, dest, hbb); 3316 return; 3317 } 3318 case COND_EXPR: 3319 { 3320 hsa_op_reg *dest 3321 = hsa_cfun->reg_for_gimple_ssa (gimple_assign_lhs (assign)); 3322 hsa_op_with_type *ctrl = NULL; 3323 tree cond = rhs1; 3324 3325 if (CONSTANT_CLASS_P (cond) || TREE_CODE (cond) == SSA_NAME) 3326 ctrl = hsa_reg_or_immed_for_gimple_op (cond, hbb); 3327 else 3328 { 3329 hsa_op_reg *r = new hsa_op_reg (BRIG_TYPE_B1); 3330 3331 gen_hsa_cmp_insn_from_gimple (TREE_CODE (cond), 3332 TREE_OPERAND (cond, 0), 3333 TREE_OPERAND (cond, 1), 3334 r, hbb); 3335 3336 ctrl = r; 3337 } 3338 3339 hsa_op_with_type *op2 = hsa_reg_or_immed_for_gimple_op (rhs2, hbb); 3340 hsa_op_with_type *op3 = hsa_reg_or_immed_for_gimple_op (rhs3, hbb); 3341 op2 = op2->extend_int_to_32bit (hbb); 3342 op3 = op3->extend_int_to_32bit (hbb); 3343 3344 BrigType16_t type = hsa_extend_inttype_to_32bit (dest->m_type); 3345 BrigType16_t utype = hsa_unsigned_type_for_type (type); 3346 if (is_a <hsa_op_immed *> (op2)) 3347 op2->m_type = utype; 3348 if (is_a <hsa_op_immed *> (op3)) 3349 op3->m_type = utype; 3350 3351 hsa_insn_basic *insn 3352 = new hsa_insn_basic (4, BRIG_OPCODE_CMOV, 3353 hsa_bittype_for_type (type), 3354 NULL, ctrl, op2, op3); 3355 3356 hbb->append_insn (insn); 3357 insn->set_output_in_type (dest, 0, hbb); 3358 return; 3359 } 3360 case COMPLEX_EXPR: 3361 { 3362 hsa_op_reg *dest 3363 = hsa_cfun->reg_for_gimple_ssa (gimple_assign_lhs (assign)); 3364 hsa_op_with_type *rhs1_reg = hsa_reg_or_immed_for_gimple_op (rhs1, hbb); 3365 rhs1_reg = rhs1_reg->extend_int_to_32bit (hbb); 3366 hsa_op_with_type *rhs2_reg = hsa_reg_or_immed_for_gimple_op (rhs2, hbb); 3367 rhs2_reg = rhs2_reg->extend_int_to_32bit (hbb); 3368 3369 if (hsa_seen_error ()) 3370 return; 3371 3372 BrigType16_t src_type = hsa_bittype_for_type (rhs1_reg->m_type); 3373 rhs1_reg = rhs1_reg->get_in_type (src_type, hbb); 3374 rhs2_reg = rhs2_reg->get_in_type (src_type, hbb); 3375 3376 hsa_insn_packed *insn 3377 = new hsa_insn_packed (3, BRIG_OPCODE_COMBINE, dest->m_type, src_type, 3378 dest, rhs1_reg, rhs2_reg); 3379 hbb->append_insn (insn); 3380 3381 return; 3382 } 3383 default: 3384 /* Implement others as we come across them. */ 3385 HSA_SORRY_ATV (gimple_location (assign), 3386 "support for HSA does not implement operation %s", 3387 get_tree_code_name (code)); 3388 return; 3389 } 3390 3391 3392 hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (lhs); 3393 hsa_op_with_type *op1 = hsa_reg_or_immed_for_gimple_op (rhs1, hbb); 3394 hsa_op_with_type *op2 3395 = rhs2 ? hsa_reg_or_immed_for_gimple_op (rhs2, hbb) : NULL; 3396 3397 if (hsa_seen_error ()) 3398 return; 3399 3400 switch (rhs_class) 3401 { 3402 case GIMPLE_TERNARY_RHS: 3403 { 3404 hsa_op_with_type *op3 = hsa_reg_or_immed_for_gimple_op (rhs3, hbb); 3405 op3 = op3->extend_int_to_32bit (hbb); 3406 hsa_insn_basic *insn = new hsa_insn_basic (4, opcode, dest->m_type, dest, 3407 op1, op2, op3); 3408 hbb->append_insn (insn); 3409 } 3410 return; 3411 3412 case GIMPLE_BINARY_RHS: 3413 gen_hsa_binary_operation (opcode, dest, op1, op2, hbb); 3414 break; 3415 3416 case GIMPLE_UNARY_RHS: 3417 gen_hsa_unary_operation (opcode, dest, op1, hbb); 3418 break; 3419 default: 3420 gcc_unreachable (); 3421 } 3422 } 3423 3424 /* Generate HSA instructions for a given gimple condition statement COND. 3425 Instructions will be appended to HBB, which also needs to be the 3426 corresponding structure to the basic_block of COND. */ 3427 3428 static void 3429 gen_hsa_insns_for_cond_stmt (gimple *cond, hsa_bb *hbb) 3430 { 3431 hsa_op_reg *ctrl = new hsa_op_reg (BRIG_TYPE_B1); 3432 hsa_insn_cbr *cbr; 3433 3434 gen_hsa_cmp_insn_from_gimple (gimple_cond_code (cond), 3435 gimple_cond_lhs (cond), 3436 gimple_cond_rhs (cond), 3437 ctrl, hbb); 3438 3439 cbr = new hsa_insn_cbr (ctrl); 3440 hbb->append_insn (cbr); 3441 } 3442 3443 /* Maximum number of elements in a jump table for an HSA SBR instruction. */ 3444 3445 #define HSA_MAXIMUM_SBR_LABELS 16 3446 3447 /* Return lowest value of a switch S that is handled in a non-default 3448 label. */ 3449 3450 static tree 3451 get_switch_low (gswitch *s) 3452 { 3453 unsigned labels = gimple_switch_num_labels (s); 3454 gcc_checking_assert (labels >= 1); 3455 3456 return CASE_LOW (gimple_switch_label (s, 1)); 3457 } 3458 3459 /* Return highest value of a switch S that is handled in a non-default 3460 label. */ 3461 3462 static tree 3463 get_switch_high (gswitch *s) 3464 { 3465 unsigned labels = gimple_switch_num_labels (s); 3466 3467 /* Compare last label to maximum number of labels. */ 3468 tree label = gimple_switch_label (s, labels - 1); 3469 tree low = CASE_LOW (label); 3470 tree high = CASE_HIGH (label); 3471 3472 return high != NULL_TREE ? high : low; 3473 } 3474 3475 static tree 3476 get_switch_size (gswitch *s) 3477 { 3478 return int_const_binop (MINUS_EXPR, get_switch_high (s), get_switch_low (s)); 3479 } 3480 3481 /* Generate HSA instructions for a given gimple switch. 3482 Instructions will be appended to HBB. */ 3483 3484 static void 3485 gen_hsa_insns_for_switch_stmt (gswitch *s, hsa_bb *hbb) 3486 { 3487 gimple_stmt_iterator it = gsi_for_stmt (s); 3488 gsi_prev (&it); 3489 3490 /* Create preambule that verifies that index - lowest_label >= 0. */ 3491 edge e = split_block (hbb->m_bb, gsi_stmt (it)); 3492 e->flags &= ~EDGE_FALLTHRU; 3493 e->flags |= EDGE_TRUE_VALUE; 3494 3495 function *func = DECL_STRUCT_FUNCTION (current_function_decl); 3496 tree index_tree = gimple_switch_index (s); 3497 tree lowest = get_switch_low (s); 3498 tree highest = get_switch_high (s); 3499 3500 hsa_op_reg *index = hsa_cfun->reg_for_gimple_ssa (index_tree); 3501 index = as_a <hsa_op_reg *> (index->extend_int_to_32bit (hbb)); 3502 3503 hsa_op_reg *cmp1_reg = new hsa_op_reg (BRIG_TYPE_B1); 3504 hsa_op_immed *cmp1_immed = new hsa_op_immed (lowest, true); 3505 hbb->append_insn (new hsa_insn_cmp (BRIG_COMPARE_GE, cmp1_reg->m_type, 3506 cmp1_reg, index, cmp1_immed)); 3507 3508 hsa_op_reg *cmp2_reg = new hsa_op_reg (BRIG_TYPE_B1); 3509 hsa_op_immed *cmp2_immed = new hsa_op_immed (highest, true); 3510 hbb->append_insn (new hsa_insn_cmp (BRIG_COMPARE_LE, cmp2_reg->m_type, 3511 cmp2_reg, index, cmp2_immed)); 3512 3513 hsa_op_reg *cmp_reg = new hsa_op_reg (BRIG_TYPE_B1); 3514 hbb->append_insn (new hsa_insn_basic (3, BRIG_OPCODE_AND, cmp_reg->m_type, 3515 cmp_reg, cmp1_reg, cmp2_reg)); 3516 3517 hbb->append_insn (new hsa_insn_cbr (cmp_reg)); 3518 3519 tree default_label = gimple_switch_default_label (s); 3520 basic_block default_label_bb = label_to_block_fn (func, 3521 CASE_LABEL (default_label)); 3522 3523 if (!gimple_seq_empty_p (phi_nodes (default_label_bb))) 3524 { 3525 default_label_bb = split_edge (find_edge (e->dest, default_label_bb)); 3526 hsa_init_new_bb (default_label_bb); 3527 } 3528 3529 make_edge (e->src, default_label_bb, EDGE_FALSE_VALUE); 3530 3531 hsa_cfun->m_modified_cfg = true; 3532 3533 /* Basic block with the SBR instruction. */ 3534 hbb = hsa_init_new_bb (e->dest); 3535 3536 hsa_op_reg *sub_index = new hsa_op_reg (index->m_type); 3537 hbb->append_insn (new hsa_insn_basic (3, BRIG_OPCODE_SUB, sub_index->m_type, 3538 sub_index, index, 3539 new hsa_op_immed (lowest, true))); 3540 3541 hsa_op_base *tmp = sub_index->get_in_type (BRIG_TYPE_U64, hbb); 3542 sub_index = as_a <hsa_op_reg *> (tmp); 3543 unsigned labels = gimple_switch_num_labels (s); 3544 unsigned HOST_WIDE_INT size = tree_to_uhwi (get_switch_size (s)); 3545 3546 hsa_insn_sbr *sbr = new hsa_insn_sbr (sub_index, size + 1); 3547 3548 /* Prepare array with default label destination. */ 3549 for (unsigned HOST_WIDE_INT i = 0; i <= size; i++) 3550 sbr->m_jump_table.safe_push (default_label_bb); 3551 3552 /* Iterate all labels and fill up the jump table. */ 3553 for (unsigned i = 1; i < labels; i++) 3554 { 3555 tree label = gimple_switch_label (s, i); 3556 basic_block bb = label_to_block_fn (func, CASE_LABEL (label)); 3557 3558 unsigned HOST_WIDE_INT sub_low 3559 = tree_to_uhwi (int_const_binop (MINUS_EXPR, CASE_LOW (label), lowest)); 3560 3561 unsigned HOST_WIDE_INT sub_high = sub_low; 3562 tree high = CASE_HIGH (label); 3563 if (high != NULL) 3564 sub_high = tree_to_uhwi (int_const_binop (MINUS_EXPR, high, lowest)); 3565 3566 for (unsigned HOST_WIDE_INT j = sub_low; j <= sub_high; j++) 3567 sbr->m_jump_table[j] = bb; 3568 } 3569 3570 hbb->append_insn (sbr); 3571 } 3572 3573 /* Verify that the function DECL can be handled by HSA. */ 3574 3575 static void 3576 verify_function_arguments (tree decl) 3577 { 3578 tree type = TREE_TYPE (decl); 3579 if (DECL_STATIC_CHAIN (decl)) 3580 { 3581 HSA_SORRY_ATV (EXPR_LOCATION (decl), 3582 "HSA does not support nested functions: %qD", decl); 3583 return; 3584 } 3585 else if (!TYPE_ARG_TYPES (type) || stdarg_p (type)) 3586 { 3587 HSA_SORRY_ATV (EXPR_LOCATION (decl), 3588 "HSA does not support functions with variadic arguments " 3589 "(or unknown return type): %qD", decl); 3590 return; 3591 } 3592 } 3593 3594 /* Return BRIG type for FORMAL_ARG_TYPE. If the formal argument type is NULL, 3595 return ACTUAL_ARG_TYPE. */ 3596 3597 static BrigType16_t 3598 get_format_argument_type (tree formal_arg_type, BrigType16_t actual_arg_type) 3599 { 3600 if (formal_arg_type == NULL) 3601 return actual_arg_type; 3602 3603 BrigType16_t decl_type 3604 = hsa_type_for_scalar_tree_type (formal_arg_type, false); 3605 return mem_type_for_type (decl_type); 3606 } 3607 3608 /* Generate HSA instructions for a direct call instruction. 3609 Instructions will be appended to HBB, which also needs to be the 3610 corresponding structure to the basic_block of STMT. 3611 If ASSIGN_LHS is false, do not copy HSA function result argument into the 3612 corresponding HSA representation of the gimple statement LHS. */ 3613 3614 static void 3615 gen_hsa_insns_for_direct_call (gimple *stmt, hsa_bb *hbb, 3616 bool assign_lhs = true) 3617 { 3618 tree decl = gimple_call_fndecl (stmt); 3619 verify_function_arguments (decl); 3620 if (hsa_seen_error ()) 3621 return; 3622 3623 hsa_insn_call *call_insn = new hsa_insn_call (decl); 3624 hsa_cfun->m_called_functions.safe_push (call_insn->m_called_function); 3625 3626 /* Argument block start. */ 3627 hsa_insn_arg_block *arg_start 3628 = new hsa_insn_arg_block (BRIG_KIND_DIRECTIVE_ARG_BLOCK_START, call_insn); 3629 hbb->append_insn (arg_start); 3630 3631 tree parm_type_chain = TYPE_ARG_TYPES (gimple_call_fntype (stmt)); 3632 3633 /* Preparation of arguments that will be passed to function. */ 3634 const unsigned args = gimple_call_num_args (stmt); 3635 for (unsigned i = 0; i < args; ++i) 3636 { 3637 tree parm = gimple_call_arg (stmt, (int)i); 3638 tree parm_decl_type = parm_type_chain != NULL_TREE 3639 ? TREE_VALUE (parm_type_chain) : NULL_TREE; 3640 hsa_op_address *addr; 3641 3642 if (AGGREGATE_TYPE_P (TREE_TYPE (parm))) 3643 { 3644 addr = gen_hsa_addr_for_arg (TREE_TYPE (parm), i); 3645 BrigAlignment8_t align; 3646 hsa_op_address *src = gen_hsa_addr_with_align (parm, hbb, &align); 3647 gen_hsa_memory_copy (hbb, addr, src, 3648 addr->m_symbol->total_byte_size (), align); 3649 } 3650 else 3651 { 3652 hsa_op_with_type *src = hsa_reg_or_immed_for_gimple_op (parm, hbb); 3653 3654 if (parm_decl_type != NULL && AGGREGATE_TYPE_P (parm_decl_type)) 3655 { 3656 HSA_SORRY_AT (gimple_location (stmt), 3657 "support for HSA does not implement an aggregate " 3658 "formal argument in a function call, while actual " 3659 "argument is not an aggregate"); 3660 return; 3661 } 3662 3663 BrigType16_t formal_arg_type 3664 = get_format_argument_type (parm_decl_type, src->m_type); 3665 if (hsa_seen_error ()) 3666 return; 3667 3668 if (src->m_type != formal_arg_type) 3669 src = src->get_in_type (formal_arg_type, hbb); 3670 3671 addr 3672 = gen_hsa_addr_for_arg (parm_decl_type != NULL_TREE ? 3673 parm_decl_type: TREE_TYPE (parm), i); 3674 hsa_insn_mem *mem = new hsa_insn_mem (BRIG_OPCODE_ST, formal_arg_type, 3675 src, addr); 3676 3677 hbb->append_insn (mem); 3678 } 3679 3680 call_insn->m_input_args.safe_push (addr->m_symbol); 3681 if (parm_type_chain) 3682 parm_type_chain = TREE_CHAIN (parm_type_chain); 3683 } 3684 3685 call_insn->m_args_code_list = new hsa_op_code_list (args); 3686 hbb->append_insn (call_insn); 3687 3688 tree result_type = TREE_TYPE (TREE_TYPE (decl)); 3689 3690 tree result = gimple_call_lhs (stmt); 3691 hsa_insn_mem *result_insn = NULL; 3692 if (!VOID_TYPE_P (result_type)) 3693 { 3694 hsa_op_address *addr = gen_hsa_addr_for_arg (result_type, -1); 3695 3696 /* Even if result of a function call is unused, we have to emit 3697 declaration for the result. */ 3698 if (result && assign_lhs) 3699 { 3700 tree lhs_type = TREE_TYPE (result); 3701 3702 if (hsa_seen_error ()) 3703 return; 3704 3705 if (AGGREGATE_TYPE_P (lhs_type)) 3706 { 3707 BrigAlignment8_t align; 3708 hsa_op_address *result_addr 3709 = gen_hsa_addr_with_align (result, hbb, &align); 3710 gen_hsa_memory_copy (hbb, result_addr, addr, 3711 addr->m_symbol->total_byte_size (), align); 3712 } 3713 else 3714 { 3715 BrigType16_t mtype 3716 = mem_type_for_type (hsa_type_for_scalar_tree_type (lhs_type, 3717 false)); 3718 3719 hsa_op_reg *dst = hsa_cfun->reg_for_gimple_ssa (result); 3720 result_insn = new hsa_insn_mem (BRIG_OPCODE_LD, mtype, dst, addr); 3721 hbb->append_insn (result_insn); 3722 } 3723 } 3724 3725 call_insn->m_output_arg = addr->m_symbol; 3726 call_insn->m_result_code_list = new hsa_op_code_list (1); 3727 } 3728 else 3729 { 3730 if (result) 3731 { 3732 HSA_SORRY_AT (gimple_location (stmt), 3733 "support for HSA does not implement an assignment of " 3734 "return value from a void function"); 3735 return; 3736 } 3737 3738 call_insn->m_result_code_list = new hsa_op_code_list (0); 3739 } 3740 3741 /* Argument block end. */ 3742 hsa_insn_arg_block *arg_end 3743 = new hsa_insn_arg_block (BRIG_KIND_DIRECTIVE_ARG_BLOCK_END, call_insn); 3744 hbb->append_insn (arg_end); 3745 } 3746 3747 /* Generate HSA instructions for a direct call of an internal fn. 3748 Instructions will be appended to HBB, which also needs to be the 3749 corresponding structure to the basic_block of STMT. */ 3750 3751 static void 3752 gen_hsa_insns_for_call_of_internal_fn (gimple *stmt, hsa_bb *hbb) 3753 { 3754 tree lhs = gimple_call_lhs (stmt); 3755 if (!lhs) 3756 return; 3757 3758 tree lhs_type = TREE_TYPE (lhs); 3759 tree rhs1 = gimple_call_arg (stmt, 0); 3760 tree rhs1_type = TREE_TYPE (rhs1); 3761 enum internal_fn fn = gimple_call_internal_fn (stmt); 3762 hsa_internal_fn *ifn 3763 = new hsa_internal_fn (fn, tree_to_uhwi (TYPE_SIZE (rhs1_type))); 3764 hsa_insn_call *call_insn = new hsa_insn_call (ifn); 3765 3766 gcc_checking_assert (FLOAT_TYPE_P (rhs1_type)); 3767 3768 if (!hsa_emitted_internal_decls->find (call_insn->m_called_internal_fn)) 3769 hsa_cfun->m_called_internal_fns.safe_push (call_insn->m_called_internal_fn); 3770 3771 hsa_insn_arg_block *arg_start 3772 = new hsa_insn_arg_block (BRIG_KIND_DIRECTIVE_ARG_BLOCK_START, call_insn); 3773 hbb->append_insn (arg_start); 3774 3775 unsigned num_args = gimple_call_num_args (stmt); 3776 3777 /* Function arguments. */ 3778 for (unsigned i = 0; i < num_args; i++) 3779 { 3780 tree parm = gimple_call_arg (stmt, (int)i); 3781 hsa_op_with_type *src = hsa_reg_or_immed_for_gimple_op (parm, hbb); 3782 3783 hsa_op_address *addr = gen_hsa_addr_for_arg (TREE_TYPE (parm), i); 3784 hsa_insn_mem *mem = new hsa_insn_mem (BRIG_OPCODE_ST, src->m_type, 3785 src, addr); 3786 3787 call_insn->m_input_args.safe_push (addr->m_symbol); 3788 hbb->append_insn (mem); 3789 } 3790 3791 call_insn->m_args_code_list = new hsa_op_code_list (num_args); 3792 hbb->append_insn (call_insn); 3793 3794 /* Assign returned value. */ 3795 hsa_op_address *addr = gen_hsa_addr_for_arg (lhs_type, -1); 3796 3797 call_insn->m_output_arg = addr->m_symbol; 3798 call_insn->m_result_code_list = new hsa_op_code_list (1); 3799 3800 /* Argument block end. */ 3801 hsa_insn_arg_block *arg_end 3802 = new hsa_insn_arg_block (BRIG_KIND_DIRECTIVE_ARG_BLOCK_END, call_insn); 3803 hbb->append_insn (arg_end); 3804 } 3805 3806 /* Generate HSA instructions for a return value instruction. 3807 Instructions will be appended to HBB, which also needs to be the 3808 corresponding structure to the basic_block of STMT. */ 3809 3810 static void 3811 gen_hsa_insns_for_return (greturn *stmt, hsa_bb *hbb) 3812 { 3813 tree retval = gimple_return_retval (stmt); 3814 if (retval) 3815 { 3816 hsa_op_address *addr = new hsa_op_address (hsa_cfun->m_output_arg); 3817 3818 if (AGGREGATE_TYPE_P (TREE_TYPE (retval))) 3819 { 3820 BrigAlignment8_t align; 3821 hsa_op_address *retval_addr = gen_hsa_addr_with_align (retval, hbb, 3822 &align); 3823 gen_hsa_memory_copy (hbb, addr, retval_addr, 3824 hsa_cfun->m_output_arg->total_byte_size (), 3825 align); 3826 } 3827 else 3828 { 3829 BrigType16_t t = hsa_type_for_scalar_tree_type (TREE_TYPE (retval), 3830 false); 3831 BrigType16_t mtype = mem_type_for_type (t); 3832 3833 /* Store of return value. */ 3834 hsa_op_with_type *src = hsa_reg_or_immed_for_gimple_op (retval, hbb); 3835 src = src->get_in_type (mtype, hbb); 3836 hsa_insn_mem *mem = new hsa_insn_mem (BRIG_OPCODE_ST, mtype, src, 3837 addr); 3838 hbb->append_insn (mem); 3839 } 3840 } 3841 3842 /* HSAIL return instruction emission. */ 3843 hsa_insn_basic *ret = new hsa_insn_basic (0, BRIG_OPCODE_RET); 3844 hbb->append_insn (ret); 3845 } 3846 3847 /* Set OP_INDEX-th operand of the instruction to DEST, as the DEST 3848 can have a different type, conversion instructions are possibly 3849 appended to HBB. */ 3850 3851 void 3852 hsa_insn_basic::set_output_in_type (hsa_op_reg *dest, unsigned op_index, 3853 hsa_bb *hbb) 3854 { 3855 gcc_checking_assert (op_output_p (op_index)); 3856 3857 if (dest->m_type == m_type) 3858 { 3859 set_op (op_index, dest); 3860 return; 3861 } 3862 3863 hsa_insn_basic *insn; 3864 hsa_op_reg *tmp; 3865 if (hsa_needs_cvt (dest->m_type, m_type)) 3866 { 3867 tmp = new hsa_op_reg (m_type); 3868 insn = new hsa_insn_cvt (dest, tmp); 3869 } 3870 else if (hsa_type_bit_size (dest->m_type) == hsa_type_bit_size (m_type)) 3871 { 3872 /* When output, HSA registers do not really have types, only sizes, so if 3873 the sizes match, we can use the register directly. */ 3874 set_op (op_index, dest); 3875 return; 3876 } 3877 else 3878 { 3879 tmp = new hsa_op_reg (m_type); 3880 insn = new hsa_insn_basic (2, BRIG_OPCODE_MOV, dest->m_type, 3881 dest, tmp->get_in_type (dest->m_type, hbb)); 3882 hsa_fixup_mov_insn_type (insn); 3883 } 3884 set_op (op_index, tmp); 3885 hbb->append_insn (insn); 3886 } 3887 3888 /* Generate instruction OPCODE to query a property of HSA grid along the 3889 given DIMENSION. Store result into DEST and append the instruction to 3890 HBB. */ 3891 3892 static void 3893 query_hsa_grid_dim (hsa_op_reg *dest, int opcode, hsa_op_immed *dimension, 3894 hsa_bb *hbb) 3895 { 3896 hsa_insn_basic *insn = new hsa_insn_basic (2, opcode, BRIG_TYPE_U32, NULL, 3897 dimension); 3898 hbb->append_insn (insn); 3899 insn->set_output_in_type (dest, 0, hbb); 3900 } 3901 3902 /* Generate instruction OPCODE to query a property of HSA grid along the given 3903 dimension which is an immediate in first argument of STMT. Store result 3904 into the register corresponding to LHS of STMT and append the instruction to 3905 HBB. */ 3906 3907 static void 3908 query_hsa_grid_dim (gimple *stmt, int opcode, hsa_bb *hbb) 3909 { 3910 tree lhs = gimple_call_lhs (dyn_cast <gcall *> (stmt)); 3911 if (lhs == NULL_TREE) 3912 return; 3913 3914 tree arg = gimple_call_arg (stmt, 0); 3915 unsigned HOST_WIDE_INT dim = 5; 3916 if (tree_fits_uhwi_p (arg)) 3917 dim = tree_to_uhwi (arg); 3918 if (dim > 2) 3919 { 3920 HSA_SORRY_AT (gimple_location (stmt), 3921 "HSA grid query dimension must be immediate constant 0, 1 " 3922 "or 2"); 3923 return; 3924 } 3925 3926 hsa_op_immed *hdim = new hsa_op_immed (dim, (BrigKind16_t) BRIG_TYPE_U32); 3927 hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (lhs); 3928 query_hsa_grid_dim (dest, opcode, hdim, hbb); 3929 } 3930 3931 /* Generate instruction OPCODE to query a property of HSA grid that is 3932 independent of any dimension. Store result into the register corresponding 3933 to LHS of STMT and append the instruction to HBB. */ 3934 3935 static void 3936 query_hsa_grid_nodim (gimple *stmt, BrigOpcode16_t opcode, hsa_bb *hbb) 3937 { 3938 tree lhs = gimple_call_lhs (dyn_cast <gcall *> (stmt)); 3939 if (lhs == NULL_TREE) 3940 return; 3941 hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (lhs); 3942 BrigType16_t brig_type = hsa_unsigned_type_for_type (dest->m_type); 3943 hsa_insn_basic *insn = new hsa_insn_basic (1, opcode, brig_type, dest); 3944 hbb->append_insn (insn); 3945 } 3946 3947 /* Emit instructions that set hsa_num_threads according to provided VALUE. 3948 Instructions are appended to basic block HBB. */ 3949 3950 static void 3951 gen_set_num_threads (tree value, hsa_bb *hbb) 3952 { 3953 hbb->append_insn (new hsa_insn_comment ("omp_set_num_threads")); 3954 hsa_op_with_type *src = hsa_reg_or_immed_for_gimple_op (value, hbb); 3955 3956 src = src->get_in_type (hsa_num_threads->m_type, hbb); 3957 hsa_op_address *addr = new hsa_op_address (hsa_num_threads); 3958 3959 hsa_insn_basic *basic 3960 = new hsa_insn_mem (BRIG_OPCODE_ST, hsa_num_threads->m_type, src, addr); 3961 hbb->append_insn (basic); 3962 } 3963 3964 /* Return byte offset of a FIELD_NAME in GOMP_hsa_kernel_dispatch which 3965 is defined in plugin-hsa.c. */ 3966 3967 static HOST_WIDE_INT 3968 get_hsa_kernel_dispatch_offset (const char *field_name) 3969 { 3970 tree *hsa_kernel_dispatch_type = hsa_get_kernel_dispatch_type (); 3971 if (*hsa_kernel_dispatch_type == NULL) 3972 { 3973 /* Collection of information needed for a dispatch of a kernel from a 3974 kernel. Keep in sync with libgomp's plugin-hsa.c. */ 3975 3976 *hsa_kernel_dispatch_type = make_node (RECORD_TYPE); 3977 tree id_f1 = build_decl (BUILTINS_LOCATION, FIELD_DECL, 3978 get_identifier ("queue"), ptr_type_node); 3979 DECL_CHAIN (id_f1) = NULL_TREE; 3980 tree id_f2 = build_decl (BUILTINS_LOCATION, FIELD_DECL, 3981 get_identifier ("omp_data_memory"), 3982 ptr_type_node); 3983 DECL_CHAIN (id_f2) = id_f1; 3984 tree id_f3 = build_decl (BUILTINS_LOCATION, FIELD_DECL, 3985 get_identifier ("kernarg_address"), 3986 ptr_type_node); 3987 DECL_CHAIN (id_f3) = id_f2; 3988 tree id_f4 = build_decl (BUILTINS_LOCATION, FIELD_DECL, 3989 get_identifier ("object"), 3990 uint64_type_node); 3991 DECL_CHAIN (id_f4) = id_f3; 3992 tree id_f5 = build_decl (BUILTINS_LOCATION, FIELD_DECL, 3993 get_identifier ("signal"), 3994 uint64_type_node); 3995 DECL_CHAIN (id_f5) = id_f4; 3996 tree id_f6 = build_decl (BUILTINS_LOCATION, FIELD_DECL, 3997 get_identifier ("private_segment_size"), 3998 uint32_type_node); 3999 DECL_CHAIN (id_f6) = id_f5; 4000 tree id_f7 = build_decl (BUILTINS_LOCATION, FIELD_DECL, 4001 get_identifier ("group_segment_size"), 4002 uint32_type_node); 4003 DECL_CHAIN (id_f7) = id_f6; 4004 tree id_f8 = build_decl (BUILTINS_LOCATION, FIELD_DECL, 4005 get_identifier ("kernel_dispatch_count"), 4006 uint64_type_node); 4007 DECL_CHAIN (id_f8) = id_f7; 4008 tree id_f9 = build_decl (BUILTINS_LOCATION, FIELD_DECL, 4009 get_identifier ("debug"), 4010 uint64_type_node); 4011 DECL_CHAIN (id_f9) = id_f8; 4012 tree id_f10 = build_decl (BUILTINS_LOCATION, FIELD_DECL, 4013 get_identifier ("omp_level"), 4014 uint64_type_node); 4015 DECL_CHAIN (id_f10) = id_f9; 4016 tree id_f11 = build_decl (BUILTINS_LOCATION, FIELD_DECL, 4017 get_identifier ("children_dispatches"), 4018 ptr_type_node); 4019 DECL_CHAIN (id_f11) = id_f10; 4020 tree id_f12 = build_decl (BUILTINS_LOCATION, FIELD_DECL, 4021 get_identifier ("omp_num_threads"), 4022 uint32_type_node); 4023 DECL_CHAIN (id_f12) = id_f11; 4024 4025 4026 finish_builtin_struct (*hsa_kernel_dispatch_type, "__hsa_kernel_dispatch", 4027 id_f12, NULL_TREE); 4028 TYPE_ARTIFICIAL (*hsa_kernel_dispatch_type) = 1; 4029 } 4030 4031 for (tree chain = TYPE_FIELDS (*hsa_kernel_dispatch_type); 4032 chain != NULL_TREE; chain = TREE_CHAIN (chain)) 4033 if (id_equal (DECL_NAME (chain), field_name)) 4034 return int_byte_position (chain); 4035 4036 gcc_unreachable (); 4037 } 4038 4039 /* Return an HSA register that will contain number of threads for 4040 a future dispatched kernel. Instructions are added to HBB. */ 4041 4042 static hsa_op_reg * 4043 gen_num_threads_for_dispatch (hsa_bb *hbb) 4044 { 4045 /* Step 1) Assign to number of threads: 4046 MIN (HSA_DEFAULT_NUM_THREADS, hsa_num_threads). */ 4047 hsa_op_reg *threads = new hsa_op_reg (hsa_num_threads->m_type); 4048 hsa_op_address *addr = new hsa_op_address (hsa_num_threads); 4049 4050 hbb->append_insn (new hsa_insn_mem (BRIG_OPCODE_LD, threads->m_type, 4051 threads, addr)); 4052 4053 hsa_op_immed *limit = new hsa_op_immed (HSA_DEFAULT_NUM_THREADS, 4054 BRIG_TYPE_U32); 4055 hsa_op_reg *r = new hsa_op_reg (BRIG_TYPE_B1); 4056 hsa_insn_cmp * cmp 4057 = new hsa_insn_cmp (BRIG_COMPARE_LT, r->m_type, r, threads, limit); 4058 hbb->append_insn (cmp); 4059 4060 BrigType16_t btype = hsa_bittype_for_type (threads->m_type); 4061 hsa_op_reg *tmp = new hsa_op_reg (threads->m_type); 4062 4063 hbb->append_insn (new hsa_insn_basic (4, BRIG_OPCODE_CMOV, btype, tmp, r, 4064 threads, limit)); 4065 4066 /* Step 2) If the number is equal to zero, 4067 return shadow->omp_num_threads. */ 4068 hsa_op_reg *shadow_reg_ptr = hsa_cfun->get_shadow_reg (); 4069 4070 hsa_op_reg *shadow_thread_count = new hsa_op_reg (BRIG_TYPE_U32); 4071 addr 4072 = new hsa_op_address (shadow_reg_ptr, 4073 get_hsa_kernel_dispatch_offset ("omp_num_threads")); 4074 hsa_insn_basic *basic 4075 = new hsa_insn_mem (BRIG_OPCODE_LD, shadow_thread_count->m_type, 4076 shadow_thread_count, addr); 4077 hbb->append_insn (basic); 4078 4079 hsa_op_reg *tmp2 = new hsa_op_reg (threads->m_type); 4080 r = new hsa_op_reg (BRIG_TYPE_B1); 4081 hsa_op_immed *imm = new hsa_op_immed (0, shadow_thread_count->m_type); 4082 hbb->append_insn (new hsa_insn_cmp (BRIG_COMPARE_EQ, r->m_type, r, tmp, imm)); 4083 hbb->append_insn (new hsa_insn_basic (4, BRIG_OPCODE_CMOV, btype, tmp2, r, 4084 shadow_thread_count, tmp)); 4085 4086 hsa_op_base *dest = tmp2->get_in_type (BRIG_TYPE_U16, hbb); 4087 4088 return as_a <hsa_op_reg *> (dest); 4089 } 4090 4091 /* Build OPCODE query for all three hsa dimensions, multiply them and store the 4092 result into DEST. */ 4093 4094 static void 4095 multiply_grid_dim_characteristics (hsa_op_reg *dest, int opcode, hsa_bb *hbb) 4096 { 4097 hsa_op_reg *dimx = new hsa_op_reg (BRIG_TYPE_U32); 4098 query_hsa_grid_dim (dimx, opcode, 4099 new hsa_op_immed (0, (BrigKind16_t) BRIG_TYPE_U32), hbb); 4100 hsa_op_reg *dimy = new hsa_op_reg (BRIG_TYPE_U32); 4101 query_hsa_grid_dim (dimy, opcode, 4102 new hsa_op_immed (1, (BrigKind16_t) BRIG_TYPE_U32), hbb); 4103 hsa_op_reg *dimz = new hsa_op_reg (BRIG_TYPE_U32); 4104 query_hsa_grid_dim (dimz, opcode, 4105 new hsa_op_immed (2, (BrigKind16_t) BRIG_TYPE_U32), hbb); 4106 hsa_op_reg *tmp = new hsa_op_reg (dest->m_type); 4107 gen_hsa_binary_operation (BRIG_OPCODE_MUL, tmp, 4108 dimx->get_in_type (dest->m_type, hbb), 4109 dimy->get_in_type (dest->m_type, hbb), hbb); 4110 gen_hsa_binary_operation (BRIG_OPCODE_MUL, dest, tmp, 4111 dimz->get_in_type (dest->m_type, hbb), hbb); 4112 } 4113 4114 /* Emit instructions that assign number of threads to lhs of gimple STMT. 4115 Instructions are appended to basic block HBB. */ 4116 4117 static void 4118 gen_get_num_threads (gimple *stmt, hsa_bb *hbb) 4119 { 4120 if (gimple_call_lhs (stmt) == NULL_TREE) 4121 return; 4122 4123 hbb->append_insn (new hsa_insn_comment ("omp_get_num_threads")); 4124 tree lhs = gimple_call_lhs (stmt); 4125 hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (lhs); 4126 multiply_grid_dim_characteristics (dest, BRIG_OPCODE_CURRENTWORKGROUPSIZE, 4127 hbb); 4128 } 4129 4130 /* Emit instructions that assign number of teams to lhs of gimple STMT. 4131 Instructions are appended to basic block HBB. */ 4132 4133 static void 4134 gen_get_num_teams (gimple *stmt, hsa_bb *hbb) 4135 { 4136 if (gimple_call_lhs (stmt) == NULL_TREE) 4137 return; 4138 4139 hbb->append_insn (new hsa_insn_comment ("omp_get_num_teams")); 4140 tree lhs = gimple_call_lhs (stmt); 4141 hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (lhs); 4142 multiply_grid_dim_characteristics (dest, BRIG_OPCODE_GRIDGROUPS, hbb); 4143 } 4144 4145 /* Emit instructions that assign a team number to lhs of gimple STMT. 4146 Instructions are appended to basic block HBB. */ 4147 4148 static void 4149 gen_get_team_num (gimple *stmt, hsa_bb *hbb) 4150 { 4151 if (gimple_call_lhs (stmt) == NULL_TREE) 4152 return; 4153 4154 hbb->append_insn (new hsa_insn_comment ("omp_get_team_num")); 4155 tree lhs = gimple_call_lhs (stmt); 4156 hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (lhs); 4157 4158 hsa_op_reg *gnum_x = new hsa_op_reg (BRIG_TYPE_U32); 4159 query_hsa_grid_dim (gnum_x, BRIG_OPCODE_GRIDGROUPS, 4160 new hsa_op_immed (0, (BrigKind16_t) BRIG_TYPE_U32), hbb); 4161 hsa_op_reg *gnum_y = new hsa_op_reg (BRIG_TYPE_U32); 4162 query_hsa_grid_dim (gnum_y, BRIG_OPCODE_GRIDGROUPS, 4163 new hsa_op_immed (1, (BrigKind16_t) BRIG_TYPE_U32), hbb); 4164 4165 hsa_op_reg *gno_z = new hsa_op_reg (BRIG_TYPE_U32); 4166 query_hsa_grid_dim (gno_z, BRIG_OPCODE_WORKGROUPID, 4167 new hsa_op_immed (2, (BrigKind16_t) BRIG_TYPE_U32), hbb); 4168 4169 hsa_op_reg *tmp1 = new hsa_op_reg (dest->m_type); 4170 gen_hsa_binary_operation (BRIG_OPCODE_MUL, tmp1, 4171 gnum_x->get_in_type (dest->m_type, hbb), 4172 gnum_y->get_in_type (dest->m_type, hbb), hbb); 4173 hsa_op_reg *tmp2 = new hsa_op_reg (dest->m_type); 4174 gen_hsa_binary_operation (BRIG_OPCODE_MUL, tmp2, tmp1, 4175 gno_z->get_in_type (dest->m_type, hbb), hbb); 4176 4177 hsa_op_reg *gno_y = new hsa_op_reg (BRIG_TYPE_U32); 4178 query_hsa_grid_dim (gno_y, BRIG_OPCODE_WORKGROUPID, 4179 new hsa_op_immed (1, (BrigKind16_t) BRIG_TYPE_U32), hbb); 4180 hsa_op_reg *tmp3 = new hsa_op_reg (dest->m_type); 4181 gen_hsa_binary_operation (BRIG_OPCODE_MUL, tmp3, 4182 gnum_x->get_in_type (dest->m_type, hbb), 4183 gno_y->get_in_type (dest->m_type, hbb), hbb); 4184 hsa_op_reg *tmp4 = new hsa_op_reg (dest->m_type); 4185 gen_hsa_binary_operation (BRIG_OPCODE_ADD, tmp4, tmp3, tmp2, hbb); 4186 hsa_op_reg *gno_x = new hsa_op_reg (BRIG_TYPE_U32); 4187 query_hsa_grid_dim (gno_x, BRIG_OPCODE_WORKGROUPID, 4188 new hsa_op_immed (0, (BrigKind16_t) BRIG_TYPE_U32), hbb); 4189 gen_hsa_binary_operation (BRIG_OPCODE_ADD, dest, tmp4, 4190 gno_x->get_in_type (dest->m_type, hbb), hbb); 4191 } 4192 4193 /* Emit instructions that get levels-var ICV to lhs of gimple STMT. 4194 Instructions are appended to basic block HBB. */ 4195 4196 static void 4197 gen_get_level (gimple *stmt, hsa_bb *hbb) 4198 { 4199 if (gimple_call_lhs (stmt) == NULL_TREE) 4200 return; 4201 4202 hbb->append_insn (new hsa_insn_comment ("omp_get_level")); 4203 4204 tree lhs = gimple_call_lhs (stmt); 4205 hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (lhs); 4206 4207 hsa_op_reg *shadow_reg_ptr = hsa_cfun->get_shadow_reg (); 4208 if (shadow_reg_ptr == NULL) 4209 { 4210 HSA_SORRY_AT (gimple_location (stmt), 4211 "support for HSA does not implement omp_get_level called " 4212 "from a function not being inlined within a kernel"); 4213 return; 4214 } 4215 4216 hsa_op_address *addr 4217 = new hsa_op_address (shadow_reg_ptr, 4218 get_hsa_kernel_dispatch_offset ("omp_level")); 4219 4220 hsa_insn_mem *mem = new hsa_insn_mem (BRIG_OPCODE_LD, BRIG_TYPE_U64, 4221 (hsa_op_base *) NULL, addr); 4222 hbb->append_insn (mem); 4223 mem->set_output_in_type (dest, 0, hbb); 4224 } 4225 4226 /* Emit instruction that implement omp_get_max_threads of gimple STMT. */ 4227 4228 static void 4229 gen_get_max_threads (gimple *stmt, hsa_bb *hbb) 4230 { 4231 tree lhs = gimple_call_lhs (stmt); 4232 if (!lhs) 4233 return; 4234 4235 hbb->append_insn (new hsa_insn_comment ("omp_get_max_threads")); 4236 4237 hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (lhs); 4238 hsa_op_with_type *num_theads_reg = gen_num_threads_for_dispatch (hbb) 4239 ->get_in_type (dest->m_type, hbb); 4240 hsa_build_append_simple_mov (dest, num_theads_reg, hbb); 4241 } 4242 4243 /* Emit instructions that implement alloca builtin gimple STMT. 4244 Instructions are appended to basic block HBB. */ 4245 4246 static void 4247 gen_hsa_alloca (gcall *call, hsa_bb *hbb) 4248 { 4249 tree lhs = gimple_call_lhs (call); 4250 if (lhs == NULL_TREE) 4251 return; 4252 4253 built_in_function fn = DECL_FUNCTION_CODE (gimple_call_fndecl (call)); 4254 4255 gcc_checking_assert (ALLOCA_FUNCTION_CODE_P (fn)); 4256 4257 unsigned bit_alignment = 0; 4258 4259 if (fn != BUILT_IN_ALLOCA) 4260 { 4261 tree alignment_tree = gimple_call_arg (call, 1); 4262 if (TREE_CODE (alignment_tree) != INTEGER_CST) 4263 { 4264 HSA_SORRY_ATV (gimple_location (call), 4265 "support for HSA does not implement " 4266 "__builtin_alloca_with_align with a non-constant " 4267 "alignment: %E", alignment_tree); 4268 } 4269 4270 bit_alignment = tree_to_uhwi (alignment_tree); 4271 } 4272 4273 tree rhs1 = gimple_call_arg (call, 0); 4274 hsa_op_with_type *size = hsa_reg_or_immed_for_gimple_op (rhs1, hbb) 4275 ->get_in_type (BRIG_TYPE_U32, hbb); 4276 hsa_op_with_type *dest = hsa_cfun->reg_for_gimple_ssa (lhs); 4277 4278 hsa_op_reg *tmp 4279 = new hsa_op_reg (hsa_get_segment_addr_type (BRIG_SEGMENT_PRIVATE)); 4280 hsa_insn_alloca *a = new hsa_insn_alloca (tmp, size, bit_alignment); 4281 hbb->append_insn (a); 4282 4283 hsa_insn_seg *seg 4284 = new hsa_insn_seg (BRIG_OPCODE_STOF, 4285 hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT), 4286 tmp->m_type, BRIG_SEGMENT_PRIVATE, dest, tmp); 4287 hbb->append_insn (seg); 4288 } 4289 4290 /* Emit instructions that implement clrsb builtin STMT: 4291 Returns the number of leading redundant sign bits in x, i.e. the number 4292 of bits following the most significant bit that are identical to it. 4293 There are no special cases for 0 or other values. 4294 Instructions are appended to basic block HBB. */ 4295 4296 static void 4297 gen_hsa_clrsb (gcall *call, hsa_bb *hbb) 4298 { 4299 tree lhs = gimple_call_lhs (call); 4300 if (lhs == NULL_TREE) 4301 return; 4302 4303 hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (lhs); 4304 tree rhs1 = gimple_call_arg (call, 0); 4305 hsa_op_with_type *arg = hsa_reg_or_immed_for_gimple_op (rhs1, hbb); 4306 arg->extend_int_to_32bit (hbb); 4307 BrigType16_t bittype = hsa_bittype_for_type (arg->m_type); 4308 unsigned bitsize = tree_to_uhwi (TYPE_SIZE (TREE_TYPE (rhs1))); 4309 4310 /* FIRSTBIT instruction is defined just for 32 and 64-bits wide integers. */ 4311 gcc_checking_assert (bitsize == 32 || bitsize == 64); 4312 4313 /* Set true to MOST_SIG if the most significant bit is set to one. */ 4314 hsa_op_immed *c = new hsa_op_immed (1ul << (bitsize - 1), 4315 hsa_uint_for_bitsize (bitsize)); 4316 4317 hsa_op_reg *and_reg = new hsa_op_reg (bittype); 4318 gen_hsa_binary_operation (BRIG_OPCODE_AND, and_reg, arg, c, hbb); 4319 4320 hsa_op_reg *most_sign = new hsa_op_reg (BRIG_TYPE_B1); 4321 hsa_insn_cmp *cmp 4322 = new hsa_insn_cmp (BRIG_COMPARE_EQ, most_sign->m_type, most_sign, 4323 and_reg, c); 4324 hbb->append_insn (cmp); 4325 4326 /* If the most significant bit is one, negate the input. Otherwise 4327 shift the input value to left by one bit. */ 4328 hsa_op_reg *arg_neg = new hsa_op_reg (arg->m_type); 4329 gen_hsa_unary_operation (BRIG_OPCODE_NEG, arg_neg, arg, hbb); 4330 4331 hsa_op_reg *shifted_arg = new hsa_op_reg (arg->m_type); 4332 gen_hsa_binary_operation (BRIG_OPCODE_SHL, shifted_arg, arg, 4333 new hsa_op_immed (1, BRIG_TYPE_U64), hbb); 4334 4335 /* Assign the value that can be used for FIRSTBIT instruction according 4336 to the most significant bit. */ 4337 hsa_op_reg *tmp = new hsa_op_reg (bittype); 4338 hsa_insn_basic *cmov 4339 = new hsa_insn_basic (4, BRIG_OPCODE_CMOV, bittype, tmp, most_sign, 4340 arg_neg, shifted_arg); 4341 hbb->append_insn (cmov); 4342 4343 hsa_op_reg *leading_bits = new hsa_op_reg (BRIG_TYPE_S32); 4344 gen_hsa_unary_operation (BRIG_OPCODE_FIRSTBIT, leading_bits, 4345 tmp->get_in_type (hsa_uint_for_bitsize (bitsize), 4346 hbb), hbb); 4347 4348 /* Set flag if the input value is equal to zero. */ 4349 hsa_op_reg *is_zero = new hsa_op_reg (BRIG_TYPE_B1); 4350 cmp = new hsa_insn_cmp (BRIG_COMPARE_EQ, is_zero->m_type, is_zero, arg, 4351 new hsa_op_immed (0, arg->m_type)); 4352 hbb->append_insn (cmp); 4353 4354 /* Return the number of leading bits, 4355 or (bitsize - 1) if the input value is zero. */ 4356 cmov = new hsa_insn_basic (4, BRIG_OPCODE_CMOV, BRIG_TYPE_B32, NULL, is_zero, 4357 new hsa_op_immed (bitsize - 1, BRIG_TYPE_U32), 4358 leading_bits->get_in_type (BRIG_TYPE_B32, hbb)); 4359 hbb->append_insn (cmov); 4360 cmov->set_output_in_type (dest, 0, hbb); 4361 } 4362 4363 /* Emit instructions that implement ffs builtin STMT: 4364 Returns one plus the index of the least significant 1-bit of x, 4365 or if x is zero, returns zero. 4366 Instructions are appended to basic block HBB. */ 4367 4368 static void 4369 gen_hsa_ffs (gcall *call, hsa_bb *hbb) 4370 { 4371 tree lhs = gimple_call_lhs (call); 4372 if (lhs == NULL_TREE) 4373 return; 4374 4375 hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (lhs); 4376 4377 tree rhs1 = gimple_call_arg (call, 0); 4378 hsa_op_with_type *arg = hsa_reg_or_immed_for_gimple_op (rhs1, hbb); 4379 arg = arg->extend_int_to_32bit (hbb); 4380 4381 hsa_op_reg *tmp = new hsa_op_reg (BRIG_TYPE_U32); 4382 hsa_insn_srctype *insn = new hsa_insn_srctype (2, BRIG_OPCODE_LASTBIT, 4383 tmp->m_type, arg->m_type, 4384 tmp, arg); 4385 hbb->append_insn (insn); 4386 4387 hsa_insn_basic *addition 4388 = new hsa_insn_basic (3, BRIG_OPCODE_ADD, tmp->m_type, NULL, tmp, 4389 new hsa_op_immed (1, tmp->m_type)); 4390 hbb->append_insn (addition); 4391 addition->set_output_in_type (dest, 0, hbb); 4392 } 4393 4394 static void 4395 gen_hsa_popcount_to_dest (hsa_op_reg *dest, hsa_op_with_type *arg, hsa_bb *hbb) 4396 { 4397 gcc_checking_assert (hsa_type_integer_p (arg->m_type)); 4398 4399 if (hsa_type_bit_size (arg->m_type) < 32) 4400 arg = arg->get_in_type (BRIG_TYPE_B32, hbb); 4401 4402 BrigType16_t srctype = hsa_bittype_for_type (arg->m_type); 4403 if (!hsa_btype_p (arg->m_type)) 4404 arg = arg->get_in_type (srctype, hbb); 4405 4406 hsa_insn_srctype *popcount 4407 = new hsa_insn_srctype (2, BRIG_OPCODE_POPCOUNT, BRIG_TYPE_U32, 4408 srctype, NULL, arg); 4409 hbb->append_insn (popcount); 4410 popcount->set_output_in_type (dest, 0, hbb); 4411 } 4412 4413 /* Emit instructions that implement parity builtin STMT: 4414 Returns the parity of x, i.e. the number of 1-bits in x modulo 2. 4415 Instructions are appended to basic block HBB. */ 4416 4417 static void 4418 gen_hsa_parity (gcall *call, hsa_bb *hbb) 4419 { 4420 tree lhs = gimple_call_lhs (call); 4421 if (lhs == NULL_TREE) 4422 return; 4423 4424 hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (lhs); 4425 tree rhs1 = gimple_call_arg (call, 0); 4426 hsa_op_with_type *arg = hsa_reg_or_immed_for_gimple_op (rhs1, hbb); 4427 4428 hsa_op_reg *popcount = new hsa_op_reg (BRIG_TYPE_U32); 4429 gen_hsa_popcount_to_dest (popcount, arg, hbb); 4430 4431 hsa_insn_basic *insn 4432 = new hsa_insn_basic (3, BRIG_OPCODE_REM, popcount->m_type, NULL, popcount, 4433 new hsa_op_immed (2, popcount->m_type)); 4434 hbb->append_insn (insn); 4435 insn->set_output_in_type (dest, 0, hbb); 4436 } 4437 4438 /* Emit instructions that implement popcount builtin STMT. 4439 Instructions are appended to basic block HBB. */ 4440 4441 static void 4442 gen_hsa_popcount (gcall *call, hsa_bb *hbb) 4443 { 4444 tree lhs = gimple_call_lhs (call); 4445 if (lhs == NULL_TREE) 4446 return; 4447 4448 hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (lhs); 4449 tree rhs1 = gimple_call_arg (call, 0); 4450 hsa_op_with_type *arg = hsa_reg_or_immed_for_gimple_op (rhs1, hbb); 4451 4452 gen_hsa_popcount_to_dest (dest, arg, hbb); 4453 } 4454 4455 /* Emit instructions that implement DIVMOD builtin STMT. 4456 Instructions are appended to basic block HBB. */ 4457 4458 static void 4459 gen_hsa_divmod (gcall *call, hsa_bb *hbb) 4460 { 4461 tree lhs = gimple_call_lhs (call); 4462 if (lhs == NULL_TREE) 4463 return; 4464 4465 tree rhs0 = gimple_call_arg (call, 0); 4466 tree rhs1 = gimple_call_arg (call, 1); 4467 4468 hsa_op_with_type *arg0 = hsa_reg_or_immed_for_gimple_op (rhs0, hbb); 4469 arg0 = arg0->extend_int_to_32bit (hbb); 4470 hsa_op_with_type *arg1 = hsa_reg_or_immed_for_gimple_op (rhs1, hbb); 4471 arg1 = arg1->extend_int_to_32bit (hbb); 4472 4473 hsa_op_reg *dest0 = new hsa_op_reg (arg0->m_type); 4474 hsa_op_reg *dest1 = new hsa_op_reg (arg1->m_type); 4475 4476 hsa_insn_basic *insn = new hsa_insn_basic (3, BRIG_OPCODE_DIV, dest0->m_type, 4477 dest0, arg0, arg1); 4478 hbb->append_insn (insn); 4479 insn = new hsa_insn_basic (3, BRIG_OPCODE_REM, dest1->m_type, dest1, arg0, 4480 arg1); 4481 hbb->append_insn (insn); 4482 4483 hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (lhs); 4484 BrigType16_t dst_type = hsa_extend_inttype_to_32bit (dest->m_type); 4485 BrigType16_t src_type = hsa_bittype_for_type (dest0->m_type); 4486 4487 insn = new hsa_insn_packed (3, BRIG_OPCODE_COMBINE, dst_type, 4488 src_type, NULL, dest0, dest1); 4489 hbb->append_insn (insn); 4490 insn->set_output_in_type (dest, 0, hbb); 4491 } 4492 4493 /* Set VALUE to a shadow kernel debug argument and append a new instruction 4494 to HBB basic block. */ 4495 4496 static void 4497 set_debug_value (hsa_bb *hbb, hsa_op_with_type *value) 4498 { 4499 hsa_op_reg *shadow_reg_ptr = hsa_cfun->get_shadow_reg (); 4500 if (shadow_reg_ptr == NULL) 4501 return; 4502 4503 hsa_op_address *addr 4504 = new hsa_op_address (shadow_reg_ptr, 4505 get_hsa_kernel_dispatch_offset ("debug")); 4506 hsa_insn_mem *mem = new hsa_insn_mem (BRIG_OPCODE_ST, BRIG_TYPE_U64, value, 4507 addr); 4508 hbb->append_insn (mem); 4509 } 4510 4511 void 4512 omp_simple_builtin::generate (gimple *stmt, hsa_bb *hbb) 4513 { 4514 if (m_sorry) 4515 { 4516 if (m_warning_message) 4517 HSA_SORRY_AT (gimple_location (stmt), m_warning_message); 4518 else 4519 HSA_SORRY_ATV (gimple_location (stmt), 4520 "Support for HSA does not implement calls to %s\n", 4521 m_name); 4522 } 4523 else if (m_warning_message != NULL) 4524 warning_at (gimple_location (stmt), OPT_Whsa, m_warning_message); 4525 4526 if (m_return_value != NULL) 4527 { 4528 tree lhs = gimple_call_lhs (stmt); 4529 if (!lhs) 4530 return; 4531 4532 hbb->append_insn (new hsa_insn_comment (m_name)); 4533 4534 hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (lhs); 4535 hsa_op_with_type *op = m_return_value->get_in_type (dest->m_type, hbb); 4536 hsa_build_append_simple_mov (dest, op, hbb); 4537 } 4538 } 4539 4540 /* If STMT is a call of a known library function, generate code to perform 4541 it and return true. */ 4542 4543 static bool 4544 gen_hsa_insns_for_known_library_call (gimple *stmt, hsa_bb *hbb) 4545 { 4546 bool handled = false; 4547 const char *name = hsa_get_declaration_name (gimple_call_fndecl (stmt)); 4548 4549 char *copy = NULL; 4550 size_t len = strlen (name); 4551 if (len > 0 && name[len - 1] == '_') 4552 { 4553 copy = XNEWVEC (char, len + 1); 4554 strcpy (copy, name); 4555 copy[len - 1] = '\0'; 4556 name = copy; 4557 } 4558 4559 /* Handle omp_* routines. */ 4560 if (strstr (name, "omp_") == name) 4561 { 4562 hsa_init_simple_builtins (); 4563 omp_simple_builtin *builtin = omp_simple_builtins->get (name); 4564 if (builtin) 4565 { 4566 builtin->generate (stmt, hbb); 4567 return true; 4568 } 4569 4570 handled = true; 4571 if (strcmp (name, "omp_set_num_threads") == 0) 4572 gen_set_num_threads (gimple_call_arg (stmt, 0), hbb); 4573 else if (strcmp (name, "omp_get_thread_num") == 0) 4574 { 4575 hbb->append_insn (new hsa_insn_comment (name)); 4576 query_hsa_grid_nodim (stmt, BRIG_OPCODE_WORKITEMFLATABSID, hbb); 4577 } 4578 else if (strcmp (name, "omp_get_num_threads") == 0) 4579 { 4580 hbb->append_insn (new hsa_insn_comment (name)); 4581 gen_get_num_threads (stmt, hbb); 4582 } 4583 else if (strcmp (name, "omp_get_num_teams") == 0) 4584 gen_get_num_teams (stmt, hbb); 4585 else if (strcmp (name, "omp_get_team_num") == 0) 4586 gen_get_team_num (stmt, hbb); 4587 else if (strcmp (name, "omp_get_level") == 0) 4588 gen_get_level (stmt, hbb); 4589 else if (strcmp (name, "omp_get_active_level") == 0) 4590 gen_get_level (stmt, hbb); 4591 else if (strcmp (name, "omp_in_parallel") == 0) 4592 gen_get_level (stmt, hbb); 4593 else if (strcmp (name, "omp_get_max_threads") == 0) 4594 gen_get_max_threads (stmt, hbb); 4595 else 4596 handled = false; 4597 4598 if (handled) 4599 { 4600 if (copy) 4601 free (copy); 4602 return true; 4603 } 4604 } 4605 4606 if (strcmp (name, "__hsa_set_debug_value") == 0) 4607 { 4608 handled = true; 4609 if (hsa_cfun->has_shadow_reg_p ()) 4610 { 4611 tree rhs1 = gimple_call_arg (stmt, 0); 4612 hsa_op_with_type *src = hsa_reg_or_immed_for_gimple_op (rhs1, hbb); 4613 4614 src = src->get_in_type (BRIG_TYPE_U64, hbb); 4615 set_debug_value (hbb, src); 4616 } 4617 } 4618 4619 if (copy) 4620 free (copy); 4621 return handled; 4622 } 4623 4624 /* Helper functions to create a single unary HSA operations out of calls to 4625 builtins. OPCODE is the HSA operation to be generated. STMT is a gimple 4626 call to a builtin. HBB is the HSA BB to which the instruction should be 4627 added. Note that nothing will be created if STMT does not have a LHS. */ 4628 4629 static void 4630 gen_hsa_unaryop_for_builtin (BrigOpcode opcode, gimple *stmt, hsa_bb *hbb) 4631 { 4632 tree lhs = gimple_call_lhs (stmt); 4633 if (!lhs) 4634 return; 4635 hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (lhs); 4636 hsa_op_with_type *op 4637 = hsa_reg_or_immed_for_gimple_op (gimple_call_arg (stmt, 0), hbb); 4638 gen_hsa_unary_operation (opcode, dest, op, hbb); 4639 } 4640 4641 /* Helper functions to create a call to standard library if LHS of the 4642 STMT is used. HBB is the HSA BB to which the instruction should be 4643 added. */ 4644 4645 static void 4646 gen_hsa_unaryop_builtin_call (gimple *stmt, hsa_bb *hbb) 4647 { 4648 tree lhs = gimple_call_lhs (stmt); 4649 if (!lhs) 4650 return; 4651 4652 if (gimple_call_internal_p (stmt)) 4653 gen_hsa_insns_for_call_of_internal_fn (stmt, hbb); 4654 else 4655 gen_hsa_insns_for_direct_call (stmt, hbb); 4656 } 4657 4658 /* Helper functions to create a single unary HSA operations out of calls to 4659 builtins (if unsafe math optimizations are enable). Otherwise, create 4660 a call to standard library function. 4661 OPCODE is the HSA operation to be generated. STMT is a gimple 4662 call to a builtin. HBB is the HSA BB to which the instruction should be 4663 added. Note that nothing will be created if STMT does not have a LHS. */ 4664 4665 static void 4666 gen_hsa_unaryop_or_call_for_builtin (BrigOpcode opcode, gimple *stmt, 4667 hsa_bb *hbb) 4668 { 4669 if (flag_unsafe_math_optimizations) 4670 gen_hsa_unaryop_for_builtin (opcode, stmt, hbb); 4671 else 4672 gen_hsa_unaryop_builtin_call (stmt, hbb); 4673 } 4674 4675 /* Generate HSA address corresponding to a value VAL (as opposed to a memory 4676 reference tree), for example an SSA_NAME or an ADDR_EXPR. HBB is the HSA BB 4677 to which the instruction should be added. */ 4678 4679 static hsa_op_address * 4680 get_address_from_value (tree val, hsa_bb *hbb) 4681 { 4682 switch (TREE_CODE (val)) 4683 { 4684 case SSA_NAME: 4685 { 4686 BrigType16_t addrtype = hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT); 4687 hsa_op_base *reg 4688 = hsa_cfun->reg_for_gimple_ssa (val)->get_in_type (addrtype, hbb); 4689 return new hsa_op_address (NULL, as_a <hsa_op_reg *> (reg), 0); 4690 } 4691 case ADDR_EXPR: 4692 return gen_hsa_addr (TREE_OPERAND (val, 0), hbb); 4693 4694 case INTEGER_CST: 4695 if (tree_fits_shwi_p (val)) 4696 return new hsa_op_address (NULL, NULL, tree_to_shwi (val)); 4697 /* fall-through */ 4698 4699 default: 4700 HSA_SORRY_ATV (EXPR_LOCATION (val), 4701 "support for HSA does not implement memory access to %E", 4702 val); 4703 return new hsa_op_address (NULL, NULL, 0); 4704 } 4705 } 4706 4707 /* Expand assignment of a result of a string BUILTIN to DST. 4708 Size of the operation is N bytes, where instructions 4709 will be append to HBB. */ 4710 4711 static void 4712 expand_lhs_of_string_op (gimple *stmt, 4713 unsigned HOST_WIDE_INT n, hsa_bb *hbb, 4714 enum built_in_function builtin) 4715 { 4716 /* If LHS is expected, we need to emit a PHI instruction. */ 4717 tree lhs = gimple_call_lhs (stmt); 4718 if (!lhs) 4719 return; 4720 4721 hsa_op_reg *lhs_reg = hsa_cfun->reg_for_gimple_ssa (lhs); 4722 4723 hsa_op_with_type *dst_reg 4724 = hsa_reg_or_immed_for_gimple_op (gimple_call_arg (stmt, 0), hbb); 4725 hsa_op_with_type *tmp; 4726 4727 switch (builtin) 4728 { 4729 case BUILT_IN_MEMPCPY: 4730 { 4731 tmp = new hsa_op_reg (dst_reg->m_type); 4732 hsa_insn_basic *add 4733 = new hsa_insn_basic (3, BRIG_OPCODE_ADD, tmp->m_type, 4734 tmp, dst_reg, 4735 new hsa_op_immed (n, dst_reg->m_type)); 4736 hbb->append_insn (add); 4737 break; 4738 } 4739 case BUILT_IN_MEMCPY: 4740 case BUILT_IN_MEMSET: 4741 tmp = dst_reg; 4742 break; 4743 default: 4744 gcc_unreachable (); 4745 } 4746 4747 hbb->append_insn (new hsa_insn_basic (2, BRIG_OPCODE_MOV, lhs_reg->m_type, 4748 lhs_reg, tmp)); 4749 } 4750 4751 #define HSA_MEMORY_BUILTINS_LIMIT 128 4752 4753 /* Expand a string builtin (from a gimple STMT) in a way that 4754 according to MISALIGNED_FLAG we process either direct emission 4755 (a bunch of memory load and store instructions), or we emit a function call 4756 of a library function (for instance 'memcpy'). Actually, a basic block 4757 for direct emission is just prepared, where caller is responsible 4758 for emission of corresponding instructions. 4759 All instruction are appended to HBB. */ 4760 4761 hsa_bb * 4762 expand_string_operation_builtin (gimple *stmt, hsa_bb *hbb, 4763 hsa_op_reg *misaligned_flag) 4764 { 4765 edge e = split_block (hbb->m_bb, stmt); 4766 basic_block condition_bb = e->src; 4767 hbb->append_insn (new hsa_insn_cbr (misaligned_flag)); 4768 4769 /* Prepare the control flow. */ 4770 edge condition_edge = EDGE_SUCC (condition_bb, 0); 4771 basic_block call_bb = split_edge (condition_edge); 4772 4773 basic_block expanded_bb = split_edge (EDGE_SUCC (call_bb, 0)); 4774 basic_block cont_bb = EDGE_SUCC (expanded_bb, 0)->dest; 4775 basic_block merge_bb = split_edge (EDGE_PRED (cont_bb, 0)); 4776 4777 condition_edge->flags &= ~EDGE_FALLTHRU; 4778 condition_edge->flags |= EDGE_TRUE_VALUE; 4779 make_edge (condition_bb, expanded_bb, EDGE_FALSE_VALUE); 4780 4781 redirect_edge_succ (EDGE_SUCC (call_bb, 0), merge_bb); 4782 4783 hsa_cfun->m_modified_cfg = true; 4784 4785 hsa_init_new_bb (expanded_bb); 4786 4787 /* Slow path: function call. */ 4788 gen_hsa_insns_for_direct_call (stmt, hsa_init_new_bb (call_bb), false); 4789 4790 return hsa_bb_for_bb (expanded_bb); 4791 } 4792 4793 /* Expand a memory copy BUILTIN (BUILT_IN_MEMCPY, BUILT_IN_MEMPCPY) from 4794 a gimple STMT and store all necessary instruction to HBB basic block. */ 4795 4796 static void 4797 expand_memory_copy (gimple *stmt, hsa_bb *hbb, enum built_in_function builtin) 4798 { 4799 tree byte_size = gimple_call_arg (stmt, 2); 4800 4801 if (!tree_fits_uhwi_p (byte_size)) 4802 { 4803 gen_hsa_insns_for_direct_call (stmt, hbb); 4804 return; 4805 } 4806 4807 unsigned HOST_WIDE_INT n = tree_to_uhwi (byte_size); 4808 4809 if (n > HSA_MEMORY_BUILTINS_LIMIT) 4810 { 4811 gen_hsa_insns_for_direct_call (stmt, hbb); 4812 return; 4813 } 4814 4815 tree dst = gimple_call_arg (stmt, 0); 4816 tree src = gimple_call_arg (stmt, 1); 4817 4818 hsa_op_address *dst_addr = get_address_from_value (dst, hbb); 4819 hsa_op_address *src_addr = get_address_from_value (src, hbb); 4820 4821 /* As gen_hsa_memory_copy relies on memory alignment 4822 greater or equal to 8 bytes, we need to verify the alignment. */ 4823 BrigType16_t addrtype = hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT); 4824 hsa_op_reg *src_addr_reg = new hsa_op_reg (addrtype); 4825 hsa_op_reg *dst_addr_reg = new hsa_op_reg (addrtype); 4826 4827 convert_addr_to_flat_segment (src_addr, src_addr_reg, hbb); 4828 convert_addr_to_flat_segment (dst_addr, dst_addr_reg, hbb); 4829 4830 /* Process BIT OR for source and destination addresses. */ 4831 hsa_op_reg *or_reg = new hsa_op_reg (addrtype); 4832 gen_hsa_binary_operation (BRIG_OPCODE_OR, or_reg, src_addr_reg, 4833 dst_addr_reg, hbb); 4834 4835 /* Process BIT AND with 0x7 to identify the desired alignment 4836 of 8 bytes. */ 4837 hsa_op_reg *masked = new hsa_op_reg (addrtype); 4838 4839 gen_hsa_binary_operation (BRIG_OPCODE_AND, masked, or_reg, 4840 new hsa_op_immed (7, addrtype), hbb); 4841 4842 hsa_op_reg *misaligned = new hsa_op_reg (BRIG_TYPE_B1); 4843 hbb->append_insn (new hsa_insn_cmp (BRIG_COMPARE_NE, misaligned->m_type, 4844 misaligned, masked, 4845 new hsa_op_immed (0, masked->m_type))); 4846 4847 hsa_bb *native_impl_bb 4848 = expand_string_operation_builtin (stmt, hbb, misaligned); 4849 4850 gen_hsa_memory_copy (native_impl_bb, dst_addr, src_addr, n, BRIG_ALIGNMENT_8); 4851 hsa_bb *merge_bb 4852 = hsa_init_new_bb (EDGE_SUCC (native_impl_bb->m_bb, 0)->dest); 4853 expand_lhs_of_string_op (stmt, n, merge_bb, builtin); 4854 } 4855 4856 4857 /* Expand a memory set BUILTIN (BUILT_IN_MEMSET, BUILT_IN_BZERO) from 4858 a gimple STMT and store all necessary instruction to HBB basic block. 4859 The operation set N bytes with a CONSTANT value. */ 4860 4861 static void 4862 expand_memory_set (gimple *stmt, unsigned HOST_WIDE_INT n, 4863 unsigned HOST_WIDE_INT constant, hsa_bb *hbb, 4864 enum built_in_function builtin) 4865 { 4866 tree dst = gimple_call_arg (stmt, 0); 4867 hsa_op_address *dst_addr = get_address_from_value (dst, hbb); 4868 4869 /* As gen_hsa_memory_set relies on memory alignment 4870 greater or equal to 8 bytes, we need to verify the alignment. */ 4871 BrigType16_t addrtype = hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT); 4872 hsa_op_reg *dst_addr_reg = new hsa_op_reg (addrtype); 4873 convert_addr_to_flat_segment (dst_addr, dst_addr_reg, hbb); 4874 4875 /* Process BIT AND with 0x7 to identify the desired alignment 4876 of 8 bytes. */ 4877 hsa_op_reg *masked = new hsa_op_reg (addrtype); 4878 4879 gen_hsa_binary_operation (BRIG_OPCODE_AND, masked, dst_addr_reg, 4880 new hsa_op_immed (7, addrtype), hbb); 4881 4882 hsa_op_reg *misaligned = new hsa_op_reg (BRIG_TYPE_B1); 4883 hbb->append_insn (new hsa_insn_cmp (BRIG_COMPARE_NE, misaligned->m_type, 4884 misaligned, masked, 4885 new hsa_op_immed (0, masked->m_type))); 4886 4887 hsa_bb *native_impl_bb 4888 = expand_string_operation_builtin (stmt, hbb, misaligned); 4889 4890 gen_hsa_memory_set (native_impl_bb, dst_addr, constant, n, BRIG_ALIGNMENT_8); 4891 hsa_bb *merge_bb 4892 = hsa_init_new_bb (EDGE_SUCC (native_impl_bb->m_bb, 0)->dest); 4893 expand_lhs_of_string_op (stmt, n, merge_bb, builtin); 4894 } 4895 4896 /* Store into MEMORDER the memory order specified by tree T, which must be an 4897 integer constant representing a C++ memory order. If it isn't, issue an HSA 4898 sorry message using LOC and return true, otherwise return false and store 4899 the name of the requested order to *MNAME. */ 4900 4901 static bool 4902 hsa_memorder_from_tree (tree t, BrigMemoryOrder *memorder, const char **mname, 4903 location_t loc) 4904 { 4905 if (!tree_fits_uhwi_p (t)) 4906 { 4907 HSA_SORRY_ATV (loc, "support for HSA does not implement memory model %E", 4908 t); 4909 return true; 4910 } 4911 4912 unsigned HOST_WIDE_INT mm = tree_to_uhwi (t); 4913 switch (mm & MEMMODEL_BASE_MASK) 4914 { 4915 case MEMMODEL_RELAXED: 4916 *memorder = BRIG_MEMORY_ORDER_RELAXED; 4917 *mname = "relaxed"; 4918 break; 4919 case MEMMODEL_CONSUME: 4920 /* HSA does not have an equivalent, but we can use the slightly stronger 4921 ACQUIRE. */ 4922 *memorder = BRIG_MEMORY_ORDER_SC_ACQUIRE; 4923 *mname = "consume"; 4924 break; 4925 case MEMMODEL_ACQUIRE: 4926 *memorder = BRIG_MEMORY_ORDER_SC_ACQUIRE; 4927 *mname = "acquire"; 4928 break; 4929 case MEMMODEL_RELEASE: 4930 *memorder = BRIG_MEMORY_ORDER_SC_RELEASE; 4931 *mname = "release"; 4932 break; 4933 case MEMMODEL_ACQ_REL: 4934 *memorder = BRIG_MEMORY_ORDER_SC_ACQUIRE_RELEASE; 4935 *mname = "acq_rel"; 4936 break; 4937 case MEMMODEL_SEQ_CST: 4938 /* Callers implementing a simple load or store need to remove the release 4939 or acquire part respectively. */ 4940 *memorder = BRIG_MEMORY_ORDER_SC_ACQUIRE_RELEASE; 4941 *mname = "seq_cst"; 4942 break; 4943 default: 4944 { 4945 HSA_SORRY_AT (loc, "support for HSA does not implement the specified " 4946 "memory model"); 4947 return true; 4948 } 4949 } 4950 return false; 4951 } 4952 4953 /* Helper function to create an HSA atomic operation instruction out of calls 4954 to atomic builtins. RET_ORIG is true if the built-in is the variant that 4955 return s the value before applying operation, and false if it should return 4956 the value after applying the operation (if it returns value at all). ACODE 4957 is the atomic operation code, STMT is a gimple call to a builtin. HBB is 4958 the HSA BB to which the instruction should be added. If SIGNAL is true, the 4959 created operation will work on HSA signals rather than atomic variables. */ 4960 4961 static void 4962 gen_hsa_atomic_for_builtin (bool ret_orig, enum BrigAtomicOperation acode, 4963 gimple *stmt, hsa_bb *hbb, bool signal) 4964 { 4965 tree lhs = gimple_call_lhs (stmt); 4966 4967 tree type = TREE_TYPE (gimple_call_arg (stmt, 1)); 4968 BrigType16_t hsa_type = hsa_type_for_scalar_tree_type (type, false); 4969 BrigType16_t mtype = mem_type_for_type (hsa_type); 4970 BrigMemoryOrder memorder; 4971 const char *mmname; 4972 4973 if (hsa_memorder_from_tree (gimple_call_arg (stmt, 2), &memorder, &mmname, 4974 gimple_location (stmt))) 4975 return; 4976 4977 /* Certain atomic insns must have Bx memory types. */ 4978 switch (acode) 4979 { 4980 case BRIG_ATOMIC_LD: 4981 case BRIG_ATOMIC_ST: 4982 case BRIG_ATOMIC_AND: 4983 case BRIG_ATOMIC_OR: 4984 case BRIG_ATOMIC_XOR: 4985 case BRIG_ATOMIC_EXCH: 4986 mtype = hsa_bittype_for_type (mtype); 4987 break; 4988 default: 4989 break; 4990 } 4991 4992 hsa_op_reg *dest; 4993 int nops, opcode; 4994 if (lhs) 4995 { 4996 if (ret_orig) 4997 dest = hsa_cfun->reg_for_gimple_ssa (lhs); 4998 else 4999 dest = new hsa_op_reg (hsa_type); 5000 opcode = signal ? BRIG_OPCODE_SIGNAL : BRIG_OPCODE_ATOMIC; 5001 nops = 3; 5002 } 5003 else 5004 { 5005 dest = NULL; 5006 opcode = signal ? BRIG_OPCODE_SIGNALNORET : BRIG_OPCODE_ATOMICNORET; 5007 nops = 2; 5008 } 5009 5010 if (acode == BRIG_ATOMIC_ST) 5011 { 5012 if (memorder == BRIG_MEMORY_ORDER_SC_ACQUIRE_RELEASE) 5013 memorder = BRIG_MEMORY_ORDER_SC_RELEASE; 5014 5015 if (memorder != BRIG_MEMORY_ORDER_RELAXED 5016 && memorder != BRIG_MEMORY_ORDER_SC_RELEASE 5017 && memorder != BRIG_MEMORY_ORDER_NONE) 5018 { 5019 HSA_SORRY_ATV (gimple_location (stmt), 5020 "support for HSA does not implement memory model for " 5021 "ATOMIC_ST: %s", mmname); 5022 return; 5023 } 5024 } 5025 5026 hsa_insn_basic *atominsn; 5027 hsa_op_base *tgt; 5028 if (signal) 5029 { 5030 atominsn = new hsa_insn_signal (nops, opcode, acode, mtype, memorder); 5031 tgt = hsa_reg_or_immed_for_gimple_op (gimple_call_arg (stmt, 0), hbb); 5032 } 5033 else 5034 { 5035 atominsn = new hsa_insn_atomic (nops, opcode, acode, mtype, memorder); 5036 hsa_op_address *addr; 5037 addr = get_address_from_value (gimple_call_arg (stmt, 0), hbb); 5038 if (addr->m_symbol && addr->m_symbol->m_segment == BRIG_SEGMENT_PRIVATE) 5039 { 5040 HSA_SORRY_AT (gimple_location (stmt), 5041 "HSA does not implement atomic operations in private " 5042 "segment"); 5043 return; 5044 } 5045 tgt = addr; 5046 } 5047 5048 hsa_op_with_type *op 5049 = hsa_reg_or_immed_for_gimple_op (gimple_call_arg (stmt, 1), hbb); 5050 if (lhs) 5051 { 5052 atominsn->set_op (0, dest); 5053 atominsn->set_op (1, tgt); 5054 atominsn->set_op (2, op); 5055 } 5056 else 5057 { 5058 atominsn->set_op (0, tgt); 5059 atominsn->set_op (1, op); 5060 } 5061 5062 hbb->append_insn (atominsn); 5063 5064 /* HSA does not natively support the variants that return the modified value, 5065 so re-do the operation again non-atomically if that is what was 5066 requested. */ 5067 if (lhs && !ret_orig) 5068 { 5069 int arith; 5070 switch (acode) 5071 { 5072 case BRIG_ATOMIC_ADD: 5073 arith = BRIG_OPCODE_ADD; 5074 break; 5075 case BRIG_ATOMIC_AND: 5076 arith = BRIG_OPCODE_AND; 5077 break; 5078 case BRIG_ATOMIC_OR: 5079 arith = BRIG_OPCODE_OR; 5080 break; 5081 case BRIG_ATOMIC_SUB: 5082 arith = BRIG_OPCODE_SUB; 5083 break; 5084 case BRIG_ATOMIC_XOR: 5085 arith = BRIG_OPCODE_XOR; 5086 break; 5087 default: 5088 gcc_unreachable (); 5089 } 5090 hsa_op_reg *real_dest = hsa_cfun->reg_for_gimple_ssa (lhs); 5091 gen_hsa_binary_operation (arith, real_dest, dest, op, hbb); 5092 } 5093 } 5094 5095 /* Generate HSA instructions for an internal fn. 5096 Instructions will be appended to HBB, which also needs to be the 5097 corresponding structure to the basic_block of STMT. */ 5098 5099 static void 5100 gen_hsa_insn_for_internal_fn_call (gcall *stmt, hsa_bb *hbb) 5101 { 5102 gcc_checking_assert (gimple_call_internal_fn (stmt)); 5103 internal_fn fn = gimple_call_internal_fn (stmt); 5104 5105 bool is_float_type_p = false; 5106 if (gimple_call_lhs (stmt) != NULL 5107 && TREE_TYPE (gimple_call_lhs (stmt)) == float_type_node) 5108 is_float_type_p = true; 5109 5110 switch (fn) 5111 { 5112 case IFN_CEIL: 5113 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_CEIL, stmt, hbb); 5114 break; 5115 5116 case IFN_FLOOR: 5117 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_FLOOR, stmt, hbb); 5118 break; 5119 5120 case IFN_RINT: 5121 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_RINT, stmt, hbb); 5122 break; 5123 5124 case IFN_SQRT: 5125 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_SQRT, stmt, hbb); 5126 break; 5127 5128 case IFN_RSQRT: 5129 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_NRSQRT, stmt, hbb); 5130 break; 5131 5132 case IFN_TRUNC: 5133 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_TRUNC, stmt, hbb); 5134 break; 5135 5136 case IFN_COS: 5137 { 5138 if (is_float_type_p) 5139 gen_hsa_unaryop_or_call_for_builtin (BRIG_OPCODE_NCOS, stmt, hbb); 5140 else 5141 gen_hsa_unaryop_builtin_call (stmt, hbb); 5142 5143 break; 5144 } 5145 case IFN_EXP2: 5146 { 5147 if (is_float_type_p) 5148 gen_hsa_unaryop_or_call_for_builtin (BRIG_OPCODE_NEXP2, stmt, hbb); 5149 else 5150 gen_hsa_unaryop_builtin_call (stmt, hbb); 5151 5152 break; 5153 } 5154 5155 case IFN_LOG2: 5156 { 5157 if (is_float_type_p) 5158 gen_hsa_unaryop_or_call_for_builtin (BRIG_OPCODE_NLOG2, stmt, hbb); 5159 else 5160 gen_hsa_unaryop_builtin_call (stmt, hbb); 5161 5162 break; 5163 } 5164 5165 case IFN_SIN: 5166 { 5167 if (is_float_type_p) 5168 gen_hsa_unaryop_or_call_for_builtin (BRIG_OPCODE_NSIN, stmt, hbb); 5169 else 5170 gen_hsa_unaryop_builtin_call (stmt, hbb); 5171 break; 5172 } 5173 5174 case IFN_CLRSB: 5175 gen_hsa_clrsb (stmt, hbb); 5176 break; 5177 5178 case IFN_CLZ: 5179 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_FIRSTBIT, stmt, hbb); 5180 break; 5181 5182 case IFN_CTZ: 5183 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_LASTBIT, stmt, hbb); 5184 break; 5185 5186 case IFN_FFS: 5187 gen_hsa_ffs (stmt, hbb); 5188 break; 5189 5190 case IFN_PARITY: 5191 gen_hsa_parity (stmt, hbb); 5192 break; 5193 5194 case IFN_POPCOUNT: 5195 gen_hsa_popcount (stmt, hbb); 5196 break; 5197 5198 case IFN_DIVMOD: 5199 gen_hsa_divmod (stmt, hbb); 5200 break; 5201 5202 case IFN_ACOS: 5203 case IFN_ASIN: 5204 case IFN_ATAN: 5205 case IFN_EXP: 5206 case IFN_EXP10: 5207 case IFN_EXPM1: 5208 case IFN_LOG: 5209 case IFN_LOG10: 5210 case IFN_LOG1P: 5211 case IFN_LOGB: 5212 case IFN_SIGNIFICAND: 5213 case IFN_TAN: 5214 case IFN_NEARBYINT: 5215 case IFN_ROUND: 5216 case IFN_ATAN2: 5217 case IFN_COPYSIGN: 5218 case IFN_FMOD: 5219 case IFN_POW: 5220 case IFN_REMAINDER: 5221 case IFN_SCALB: 5222 case IFN_FMIN: 5223 case IFN_FMAX: 5224 gen_hsa_insns_for_call_of_internal_fn (stmt, hbb); 5225 break; 5226 5227 default: 5228 HSA_SORRY_ATV (gimple_location (stmt), 5229 "support for HSA does not implement internal function: %s", 5230 internal_fn_name (fn)); 5231 break; 5232 } 5233 } 5234 5235 /* Generate HSA instructions for the given call statement STMT. Instructions 5236 will be appended to HBB. */ 5237 5238 static void 5239 gen_hsa_insns_for_call (gimple *stmt, hsa_bb *hbb) 5240 { 5241 gcall *call = as_a <gcall *> (stmt); 5242 tree lhs = gimple_call_lhs (stmt); 5243 hsa_op_reg *dest; 5244 5245 if (gimple_call_internal_p (stmt)) 5246 { 5247 gen_hsa_insn_for_internal_fn_call (call, hbb); 5248 return; 5249 } 5250 5251 if (!gimple_call_builtin_p (stmt, BUILT_IN_NORMAL)) 5252 { 5253 tree function_decl = gimple_call_fndecl (stmt); 5254 /* Prefetch pass can create type-mismatching prefetch builtin calls which 5255 fail the gimple_call_builtin_p test above. Handle them here. */ 5256 if (DECL_BUILT_IN_CLASS (function_decl) 5257 && DECL_FUNCTION_CODE (function_decl) == BUILT_IN_PREFETCH) 5258 return; 5259 5260 if (function_decl == NULL_TREE) 5261 { 5262 HSA_SORRY_AT (gimple_location (stmt), 5263 "support for HSA does not implement indirect calls"); 5264 return; 5265 } 5266 5267 if (hsa_callable_function_p (function_decl)) 5268 gen_hsa_insns_for_direct_call (stmt, hbb); 5269 else if (!gen_hsa_insns_for_known_library_call (stmt, hbb)) 5270 HSA_SORRY_AT (gimple_location (stmt), 5271 "HSA supports only calls of functions marked with pragma " 5272 "omp declare target"); 5273 return; 5274 } 5275 5276 tree fndecl = gimple_call_fndecl (stmt); 5277 enum built_in_function builtin = DECL_FUNCTION_CODE (fndecl); 5278 switch (builtin) 5279 { 5280 case BUILT_IN_FABS: 5281 case BUILT_IN_FABSF: 5282 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_ABS, stmt, hbb); 5283 break; 5284 5285 case BUILT_IN_CEIL: 5286 case BUILT_IN_CEILF: 5287 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_CEIL, stmt, hbb); 5288 break; 5289 5290 case BUILT_IN_FLOOR: 5291 case BUILT_IN_FLOORF: 5292 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_FLOOR, stmt, hbb); 5293 break; 5294 5295 case BUILT_IN_RINT: 5296 case BUILT_IN_RINTF: 5297 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_RINT, stmt, hbb); 5298 break; 5299 5300 case BUILT_IN_SQRT: 5301 case BUILT_IN_SQRTF: 5302 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_SQRT, stmt, hbb); 5303 break; 5304 5305 case BUILT_IN_TRUNC: 5306 case BUILT_IN_TRUNCF: 5307 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_TRUNC, stmt, hbb); 5308 break; 5309 5310 case BUILT_IN_COS: 5311 case BUILT_IN_SIN: 5312 case BUILT_IN_EXP2: 5313 case BUILT_IN_LOG2: 5314 /* HSAIL does not provide an instruction for double argument type. */ 5315 gen_hsa_unaryop_builtin_call (stmt, hbb); 5316 break; 5317 5318 case BUILT_IN_COSF: 5319 gen_hsa_unaryop_or_call_for_builtin (BRIG_OPCODE_NCOS, stmt, hbb); 5320 break; 5321 5322 case BUILT_IN_EXP2F: 5323 gen_hsa_unaryop_or_call_for_builtin (BRIG_OPCODE_NEXP2, stmt, hbb); 5324 break; 5325 5326 case BUILT_IN_LOG2F: 5327 gen_hsa_unaryop_or_call_for_builtin (BRIG_OPCODE_NLOG2, stmt, hbb); 5328 break; 5329 5330 case BUILT_IN_SINF: 5331 gen_hsa_unaryop_or_call_for_builtin (BRIG_OPCODE_NSIN, stmt, hbb); 5332 break; 5333 5334 case BUILT_IN_CLRSB: 5335 case BUILT_IN_CLRSBL: 5336 case BUILT_IN_CLRSBLL: 5337 gen_hsa_clrsb (call, hbb); 5338 break; 5339 5340 case BUILT_IN_CLZ: 5341 case BUILT_IN_CLZL: 5342 case BUILT_IN_CLZLL: 5343 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_FIRSTBIT, stmt, hbb); 5344 break; 5345 5346 case BUILT_IN_CTZ: 5347 case BUILT_IN_CTZL: 5348 case BUILT_IN_CTZLL: 5349 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_LASTBIT, stmt, hbb); 5350 break; 5351 5352 case BUILT_IN_FFS: 5353 case BUILT_IN_FFSL: 5354 case BUILT_IN_FFSLL: 5355 gen_hsa_ffs (call, hbb); 5356 break; 5357 5358 case BUILT_IN_PARITY: 5359 case BUILT_IN_PARITYL: 5360 case BUILT_IN_PARITYLL: 5361 gen_hsa_parity (call, hbb); 5362 break; 5363 5364 case BUILT_IN_POPCOUNT: 5365 case BUILT_IN_POPCOUNTL: 5366 case BUILT_IN_POPCOUNTLL: 5367 gen_hsa_popcount (call, hbb); 5368 break; 5369 5370 case BUILT_IN_ATOMIC_LOAD_1: 5371 case BUILT_IN_ATOMIC_LOAD_2: 5372 case BUILT_IN_ATOMIC_LOAD_4: 5373 case BUILT_IN_ATOMIC_LOAD_8: 5374 case BUILT_IN_ATOMIC_LOAD_16: 5375 { 5376 BrigType16_t mtype; 5377 hsa_op_base *src; 5378 src = get_address_from_value (gimple_call_arg (stmt, 0), hbb); 5379 5380 BrigMemoryOrder memorder; 5381 const char *mmname; 5382 if (hsa_memorder_from_tree (gimple_call_arg (stmt, 1), &memorder, 5383 &mmname, gimple_location (stmt))) 5384 return; 5385 5386 if (memorder == BRIG_MEMORY_ORDER_SC_ACQUIRE_RELEASE) 5387 memorder = BRIG_MEMORY_ORDER_SC_ACQUIRE; 5388 5389 if (memorder != BRIG_MEMORY_ORDER_RELAXED 5390 && memorder != BRIG_MEMORY_ORDER_SC_ACQUIRE 5391 && memorder != BRIG_MEMORY_ORDER_NONE) 5392 { 5393 HSA_SORRY_ATV (gimple_location (stmt), 5394 "support for HSA does not implement " 5395 "memory model for atomic loads: %s", mmname); 5396 return; 5397 } 5398 5399 if (lhs) 5400 { 5401 BrigType16_t t = hsa_type_for_scalar_tree_type (TREE_TYPE (lhs), 5402 false); 5403 mtype = mem_type_for_type (t); 5404 mtype = hsa_bittype_for_type (mtype); 5405 dest = hsa_cfun->reg_for_gimple_ssa (lhs); 5406 } 5407 else 5408 { 5409 mtype = BRIG_TYPE_B64; 5410 dest = new hsa_op_reg (mtype); 5411 } 5412 5413 hsa_insn_basic *atominsn; 5414 atominsn = new hsa_insn_atomic (2, BRIG_OPCODE_ATOMIC, BRIG_ATOMIC_LD, 5415 mtype, memorder, dest, src); 5416 5417 hbb->append_insn (atominsn); 5418 break; 5419 } 5420 5421 case BUILT_IN_ATOMIC_EXCHANGE_1: 5422 case BUILT_IN_ATOMIC_EXCHANGE_2: 5423 case BUILT_IN_ATOMIC_EXCHANGE_4: 5424 case BUILT_IN_ATOMIC_EXCHANGE_8: 5425 case BUILT_IN_ATOMIC_EXCHANGE_16: 5426 gen_hsa_atomic_for_builtin (true, BRIG_ATOMIC_EXCH, stmt, hbb, false); 5427 break; 5428 break; 5429 5430 case BUILT_IN_ATOMIC_FETCH_ADD_1: 5431 case BUILT_IN_ATOMIC_FETCH_ADD_2: 5432 case BUILT_IN_ATOMIC_FETCH_ADD_4: 5433 case BUILT_IN_ATOMIC_FETCH_ADD_8: 5434 case BUILT_IN_ATOMIC_FETCH_ADD_16: 5435 gen_hsa_atomic_for_builtin (true, BRIG_ATOMIC_ADD, stmt, hbb, false); 5436 break; 5437 break; 5438 5439 case BUILT_IN_ATOMIC_FETCH_SUB_1: 5440 case BUILT_IN_ATOMIC_FETCH_SUB_2: 5441 case BUILT_IN_ATOMIC_FETCH_SUB_4: 5442 case BUILT_IN_ATOMIC_FETCH_SUB_8: 5443 case BUILT_IN_ATOMIC_FETCH_SUB_16: 5444 gen_hsa_atomic_for_builtin (true, BRIG_ATOMIC_SUB, stmt, hbb, false); 5445 break; 5446 break; 5447 5448 case BUILT_IN_ATOMIC_FETCH_AND_1: 5449 case BUILT_IN_ATOMIC_FETCH_AND_2: 5450 case BUILT_IN_ATOMIC_FETCH_AND_4: 5451 case BUILT_IN_ATOMIC_FETCH_AND_8: 5452 case BUILT_IN_ATOMIC_FETCH_AND_16: 5453 gen_hsa_atomic_for_builtin (true, BRIG_ATOMIC_AND, stmt, hbb, false); 5454 break; 5455 break; 5456 5457 case BUILT_IN_ATOMIC_FETCH_XOR_1: 5458 case BUILT_IN_ATOMIC_FETCH_XOR_2: 5459 case BUILT_IN_ATOMIC_FETCH_XOR_4: 5460 case BUILT_IN_ATOMIC_FETCH_XOR_8: 5461 case BUILT_IN_ATOMIC_FETCH_XOR_16: 5462 gen_hsa_atomic_for_builtin (true, BRIG_ATOMIC_XOR, stmt, hbb, false); 5463 break; 5464 break; 5465 5466 case BUILT_IN_ATOMIC_FETCH_OR_1: 5467 case BUILT_IN_ATOMIC_FETCH_OR_2: 5468 case BUILT_IN_ATOMIC_FETCH_OR_4: 5469 case BUILT_IN_ATOMIC_FETCH_OR_8: 5470 case BUILT_IN_ATOMIC_FETCH_OR_16: 5471 gen_hsa_atomic_for_builtin (true, BRIG_ATOMIC_OR, stmt, hbb, false); 5472 break; 5473 break; 5474 5475 case BUILT_IN_ATOMIC_STORE_1: 5476 case BUILT_IN_ATOMIC_STORE_2: 5477 case BUILT_IN_ATOMIC_STORE_4: 5478 case BUILT_IN_ATOMIC_STORE_8: 5479 case BUILT_IN_ATOMIC_STORE_16: 5480 /* Since there cannot be any LHS, the first parameter is meaningless. */ 5481 gen_hsa_atomic_for_builtin (true, BRIG_ATOMIC_ST, stmt, hbb, false); 5482 break; 5483 break; 5484 5485 case BUILT_IN_ATOMIC_ADD_FETCH_1: 5486 case BUILT_IN_ATOMIC_ADD_FETCH_2: 5487 case BUILT_IN_ATOMIC_ADD_FETCH_4: 5488 case BUILT_IN_ATOMIC_ADD_FETCH_8: 5489 case BUILT_IN_ATOMIC_ADD_FETCH_16: 5490 gen_hsa_atomic_for_builtin (false, BRIG_ATOMIC_ADD, stmt, hbb, false); 5491 break; 5492 5493 case BUILT_IN_ATOMIC_SUB_FETCH_1: 5494 case BUILT_IN_ATOMIC_SUB_FETCH_2: 5495 case BUILT_IN_ATOMIC_SUB_FETCH_4: 5496 case BUILT_IN_ATOMIC_SUB_FETCH_8: 5497 case BUILT_IN_ATOMIC_SUB_FETCH_16: 5498 gen_hsa_atomic_for_builtin (false, BRIG_ATOMIC_SUB, stmt, hbb, false); 5499 break; 5500 5501 case BUILT_IN_ATOMIC_AND_FETCH_1: 5502 case BUILT_IN_ATOMIC_AND_FETCH_2: 5503 case BUILT_IN_ATOMIC_AND_FETCH_4: 5504 case BUILT_IN_ATOMIC_AND_FETCH_8: 5505 case BUILT_IN_ATOMIC_AND_FETCH_16: 5506 gen_hsa_atomic_for_builtin (false, BRIG_ATOMIC_AND, stmt, hbb, false); 5507 break; 5508 5509 case BUILT_IN_ATOMIC_XOR_FETCH_1: 5510 case BUILT_IN_ATOMIC_XOR_FETCH_2: 5511 case BUILT_IN_ATOMIC_XOR_FETCH_4: 5512 case BUILT_IN_ATOMIC_XOR_FETCH_8: 5513 case BUILT_IN_ATOMIC_XOR_FETCH_16: 5514 gen_hsa_atomic_for_builtin (false, BRIG_ATOMIC_XOR, stmt, hbb, false); 5515 break; 5516 5517 case BUILT_IN_ATOMIC_OR_FETCH_1: 5518 case BUILT_IN_ATOMIC_OR_FETCH_2: 5519 case BUILT_IN_ATOMIC_OR_FETCH_4: 5520 case BUILT_IN_ATOMIC_OR_FETCH_8: 5521 case BUILT_IN_ATOMIC_OR_FETCH_16: 5522 gen_hsa_atomic_for_builtin (false, BRIG_ATOMIC_OR, stmt, hbb, false); 5523 break; 5524 5525 case BUILT_IN_SYNC_VAL_COMPARE_AND_SWAP_1: 5526 case BUILT_IN_SYNC_VAL_COMPARE_AND_SWAP_2: 5527 case BUILT_IN_SYNC_VAL_COMPARE_AND_SWAP_4: 5528 case BUILT_IN_SYNC_VAL_COMPARE_AND_SWAP_8: 5529 case BUILT_IN_SYNC_VAL_COMPARE_AND_SWAP_16: 5530 { 5531 tree type = TREE_TYPE (gimple_call_arg (stmt, 1)); 5532 BrigType16_t atype 5533 = hsa_bittype_for_type (hsa_type_for_scalar_tree_type (type, false)); 5534 BrigMemoryOrder memorder = BRIG_MEMORY_ORDER_SC_ACQUIRE_RELEASE; 5535 hsa_insn_basic *atominsn; 5536 hsa_op_base *tgt; 5537 atominsn = new hsa_insn_atomic (4, BRIG_OPCODE_ATOMIC, 5538 BRIG_ATOMIC_CAS, atype, memorder); 5539 tgt = get_address_from_value (gimple_call_arg (stmt, 0), hbb); 5540 5541 if (lhs != NULL) 5542 dest = hsa_cfun->reg_for_gimple_ssa (lhs); 5543 else 5544 dest = new hsa_op_reg (atype); 5545 5546 atominsn->set_op (0, dest); 5547 atominsn->set_op (1, tgt); 5548 5549 hsa_op_with_type *op 5550 = hsa_reg_or_immed_for_gimple_op (gimple_call_arg (stmt, 1), hbb); 5551 atominsn->set_op (2, op); 5552 op = hsa_reg_or_immed_for_gimple_op (gimple_call_arg (stmt, 2), hbb); 5553 atominsn->set_op (3, op); 5554 5555 hbb->append_insn (atominsn); 5556 break; 5557 } 5558 5559 case BUILT_IN_HSA_WORKGROUPID: 5560 query_hsa_grid_dim (stmt, BRIG_OPCODE_WORKGROUPID, hbb); 5561 break; 5562 case BUILT_IN_HSA_WORKITEMID: 5563 query_hsa_grid_dim (stmt, BRIG_OPCODE_WORKITEMID, hbb); 5564 break; 5565 case BUILT_IN_HSA_WORKITEMABSID: 5566 query_hsa_grid_dim (stmt, BRIG_OPCODE_WORKITEMABSID, hbb); 5567 break; 5568 case BUILT_IN_HSA_GRIDSIZE: 5569 query_hsa_grid_dim (stmt, BRIG_OPCODE_GRIDSIZE, hbb); 5570 break; 5571 case BUILT_IN_HSA_CURRENTWORKGROUPSIZE: 5572 query_hsa_grid_dim (stmt, BRIG_OPCODE_CURRENTWORKGROUPSIZE, hbb); 5573 break; 5574 5575 case BUILT_IN_GOMP_BARRIER: 5576 hbb->append_insn (new hsa_insn_br (0, BRIG_OPCODE_BARRIER, BRIG_TYPE_NONE, 5577 BRIG_WIDTH_ALL)); 5578 break; 5579 case BUILT_IN_GOMP_PARALLEL: 5580 HSA_SORRY_AT (gimple_location (stmt), 5581 "support for HSA does not implement non-gridified " 5582 "OpenMP parallel constructs."); 5583 break; 5584 5585 case BUILT_IN_OMP_GET_THREAD_NUM: 5586 { 5587 query_hsa_grid_nodim (stmt, BRIG_OPCODE_WORKITEMFLATABSID, hbb); 5588 break; 5589 } 5590 5591 case BUILT_IN_OMP_GET_NUM_THREADS: 5592 { 5593 gen_get_num_threads (stmt, hbb); 5594 break; 5595 } 5596 case BUILT_IN_GOMP_TEAMS: 5597 { 5598 gen_set_num_threads (gimple_call_arg (stmt, 1), hbb); 5599 break; 5600 } 5601 case BUILT_IN_OMP_GET_NUM_TEAMS: 5602 { 5603 gen_get_num_teams (stmt, hbb); 5604 break; 5605 } 5606 case BUILT_IN_OMP_GET_TEAM_NUM: 5607 { 5608 gen_get_team_num (stmt, hbb); 5609 break; 5610 } 5611 case BUILT_IN_MEMCPY: 5612 case BUILT_IN_MEMPCPY: 5613 { 5614 expand_memory_copy (stmt, hbb, builtin); 5615 break; 5616 } 5617 case BUILT_IN_MEMSET: 5618 { 5619 tree c = gimple_call_arg (stmt, 1); 5620 5621 if (TREE_CODE (c) != INTEGER_CST) 5622 { 5623 gen_hsa_insns_for_direct_call (stmt, hbb); 5624 return; 5625 } 5626 5627 tree byte_size = gimple_call_arg (stmt, 2); 5628 5629 if (!tree_fits_uhwi_p (byte_size)) 5630 { 5631 gen_hsa_insns_for_direct_call (stmt, hbb); 5632 return; 5633 } 5634 5635 unsigned HOST_WIDE_INT n = tree_to_uhwi (byte_size); 5636 5637 if (n > HSA_MEMORY_BUILTINS_LIMIT) 5638 { 5639 gen_hsa_insns_for_direct_call (stmt, hbb); 5640 return; 5641 } 5642 5643 unsigned HOST_WIDE_INT constant 5644 = tree_to_uhwi (fold_convert (unsigned_char_type_node, c)); 5645 5646 expand_memory_set (stmt, n, constant, hbb, builtin); 5647 5648 break; 5649 } 5650 case BUILT_IN_BZERO: 5651 { 5652 tree byte_size = gimple_call_arg (stmt, 1); 5653 5654 if (!tree_fits_uhwi_p (byte_size)) 5655 { 5656 gen_hsa_insns_for_direct_call (stmt, hbb); 5657 return; 5658 } 5659 5660 unsigned HOST_WIDE_INT n = tree_to_uhwi (byte_size); 5661 5662 if (n > HSA_MEMORY_BUILTINS_LIMIT) 5663 { 5664 gen_hsa_insns_for_direct_call (stmt, hbb); 5665 return; 5666 } 5667 5668 expand_memory_set (stmt, n, 0, hbb, builtin); 5669 5670 break; 5671 } 5672 CASE_BUILT_IN_ALLOCA: 5673 { 5674 gen_hsa_alloca (call, hbb); 5675 break; 5676 } 5677 case BUILT_IN_PREFETCH: 5678 break; 5679 default: 5680 { 5681 tree name_tree = DECL_NAME (fndecl); 5682 const char *s = IDENTIFIER_POINTER (name_tree); 5683 size_t len = strlen (s); 5684 if (len > 4 && (strncmp (s, "__builtin_GOMP_", 15) == 0)) 5685 HSA_SORRY_ATV (gimple_location (stmt), 5686 "support for HSA does not implement GOMP function %s", 5687 s); 5688 else 5689 gen_hsa_insns_for_direct_call (stmt, hbb); 5690 return; 5691 } 5692 } 5693 } 5694 5695 /* Generate HSA instructions for a given gimple statement. Instructions will be 5696 appended to HBB. */ 5697 5698 static void 5699 gen_hsa_insns_for_gimple_stmt (gimple *stmt, hsa_bb *hbb) 5700 { 5701 switch (gimple_code (stmt)) 5702 { 5703 case GIMPLE_ASSIGN: 5704 if (gimple_clobber_p (stmt)) 5705 break; 5706 5707 if (gimple_assign_single_p (stmt)) 5708 { 5709 tree lhs = gimple_assign_lhs (stmt); 5710 tree rhs = gimple_assign_rhs1 (stmt); 5711 gen_hsa_insns_for_single_assignment (lhs, rhs, hbb); 5712 } 5713 else 5714 gen_hsa_insns_for_operation_assignment (stmt, hbb); 5715 break; 5716 case GIMPLE_RETURN: 5717 gen_hsa_insns_for_return (as_a <greturn *> (stmt), hbb); 5718 break; 5719 case GIMPLE_COND: 5720 gen_hsa_insns_for_cond_stmt (stmt, hbb); 5721 break; 5722 case GIMPLE_CALL: 5723 gen_hsa_insns_for_call (stmt, hbb); 5724 break; 5725 case GIMPLE_DEBUG: 5726 /* ??? HSA supports some debug facilities. */ 5727 break; 5728 case GIMPLE_LABEL: 5729 { 5730 tree label = gimple_label_label (as_a <glabel *> (stmt)); 5731 if (FORCED_LABEL (label)) 5732 HSA_SORRY_AT (gimple_location (stmt), 5733 "support for HSA does not implement gimple label with " 5734 "address taken"); 5735 5736 break; 5737 } 5738 case GIMPLE_NOP: 5739 { 5740 hbb->append_insn (new hsa_insn_basic (0, BRIG_OPCODE_NOP)); 5741 break; 5742 } 5743 case GIMPLE_SWITCH: 5744 { 5745 gen_hsa_insns_for_switch_stmt (as_a <gswitch *> (stmt), hbb); 5746 break; 5747 } 5748 default: 5749 HSA_SORRY_ATV (gimple_location (stmt), 5750 "support for HSA does not implement gimple statement %s", 5751 gimple_code_name[(int) gimple_code (stmt)]); 5752 } 5753 } 5754 5755 /* Generate a HSA PHI from a gimple PHI. */ 5756 5757 static void 5758 gen_hsa_phi_from_gimple_phi (gimple *phi_stmt, hsa_bb *hbb) 5759 { 5760 hsa_insn_phi *hphi; 5761 unsigned count = gimple_phi_num_args (phi_stmt); 5762 5763 hsa_op_reg *dest 5764 = hsa_cfun->reg_for_gimple_ssa (gimple_phi_result (phi_stmt)); 5765 hphi = new hsa_insn_phi (count, dest); 5766 hphi->m_bb = hbb->m_bb; 5767 5768 auto_vec <tree, 8> aexprs; 5769 auto_vec <hsa_op_reg *, 8> aregs; 5770 5771 /* Calling split_edge when processing a PHI node messes up with the order of 5772 gimple phi node arguments (it moves the one associated with the edge to 5773 the end). We need to keep the order of edges and arguments of HSA phi 5774 node arguments consistent, so we do all required splitting as the first 5775 step, and in reverse order as to not be affected by the re-orderings. */ 5776 for (unsigned j = count; j != 0; j--) 5777 { 5778 unsigned i = j - 1; 5779 tree op = gimple_phi_arg_def (phi_stmt, i); 5780 if (TREE_CODE (op) != ADDR_EXPR) 5781 continue; 5782 5783 edge e = gimple_phi_arg_edge (as_a <gphi *> (phi_stmt), i); 5784 hsa_bb *hbb_src = hsa_init_new_bb (split_edge (e)); 5785 hsa_op_address *addr = gen_hsa_addr (TREE_OPERAND (op, 0), 5786 hbb_src); 5787 5788 hsa_op_reg *dest 5789 = new hsa_op_reg (hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT)); 5790 hsa_insn_basic *insn 5791 = new hsa_insn_basic (2, BRIG_OPCODE_LDA, BRIG_TYPE_U64, 5792 dest, addr); 5793 hbb_src->append_insn (insn); 5794 aexprs.safe_push (op); 5795 aregs.safe_push (dest); 5796 } 5797 5798 tree lhs = gimple_phi_result (phi_stmt); 5799 for (unsigned i = 0; i < count; i++) 5800 { 5801 tree op = gimple_phi_arg_def (phi_stmt, i); 5802 5803 if (TREE_CODE (op) == SSA_NAME) 5804 { 5805 hsa_op_reg *hreg = hsa_cfun->reg_for_gimple_ssa (op); 5806 hphi->set_op (i, hreg); 5807 } 5808 else 5809 { 5810 gcc_assert (is_gimple_min_invariant (op)); 5811 tree t = TREE_TYPE (op); 5812 if (!POINTER_TYPE_P (t) 5813 || (TREE_CODE (op) == STRING_CST 5814 && TREE_CODE (TREE_TYPE (t)) == INTEGER_TYPE)) 5815 hphi->set_op (i, new hsa_op_immed (op)); 5816 else if (POINTER_TYPE_P (TREE_TYPE (lhs)) 5817 && TREE_CODE (op) == INTEGER_CST) 5818 { 5819 /* Handle assignment of NULL value to a pointer type. */ 5820 hphi->set_op (i, new hsa_op_immed (op)); 5821 } 5822 else if (TREE_CODE (op) == ADDR_EXPR) 5823 { 5824 hsa_op_reg *dest = NULL; 5825 for (unsigned a_idx = 0; a_idx < aexprs.length (); a_idx++) 5826 if (aexprs[a_idx] == op) 5827 { 5828 dest = aregs[a_idx]; 5829 break; 5830 } 5831 gcc_assert (dest); 5832 hphi->set_op (i, dest); 5833 } 5834 else 5835 { 5836 HSA_SORRY_AT (gimple_location (phi_stmt), 5837 "support for HSA does not handle PHI nodes with " 5838 "constant address operands"); 5839 return; 5840 } 5841 } 5842 } 5843 5844 hbb->append_phi (hphi); 5845 } 5846 5847 /* Constructor of class containing HSA-specific information about a basic 5848 block. CFG_BB is the CFG BB this HSA BB is associated with. IDX is the new 5849 index of this BB (so that the constructor does not attempt to use 5850 hsa_cfun during its construction). */ 5851 5852 hsa_bb::hsa_bb (basic_block cfg_bb, int idx) 5853 : m_bb (cfg_bb), m_first_insn (NULL), m_last_insn (NULL), m_first_phi (NULL), 5854 m_last_phi (NULL), m_index (idx) 5855 { 5856 gcc_assert (!cfg_bb->aux); 5857 cfg_bb->aux = this; 5858 } 5859 5860 /* Constructor of class containing HSA-specific information about a basic 5861 block. CFG_BB is the CFG BB this HSA BB is associated with. */ 5862 5863 hsa_bb::hsa_bb (basic_block cfg_bb) 5864 : m_bb (cfg_bb), m_first_insn (NULL), m_last_insn (NULL), m_first_phi (NULL), 5865 m_last_phi (NULL), m_index (hsa_cfun->m_hbb_count++) 5866 { 5867 gcc_assert (!cfg_bb->aux); 5868 cfg_bb->aux = this; 5869 } 5870 5871 /* Create and initialize and return a new hsa_bb structure for a given CFG 5872 basic block BB. */ 5873 5874 hsa_bb * 5875 hsa_init_new_bb (basic_block bb) 5876 { 5877 void *m = obstack_alloc (&hsa_obstack, sizeof (hsa_bb)); 5878 return new (m) hsa_bb (bb); 5879 } 5880 5881 /* Initialize OMP in an HSA basic block PROLOGUE. */ 5882 5883 static void 5884 init_prologue (void) 5885 { 5886 if (!hsa_cfun->m_kern_p) 5887 return; 5888 5889 hsa_bb *prologue = hsa_bb_for_bb (ENTRY_BLOCK_PTR_FOR_FN (cfun)); 5890 5891 /* Create a magic number that is going to be printed by libgomp. */ 5892 unsigned index = hsa_get_number_decl_kernel_mappings (); 5893 5894 /* Emit store to debug argument. */ 5895 if (PARAM_VALUE (PARAM_HSA_GEN_DEBUG_STORES) > 0) 5896 set_debug_value (prologue, new hsa_op_immed (1000 + index, BRIG_TYPE_U64)); 5897 } 5898 5899 /* Initialize hsa_num_threads to a default value. */ 5900 5901 static void 5902 init_hsa_num_threads (void) 5903 { 5904 hsa_bb *prologue = hsa_bb_for_bb (ENTRY_BLOCK_PTR_FOR_FN (cfun)); 5905 5906 /* Save the default value to private variable hsa_num_threads. */ 5907 hsa_insn_basic *basic 5908 = new hsa_insn_mem (BRIG_OPCODE_ST, hsa_num_threads->m_type, 5909 new hsa_op_immed (0, hsa_num_threads->m_type), 5910 new hsa_op_address (hsa_num_threads)); 5911 prologue->append_insn (basic); 5912 } 5913 5914 /* Go over gimple representation and generate our internal HSA one. */ 5915 5916 static void 5917 gen_body_from_gimple () 5918 { 5919 basic_block bb; 5920 5921 /* Verify CFG for complex edges we are unable to handle. */ 5922 edge_iterator ei; 5923 edge e; 5924 5925 FOR_EACH_BB_FN (bb, cfun) 5926 { 5927 FOR_EACH_EDGE (e, ei, bb->succs) 5928 { 5929 /* Verify all unsupported flags for edges that point 5930 to the same basic block. */ 5931 if (e->flags & EDGE_EH) 5932 { 5933 HSA_SORRY_AT (UNKNOWN_LOCATION, 5934 "support for HSA does not implement exception " 5935 "handling"); 5936 return; 5937 } 5938 } 5939 } 5940 5941 FOR_EACH_BB_FN (bb, cfun) 5942 { 5943 gimple_stmt_iterator gsi; 5944 hsa_bb *hbb = hsa_bb_for_bb (bb); 5945 if (hbb) 5946 continue; 5947 5948 hbb = hsa_init_new_bb (bb); 5949 5950 for (gsi = gsi_start_bb (bb); !gsi_end_p (gsi); gsi_next (&gsi)) 5951 { 5952 gen_hsa_insns_for_gimple_stmt (gsi_stmt (gsi), hbb); 5953 if (hsa_seen_error ()) 5954 return; 5955 } 5956 } 5957 5958 FOR_EACH_BB_FN (bb, cfun) 5959 { 5960 gimple_stmt_iterator gsi; 5961 hsa_bb *hbb = hsa_bb_for_bb (bb); 5962 gcc_assert (hbb != NULL); 5963 5964 for (gsi = gsi_start_phis (bb); !gsi_end_p (gsi); gsi_next (&gsi)) 5965 if (!virtual_operand_p (gimple_phi_result (gsi_stmt (gsi)))) 5966 gen_hsa_phi_from_gimple_phi (gsi_stmt (gsi), hbb); 5967 } 5968 5969 if (dump_file && (dump_flags & TDF_DETAILS)) 5970 { 5971 fprintf (dump_file, "------- Generated SSA form -------\n"); 5972 dump_hsa_cfun (dump_file); 5973 } 5974 } 5975 5976 static void 5977 gen_function_decl_parameters (hsa_function_representation *f, 5978 tree decl) 5979 { 5980 tree parm; 5981 unsigned i; 5982 5983 for (parm = TYPE_ARG_TYPES (TREE_TYPE (decl)), i = 0; 5984 parm; 5985 parm = TREE_CHAIN (parm), i++) 5986 { 5987 /* Result type if last in the tree list. */ 5988 if (TREE_CHAIN (parm) == NULL) 5989 break; 5990 5991 tree v = TREE_VALUE (parm); 5992 5993 hsa_symbol *arg = new hsa_symbol (BRIG_TYPE_NONE, BRIG_SEGMENT_ARG, 5994 BRIG_LINKAGE_NONE); 5995 arg->m_type = hsa_type_for_tree_type (v, &arg->m_dim); 5996 arg->m_name_number = i; 5997 5998 f->m_input_args.safe_push (arg); 5999 } 6000 6001 tree result_type = TREE_TYPE (TREE_TYPE (decl)); 6002 if (!VOID_TYPE_P (result_type)) 6003 { 6004 f->m_output_arg = new hsa_symbol (BRIG_TYPE_NONE, BRIG_SEGMENT_ARG, 6005 BRIG_LINKAGE_NONE); 6006 f->m_output_arg->m_type 6007 = hsa_type_for_tree_type (result_type, &f->m_output_arg->m_dim); 6008 f->m_output_arg->m_name = "res"; 6009 } 6010 } 6011 6012 /* Generate the vector of parameters of the HSA representation of the current 6013 function. This also includes the output parameter representing the 6014 result. */ 6015 6016 static void 6017 gen_function_def_parameters () 6018 { 6019 tree parm; 6020 6021 hsa_bb *prologue = hsa_bb_for_bb (ENTRY_BLOCK_PTR_FOR_FN (cfun)); 6022 6023 for (parm = DECL_ARGUMENTS (cfun->decl); parm; 6024 parm = DECL_CHAIN (parm)) 6025 { 6026 struct hsa_symbol **slot; 6027 6028 hsa_symbol *arg 6029 = new hsa_symbol (BRIG_TYPE_NONE, hsa_cfun->m_kern_p 6030 ? BRIG_SEGMENT_KERNARG : BRIG_SEGMENT_ARG, 6031 BRIG_LINKAGE_FUNCTION); 6032 arg->fillup_for_decl (parm); 6033 6034 hsa_cfun->m_input_args.safe_push (arg); 6035 6036 if (hsa_seen_error ()) 6037 return; 6038 6039 arg->m_name = hsa_get_declaration_name (parm); 6040 6041 /* Copy all input arguments and create corresponding private symbols 6042 for them. */ 6043 hsa_symbol *private_arg; 6044 hsa_op_address *parm_addr = new hsa_op_address (arg); 6045 6046 if (TREE_ADDRESSABLE (parm) 6047 || (!is_gimple_reg (parm) && !TREE_READONLY (parm))) 6048 { 6049 private_arg = hsa_cfun->create_hsa_temporary (arg->m_type); 6050 private_arg->fillup_for_decl (parm); 6051 6052 BrigAlignment8_t align = MIN (arg->m_align, private_arg->m_align); 6053 6054 hsa_op_address *private_arg_addr = new hsa_op_address (private_arg); 6055 gen_hsa_memory_copy (prologue, private_arg_addr, parm_addr, 6056 arg->total_byte_size (), align); 6057 } 6058 else 6059 private_arg = arg; 6060 6061 slot = hsa_cfun->m_local_symbols->find_slot (private_arg, INSERT); 6062 gcc_assert (!*slot); 6063 *slot = private_arg; 6064 6065 if (is_gimple_reg (parm)) 6066 { 6067 tree ddef = ssa_default_def (cfun, parm); 6068 if (ddef && !has_zero_uses (ddef)) 6069 { 6070 BrigType16_t t = hsa_type_for_scalar_tree_type (TREE_TYPE (ddef), 6071 false); 6072 BrigType16_t mtype = mem_type_for_type (t); 6073 hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (ddef); 6074 hsa_insn_mem *mem = new hsa_insn_mem (BRIG_OPCODE_LD, mtype, 6075 dest, parm_addr); 6076 gcc_assert (!parm_addr->m_reg); 6077 prologue->append_insn (mem); 6078 } 6079 } 6080 } 6081 6082 if (!VOID_TYPE_P (TREE_TYPE (TREE_TYPE (cfun->decl)))) 6083 { 6084 struct hsa_symbol **slot; 6085 6086 hsa_cfun->m_output_arg = new hsa_symbol (BRIG_TYPE_NONE, BRIG_SEGMENT_ARG, 6087 BRIG_LINKAGE_FUNCTION); 6088 hsa_cfun->m_output_arg->fillup_for_decl (DECL_RESULT (cfun->decl)); 6089 6090 if (hsa_seen_error ()) 6091 return; 6092 6093 hsa_cfun->m_output_arg->m_name = "res"; 6094 slot = hsa_cfun->m_local_symbols->find_slot (hsa_cfun->m_output_arg, 6095 INSERT); 6096 gcc_assert (!*slot); 6097 *slot = hsa_cfun->m_output_arg; 6098 } 6099 } 6100 6101 /* Generate function representation that corresponds to 6102 a function declaration. */ 6103 6104 hsa_function_representation * 6105 hsa_generate_function_declaration (tree decl) 6106 { 6107 hsa_function_representation *fun 6108 = new hsa_function_representation (decl, false, 0); 6109 6110 fun->m_declaration_p = true; 6111 fun->m_name = get_brig_function_name (decl); 6112 gen_function_decl_parameters (fun, decl); 6113 6114 return fun; 6115 } 6116 6117 6118 /* Generate function representation that corresponds to 6119 an internal FN. */ 6120 6121 hsa_function_representation * 6122 hsa_generate_internal_fn_decl (hsa_internal_fn *fn) 6123 { 6124 hsa_function_representation *fun = new hsa_function_representation (fn); 6125 6126 fun->m_name = fn->name (); 6127 6128 for (unsigned i = 0; i < fn->get_arity (); i++) 6129 { 6130 hsa_symbol *arg 6131 = new hsa_symbol (fn->get_argument_type (i), BRIG_SEGMENT_ARG, 6132 BRIG_LINKAGE_NONE); 6133 arg->m_name_number = i; 6134 fun->m_input_args.safe_push (arg); 6135 } 6136 6137 fun->m_output_arg = new hsa_symbol (fn->get_argument_type (-1), 6138 BRIG_SEGMENT_ARG, BRIG_LINKAGE_NONE); 6139 fun->m_output_arg->m_name = "res"; 6140 6141 return fun; 6142 } 6143 6144 /* Return true if switch statement S can be transformed 6145 to a SBR instruction in HSAIL. */ 6146 6147 static bool 6148 transformable_switch_to_sbr_p (gswitch *s) 6149 { 6150 /* Identify if a switch statement can be transformed to 6151 SBR instruction, like: 6152 6153 sbr_u32 $s1 [@label1, @label2, @label3]; 6154 */ 6155 6156 tree size = get_switch_size (s); 6157 if (!tree_fits_uhwi_p (size)) 6158 return false; 6159 6160 if (tree_to_uhwi (size) > HSA_MAXIMUM_SBR_LABELS) 6161 return false; 6162 6163 return true; 6164 } 6165 6166 /* Structure hold connection between PHI nodes and immediate 6167 values hold by there nodes. */ 6168 6169 struct phi_definition 6170 { 6171 phi_definition (unsigned phi_i, unsigned label_i, tree imm): 6172 phi_index (phi_i), label_index (label_i), phi_value (imm) 6173 {} 6174 6175 unsigned phi_index; 6176 unsigned label_index; 6177 tree phi_value; 6178 }; 6179 6180 /* Sum slice of a vector V, starting from index START and ending 6181 at the index END - 1. */ 6182 6183 template <typename T> 6184 static 6185 T sum_slice (const auto_vec <T> &v, unsigned start, unsigned end, 6186 T zero) 6187 { 6188 T s = zero; 6189 6190 for (unsigned i = start; i < end; i++) 6191 s += v[i]; 6192 6193 return s; 6194 } 6195 6196 /* Function transforms GIMPLE SWITCH statements to a series of IF statements. 6197 Let's assume following example: 6198 6199 L0: 6200 switch (index) 6201 case C1: 6202 L1: hard_work_1 (); 6203 break; 6204 case C2..C3: 6205 L2: hard_work_2 (); 6206 break; 6207 default: 6208 LD: hard_work_3 (); 6209 break; 6210 6211 The transformation encompasses following steps: 6212 1) all immediate values used by edges coming from the switch basic block 6213 are saved 6214 2) all these edges are removed 6215 3) the switch statement (in L0) is replaced by: 6216 if (index == C1) 6217 goto L1; 6218 else 6219 goto L1'; 6220 6221 4) newly created basic block Lx' is used for generation of 6222 a next condition 6223 5) else branch of the last condition goes to LD 6224 6) fix all immediate values in PHI nodes that were propagated though 6225 edges that were removed in step 2 6226 6227 Note: if a case is made by a range C1..C2, then process 6228 following transformation: 6229 6230 switch_cond_op1 = C1 <= index; 6231 switch_cond_op2 = index <= C2; 6232 switch_cond_and = switch_cond_op1 & switch_cond_op2; 6233 if (switch_cond_and != 0) 6234 goto Lx; 6235 else 6236 goto Ly; 6237 6238 */ 6239 6240 static bool 6241 convert_switch_statements (void) 6242 { 6243 function *func = DECL_STRUCT_FUNCTION (current_function_decl); 6244 basic_block bb; 6245 6246 bool modified_cfg = false; 6247 6248 FOR_EACH_BB_FN (bb, func) 6249 { 6250 gimple_stmt_iterator gsi = gsi_last_bb (bb); 6251 if (gsi_end_p (gsi)) 6252 continue; 6253 6254 gimple *stmt = gsi_stmt (gsi); 6255 6256 if (gimple_code (stmt) == GIMPLE_SWITCH) 6257 { 6258 gswitch *s = as_a <gswitch *> (stmt); 6259 6260 /* If the switch can utilize SBR insn, skip the statement. */ 6261 if (transformable_switch_to_sbr_p (s)) 6262 continue; 6263 6264 modified_cfg = true; 6265 6266 unsigned labels = gimple_switch_num_labels (s); 6267 tree index = gimple_switch_index (s); 6268 tree index_type = TREE_TYPE (index); 6269 tree default_label = gimple_switch_default_label (s); 6270 basic_block default_label_bb 6271 = label_to_block_fn (func, CASE_LABEL (default_label)); 6272 basic_block cur_bb = bb; 6273 6274 auto_vec <edge> new_edges; 6275 auto_vec <phi_definition *> phi_todo_list; 6276 auto_vec <profile_count> edge_counts; 6277 auto_vec <profile_probability> edge_probabilities; 6278 6279 /* Investigate all labels that and PHI nodes in these edges which 6280 should be fixed after we add new collection of edges. */ 6281 for (unsigned i = 0; i < labels; i++) 6282 { 6283 tree label = gimple_switch_label (s, i); 6284 basic_block label_bb = label_to_block_fn (func, CASE_LABEL (label)); 6285 edge e = find_edge (bb, label_bb); 6286 edge_counts.safe_push (e->count ()); 6287 edge_probabilities.safe_push (e->probability); 6288 gphi_iterator phi_gsi; 6289 6290 /* Save PHI definitions that will be destroyed because of an edge 6291 is going to be removed. */ 6292 unsigned phi_index = 0; 6293 for (phi_gsi = gsi_start_phis (e->dest); 6294 !gsi_end_p (phi_gsi); gsi_next (&phi_gsi)) 6295 { 6296 gphi *phi = phi_gsi.phi (); 6297 for (unsigned j = 0; j < gimple_phi_num_args (phi); j++) 6298 { 6299 if (gimple_phi_arg_edge (phi, j) == e) 6300 { 6301 tree imm = gimple_phi_arg_def (phi, j); 6302 phi_definition *p = new phi_definition (phi_index, i, 6303 imm); 6304 phi_todo_list.safe_push (p); 6305 break; 6306 } 6307 } 6308 phi_index++; 6309 } 6310 } 6311 6312 /* Remove all edges for the current basic block. */ 6313 for (int i = EDGE_COUNT (bb->succs) - 1; i >= 0; i--) 6314 { 6315 edge e = EDGE_SUCC (bb, i); 6316 remove_edge (e); 6317 } 6318 6319 /* Iterate all non-default labels. */ 6320 for (unsigned i = 1; i < labels; i++) 6321 { 6322 tree label = gimple_switch_label (s, i); 6323 tree low = CASE_LOW (label); 6324 tree high = CASE_HIGH (label); 6325 6326 if (!useless_type_conversion_p (TREE_TYPE (low), index_type)) 6327 low = fold_convert (index_type, low); 6328 6329 gimple_stmt_iterator cond_gsi = gsi_last_bb (cur_bb); 6330 gimple *c = NULL; 6331 if (high) 6332 { 6333 tree tmp1 = make_temp_ssa_name (boolean_type_node, NULL, 6334 "switch_cond_op1"); 6335 6336 gimple *assign1 = gimple_build_assign (tmp1, LE_EXPR, low, 6337 index); 6338 6339 tree tmp2 = make_temp_ssa_name (boolean_type_node, NULL, 6340 "switch_cond_op2"); 6341 6342 if (!useless_type_conversion_p (TREE_TYPE (high), index_type)) 6343 high = fold_convert (index_type, high); 6344 gimple *assign2 = gimple_build_assign (tmp2, LE_EXPR, index, 6345 high); 6346 6347 tree tmp3 = make_temp_ssa_name (boolean_type_node, NULL, 6348 "switch_cond_and"); 6349 gimple *assign3 = gimple_build_assign (tmp3, BIT_AND_EXPR, tmp1, 6350 tmp2); 6351 6352 gsi_insert_before (&cond_gsi, assign1, GSI_SAME_STMT); 6353 gsi_insert_before (&cond_gsi, assign2, GSI_SAME_STMT); 6354 gsi_insert_before (&cond_gsi, assign3, GSI_SAME_STMT); 6355 6356 tree b = constant_boolean_node (false, boolean_type_node); 6357 c = gimple_build_cond (NE_EXPR, tmp3, b, NULL, NULL); 6358 } 6359 else 6360 c = gimple_build_cond (EQ_EXPR, index, low, NULL, NULL); 6361 6362 gimple_set_location (c, gimple_location (stmt)); 6363 6364 gsi_insert_before (&cond_gsi, c, GSI_SAME_STMT); 6365 6366 basic_block label_bb 6367 = label_to_block_fn (func, CASE_LABEL (label)); 6368 edge new_edge = make_edge (cur_bb, label_bb, EDGE_TRUE_VALUE); 6369 profile_probability prob_sum = sum_slice <profile_probability> 6370 (edge_probabilities, i, labels, profile_probability::never ()) 6371 + edge_probabilities[0]; 6372 6373 if (prob_sum.initialized_p ()) 6374 new_edge->probability = edge_probabilities[i] / prob_sum; 6375 6376 new_edges.safe_push (new_edge); 6377 6378 if (i < labels - 1) 6379 { 6380 /* Prepare another basic block that will contain 6381 next condition. */ 6382 basic_block next_bb = create_empty_bb (cur_bb); 6383 if (current_loops) 6384 { 6385 add_bb_to_loop (next_bb, cur_bb->loop_father); 6386 loops_state_set (LOOPS_NEED_FIXUP); 6387 } 6388 6389 edge next_edge = make_edge (cur_bb, next_bb, EDGE_FALSE_VALUE); 6390 next_edge->probability = new_edge->probability.invert (); 6391 next_bb->count = next_edge->count (); 6392 cur_bb = next_bb; 6393 } 6394 else /* Link last IF statement and default label 6395 of the switch. */ 6396 { 6397 edge e = make_edge (cur_bb, default_label_bb, EDGE_FALSE_VALUE); 6398 e->probability = new_edge->probability.invert (); 6399 new_edges.safe_insert (0, e); 6400 } 6401 } 6402 6403 /* Restore original PHI immediate value. */ 6404 for (unsigned i = 0; i < phi_todo_list.length (); i++) 6405 { 6406 phi_definition *phi_def = phi_todo_list[i]; 6407 edge new_edge = new_edges[phi_def->label_index]; 6408 6409 gphi_iterator it = gsi_start_phis (new_edge->dest); 6410 for (unsigned i = 0; i < phi_def->phi_index; i++) 6411 gsi_next (&it); 6412 6413 gphi *phi = it.phi (); 6414 add_phi_arg (phi, phi_def->phi_value, new_edge, UNKNOWN_LOCATION); 6415 delete phi_def; 6416 } 6417 6418 /* Remove the original GIMPLE switch statement. */ 6419 gsi_remove (&gsi, true); 6420 } 6421 } 6422 6423 if (dump_file) 6424 dump_function_to_file (current_function_decl, dump_file, TDF_DETAILS); 6425 6426 return modified_cfg; 6427 } 6428 6429 /* Expand builtins that can't be handled by HSA back-end. */ 6430 6431 static void 6432 expand_builtins () 6433 { 6434 function *func = DECL_STRUCT_FUNCTION (current_function_decl); 6435 basic_block bb; 6436 6437 FOR_EACH_BB_FN (bb, func) 6438 { 6439 for (gimple_stmt_iterator gsi = gsi_start_bb (bb); !gsi_end_p (gsi); 6440 gsi_next (&gsi)) 6441 { 6442 gimple *stmt = gsi_stmt (gsi); 6443 6444 if (gimple_code (stmt) != GIMPLE_CALL) 6445 continue; 6446 6447 gcall *call = as_a <gcall *> (stmt); 6448 6449 if (!gimple_call_builtin_p (call, BUILT_IN_NORMAL)) 6450 continue; 6451 6452 tree fndecl = gimple_call_fndecl (stmt); 6453 enum built_in_function fn = DECL_FUNCTION_CODE (fndecl); 6454 switch (fn) 6455 { 6456 case BUILT_IN_CEXPF: 6457 case BUILT_IN_CEXPIF: 6458 case BUILT_IN_CEXPI: 6459 { 6460 /* Similar to builtins.c (expand_builtin_cexpi), the builtin 6461 can be transformed to: cexp(I * z) = ccos(z) + I * csin(z). */ 6462 tree lhs = gimple_call_lhs (stmt); 6463 tree rhs = gimple_call_arg (stmt, 0); 6464 tree rhs_type = TREE_TYPE (rhs); 6465 bool float_type_p = rhs_type == float_type_node; 6466 tree real_part = make_temp_ssa_name (rhs_type, NULL, 6467 "cexp_real_part"); 6468 tree imag_part = make_temp_ssa_name (rhs_type, NULL, 6469 "cexp_imag_part"); 6470 6471 tree cos_fndecl 6472 = mathfn_built_in (rhs_type, fn == float_type_p 6473 ? BUILT_IN_COSF : BUILT_IN_COS); 6474 gcall *cos = gimple_build_call (cos_fndecl, 1, rhs); 6475 gimple_call_set_lhs (cos, real_part); 6476 gsi_insert_before (&gsi, cos, GSI_SAME_STMT); 6477 6478 tree sin_fndecl 6479 = mathfn_built_in (rhs_type, fn == float_type_p 6480 ? BUILT_IN_SINF : BUILT_IN_SIN); 6481 gcall *sin = gimple_build_call (sin_fndecl, 1, rhs); 6482 gimple_call_set_lhs (sin, imag_part); 6483 gsi_insert_before (&gsi, sin, GSI_SAME_STMT); 6484 6485 6486 gassign *assign = gimple_build_assign (lhs, COMPLEX_EXPR, 6487 real_part, imag_part); 6488 gsi_insert_before (&gsi, assign, GSI_SAME_STMT); 6489 gsi_remove (&gsi, true); 6490 6491 break; 6492 } 6493 default: 6494 break; 6495 } 6496 } 6497 } 6498 } 6499 6500 /* Emit HSA module variables that are global for the entire module. */ 6501 6502 static void 6503 emit_hsa_module_variables (void) 6504 { 6505 hsa_num_threads = new hsa_symbol (BRIG_TYPE_U32, BRIG_SEGMENT_PRIVATE, 6506 BRIG_LINKAGE_MODULE, true); 6507 6508 hsa_num_threads->m_name = "hsa_num_threads"; 6509 6510 hsa_brig_emit_omp_symbols (); 6511 } 6512 6513 /* Generate HSAIL representation of the current function and write into a 6514 special section of the output file. If KERNEL is set, the function will be 6515 considered an HSA kernel callable from the host, otherwise it will be 6516 compiled as an HSA function callable from other HSA code. */ 6517 6518 static void 6519 generate_hsa (bool kernel) 6520 { 6521 hsa_init_data_for_cfun (); 6522 6523 if (hsa_num_threads == NULL) 6524 emit_hsa_module_variables (); 6525 6526 bool modified_cfg = convert_switch_statements (); 6527 /* Initialize hsa_cfun. */ 6528 hsa_cfun = new hsa_function_representation (cfun->decl, kernel, 6529 SSANAMES (cfun)->length (), 6530 modified_cfg); 6531 hsa_cfun->init_extra_bbs (); 6532 6533 if (flag_tm) 6534 { 6535 HSA_SORRY_AT (UNKNOWN_LOCATION, 6536 "support for HSA does not implement transactional memory"); 6537 goto fail; 6538 } 6539 6540 verify_function_arguments (cfun->decl); 6541 if (hsa_seen_error ()) 6542 goto fail; 6543 6544 hsa_cfun->m_name = get_brig_function_name (cfun->decl); 6545 6546 gen_function_def_parameters (); 6547 if (hsa_seen_error ()) 6548 goto fail; 6549 6550 init_prologue (); 6551 6552 gen_body_from_gimple (); 6553 if (hsa_seen_error ()) 6554 goto fail; 6555 6556 if (hsa_cfun->m_kernel_dispatch_count) 6557 init_hsa_num_threads (); 6558 6559 if (hsa_cfun->m_kern_p) 6560 { 6561 hsa_function_summary *s 6562 = hsa_summaries->get (cgraph_node::get (hsa_cfun->m_decl)); 6563 hsa_add_kern_decl_mapping (current_function_decl, hsa_cfun->m_name, 6564 hsa_cfun->m_maximum_omp_data_size, 6565 s->m_gridified_kernel_p); 6566 } 6567 6568 if (flag_checking) 6569 { 6570 for (unsigned i = 0; i < hsa_cfun->m_ssa_map.length (); i++) 6571 if (hsa_cfun->m_ssa_map[i]) 6572 hsa_cfun->m_ssa_map[i]->verify_ssa (); 6573 6574 basic_block bb; 6575 FOR_EACH_BB_FN (bb, cfun) 6576 { 6577 hsa_bb *hbb = hsa_bb_for_bb (bb); 6578 6579 for (hsa_insn_basic *insn = hbb->m_first_insn; insn; 6580 insn = insn->m_next) 6581 insn->verify (); 6582 } 6583 } 6584 6585 hsa_regalloc (); 6586 hsa_brig_emit_function (); 6587 6588 fail: 6589 hsa_deinit_data_for_cfun (); 6590 } 6591 6592 namespace { 6593 6594 const pass_data pass_data_gen_hsail = 6595 { 6596 GIMPLE_PASS, 6597 "hsagen", /* name */ 6598 OPTGROUP_OMP, /* optinfo_flags */ 6599 TV_NONE, /* tv_id */ 6600 PROP_cfg | PROP_ssa, /* properties_required */ 6601 0, /* properties_provided */ 6602 0, /* properties_destroyed */ 6603 0, /* todo_flags_start */ 6604 0 /* todo_flags_finish */ 6605 }; 6606 6607 class pass_gen_hsail : public gimple_opt_pass 6608 { 6609 public: 6610 pass_gen_hsail (gcc::context *ctxt) 6611 : gimple_opt_pass(pass_data_gen_hsail, ctxt) 6612 {} 6613 6614 /* opt_pass methods: */ 6615 bool gate (function *); 6616 unsigned int execute (function *); 6617 6618 }; // class pass_gen_hsail 6619 6620 /* Determine whether or not to run generation of HSAIL. */ 6621 6622 bool 6623 pass_gen_hsail::gate (function *f) 6624 { 6625 return hsa_gen_requested_p () 6626 && hsa_gpu_implementation_p (f->decl); 6627 } 6628 6629 unsigned int 6630 pass_gen_hsail::execute (function *) 6631 { 6632 hsa_function_summary *s 6633 = hsa_summaries->get (cgraph_node::get_create (current_function_decl)); 6634 6635 expand_builtins (); 6636 generate_hsa (s->m_kind == HSA_KERNEL); 6637 TREE_ASM_WRITTEN (current_function_decl) = 1; 6638 return TODO_discard_function; 6639 } 6640 6641 } // anon namespace 6642 6643 /* Create the instance of hsa gen pass. */ 6644 6645 gimple_opt_pass * 6646 make_pass_gen_hsail (gcc::context *ctxt) 6647 { 6648 return new pass_gen_hsail (ctxt); 6649 } 6650