1 /* Copyright (C) 2016-2021 Free Software Foundation, Inc.
3 This file is free software; you can redistribute it and/or modify it under
4 the terms of the GNU General Public License as published by the Free
5 Software Foundation; either version 3 of the License, or (at your option)
8 This file is distributed in the hope that it will be useful, but WITHOUT
9 ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
10 FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License
13 You should have received a copy of the GNU General Public License
14 along with GCC; see the file COPYING3. If not see
15 <http://www.gnu.org/licenses/>. */
19 /* We want GET_MODE_SIZE et al to return integers, please. */
20 #define IN_TARGET_CODE 1
24 #include "coretypes.h"
32 #include "stringpool.h"
37 #include "diagnostic-core.h"
38 #include "insn-attr.h"
39 #include "fold-const.h"
45 #include "langhooks.h"
47 #include "omp-general.h"
48 #include "print-rtl.h"
54 /* This file should be included last. */
55 #include "target-def.h"
58 /* {{{ Global variables. */
60 /* Constants used by FP instructions. */
62 static REAL_VALUE_TYPE dconst4
, dconst1over2pi
;
63 static bool ext_gcn_constants_init
= 0;
65 /* Holds the ISA variant, derived from the command line parameters. */
67 int gcn_isa
= 3; /* Default to GCN3. */
69 /* Reserve this much space for LDS (for propagating variables from
70 worker-single mode to worker-partitioned mode), per workgroup. Global
71 analysis could calculate an exact bound, but we don't do that yet.
73 We want to permit full occupancy, so size accordingly. */
75 #define OMP_LDS_SIZE 0x600 /* 0x600 is 1/40 total, rounded down. */
76 #define ACC_LDS_SIZE 32768 /* Half of the total should be fine. */
77 #define OTHER_LDS_SIZE 65536 /* If in doubt, reserve all of it. */
79 #define LDS_SIZE (flag_openacc ? ACC_LDS_SIZE \
80 : flag_openmp ? OMP_LDS_SIZE \
83 /* The number of registers usable by normal non-kernel functions.
84 The SGPR count includes any special extra registers such as VCC. */
86 #define MAX_NORMAL_SGPR_COUNT 62 // i.e. 64 with VCC
87 #define MAX_NORMAL_VGPR_COUNT 24
90 /* {{{ Initialization and options. */
92 /* Initialize machine_function. */
94 static struct machine_function
*
95 gcn_init_machine_status (void)
97 struct machine_function
*f
;
99 f
= ggc_cleared_alloc
<machine_function
> ();
101 /* Set up LDS allocation for broadcasting for this function. */
102 f
->lds_allocated
= 32;
103 f
->lds_allocs
= hash_map
<tree
, int>::create_ggc (64);
105 /* And LDS temporary decls for worker reductions. */
106 vec_alloc (f
->reduc_decls
, 0);
109 f
->use_flat_addressing
= true;
114 /* Implement TARGET_OPTION_OVERRIDE.
116 Override option settings where defaults are variable, or we have specific
117 needs to consider. */
120 gcn_option_override (void)
122 init_machine_status
= gcn_init_machine_status
;
124 /* The HSA runtime does not respect ELF load addresses, so force PIE. */
130 gcn_isa
= gcn_arch
== PROCESSOR_FIJI
? 3 : 5;
132 /* The default stack size needs to be small for offload kernels because
133 there may be many, many threads. Also, a smaller stack gives a
134 measureable performance boost. But, a small stack is insufficient
135 for running the testsuite, so we use a larger default for the stand
137 if (stack_size_opt
== -1)
139 if (flag_openacc
|| flag_openmp
)
140 /* 512 bytes per work item = 32kB total. */
141 stack_size_opt
= 512 * 64;
144 stack_size_opt
= 1048576;
149 /* {{{ Attributes. */
151 /* This table defines the arguments that are permitted in
152 __attribute__ ((amdgpu_hsa_kernel (...))).
154 The names and values correspond to the HSA metadata that is encoded
155 into the assembler file and binary. */
157 static const struct gcn_kernel_arg_type
160 const char *header_pseudo
;
163 /* This should be set to -1 or -2 for a dynamically allocated register
164 number. Use -1 if this argument contributes to the user_sgpr_count,
167 } gcn_kernel_arg_types
[] = {
168 {"exec", NULL
, DImode
, EXEC_REG
},
169 #define PRIVATE_SEGMENT_BUFFER_ARG 1
170 {"private_segment_buffer",
171 ".amdhsa_user_sgpr_private_segment_buffer", TImode
, -1},
172 #define DISPATCH_PTR_ARG 2
173 {"dispatch_ptr", ".amdhsa_user_sgpr_dispatch_ptr", DImode
, -1},
174 #define QUEUE_PTR_ARG 3
175 {"queue_ptr", ".amdhsa_user_sgpr_queue_ptr", DImode
, -1},
176 #define KERNARG_SEGMENT_PTR_ARG 4
177 {"kernarg_segment_ptr", ".amdhsa_user_sgpr_kernarg_segment_ptr", DImode
, -1},
178 {"dispatch_id", ".amdhsa_user_sgpr_dispatch_id", DImode
, -1},
179 #define FLAT_SCRATCH_INIT_ARG 6
180 {"flat_scratch_init", ".amdhsa_user_sgpr_flat_scratch_init", DImode
, -1},
181 #define FLAT_SCRATCH_SEGMENT_SIZE_ARG 7
182 {"private_segment_size", ".amdhsa_user_sgpr_private_segment_size", SImode
, -1},
183 #define WORKGROUP_ID_X_ARG 8
184 {"workgroup_id_X", ".amdhsa_system_sgpr_workgroup_id_x", SImode
, -2},
185 {"workgroup_id_Y", ".amdhsa_system_sgpr_workgroup_id_y", SImode
, -2},
186 {"workgroup_id_Z", ".amdhsa_system_sgpr_workgroup_id_z", SImode
, -2},
187 {"workgroup_info", ".amdhsa_system_sgpr_workgroup_info", SImode
, -1},
188 #define PRIVATE_SEGMENT_WAVE_OFFSET_ARG 12
189 {"private_segment_wave_offset",
190 ".amdhsa_system_sgpr_private_segment_wavefront_offset", SImode
, -2},
191 #define WORK_ITEM_ID_X_ARG 13
192 {"work_item_id_X", NULL
, V64SImode
, FIRST_VGPR_REG
},
193 #define WORK_ITEM_ID_Y_ARG 14
194 {"work_item_id_Y", NULL
, V64SImode
, FIRST_VGPR_REG
+ 1},
195 #define WORK_ITEM_ID_Z_ARG 15
196 {"work_item_id_Z", NULL
, V64SImode
, FIRST_VGPR_REG
+ 2}
199 static const long default_requested_args
200 = (1 << PRIVATE_SEGMENT_BUFFER_ARG
)
201 | (1 << DISPATCH_PTR_ARG
)
202 | (1 << QUEUE_PTR_ARG
)
203 | (1 << KERNARG_SEGMENT_PTR_ARG
)
204 | (1 << PRIVATE_SEGMENT_WAVE_OFFSET_ARG
)
205 | (1 << WORKGROUP_ID_X_ARG
)
206 | (1 << WORK_ITEM_ID_X_ARG
)
207 | (1 << WORK_ITEM_ID_Y_ARG
)
208 | (1 << WORK_ITEM_ID_Z_ARG
);
210 /* Extract parameter settings from __attribute__((amdgpu_hsa_kernel ())).
211 This function also sets the default values for some arguments.
213 Return true on success, with ARGS populated. */
216 gcn_parse_amdgpu_hsa_kernel_attribute (struct gcn_kernel_args
*args
,
220 args
->requested
= default_requested_args
;
223 for (int a
= 0; a
< GCN_KERNEL_ARG_TYPES
; a
++)
226 for (; list
; list
= TREE_CHAIN (list
))
229 if (TREE_CODE (TREE_VALUE (list
)) != STRING_CST
)
231 error ("amdgpu_hsa_kernel attribute requires string constant "
235 str
= TREE_STRING_POINTER (TREE_VALUE (list
));
237 for (a
= 0; a
< GCN_KERNEL_ARG_TYPES
; a
++)
239 if (!strcmp (str
, gcn_kernel_arg_types
[a
].name
))
242 if (a
== GCN_KERNEL_ARG_TYPES
)
244 error ("unknown specifier %s in amdgpu_hsa_kernel attribute", str
);
248 if (args
->requested
& (1 << a
))
250 error ("duplicated parameter specifier %s in amdgpu_hsa_kernel "
255 args
->requested
|= (1 << a
);
256 args
->order
[args
->nargs
++] = a
;
259 /* Requesting WORK_ITEM_ID_Z_ARG implies requesting WORK_ITEM_ID_X_ARG and
260 WORK_ITEM_ID_Y_ARG. Similarly, requesting WORK_ITEM_ID_Y_ARG implies
261 requesting WORK_ITEM_ID_X_ARG. */
262 if (args
->requested
& (1 << WORK_ITEM_ID_Z_ARG
))
263 args
->requested
|= (1 << WORK_ITEM_ID_Y_ARG
);
264 if (args
->requested
& (1 << WORK_ITEM_ID_Y_ARG
))
265 args
->requested
|= (1 << WORK_ITEM_ID_X_ARG
);
267 int sgpr_regno
= FIRST_SGPR_REG
;
269 for (int a
= 0; a
< GCN_KERNEL_ARG_TYPES
; a
++)
271 if (!(args
->requested
& (1 << a
)))
274 if (gcn_kernel_arg_types
[a
].fixed_regno
>= 0)
275 args
->reg
[a
] = gcn_kernel_arg_types
[a
].fixed_regno
;
280 switch (gcn_kernel_arg_types
[a
].mode
)
294 args
->reg
[a
] = sgpr_regno
;
295 sgpr_regno
+= reg_count
;
296 if (gcn_kernel_arg_types
[a
].fixed_regno
== -1)
297 args
->nsgprs
+= reg_count
;
300 if (sgpr_regno
> FIRST_SGPR_REG
+ 16)
302 error ("too many arguments passed in sgpr registers");
307 /* Referenced by TARGET_ATTRIBUTE_TABLE.
309 Validates target specific attributes. */
312 gcn_handle_amdgpu_hsa_kernel_attribute (tree
*node
, tree name
,
313 tree args
, int, bool *no_add_attrs
)
315 if (!FUNC_OR_METHOD_TYPE_P (*node
))
317 warning (OPT_Wattributes
, "%qE attribute only applies to functions",
319 *no_add_attrs
= true;
323 /* Can combine regparm with all attributes but fastcall, and thiscall. */
324 if (is_attribute_p ("gcnhsa_kernel", name
))
326 struct gcn_kernel_args kernelarg
;
328 if (gcn_parse_amdgpu_hsa_kernel_attribute (&kernelarg
, args
))
329 *no_add_attrs
= true;
337 /* Implement TARGET_ATTRIBUTE_TABLE.
339 Create target-specific __attribute__ types. */
341 static const struct attribute_spec gcn_attribute_table
[] = {
342 /* { name, min_len, max_len, decl_req, type_req, fn_type_req, handler,
343 affects_type_identity } */
344 {"amdgpu_hsa_kernel", 0, GCN_KERNEL_ARG_TYPES
, false, true,
345 true, true, gcn_handle_amdgpu_hsa_kernel_attribute
, NULL
},
347 {NULL
, 0, 0, false, false, false, false, NULL
, NULL
}
351 /* {{{ Registers and modes. */
353 /* Implement TARGET_SCALAR_MODE_SUPPORTED_P. */
356 gcn_scalar_mode_supported_p (scalar_mode mode
)
358 return (mode
== BImode
360 || mode
== HImode
/* || mode == HFmode */
361 || mode
== SImode
|| mode
== SFmode
362 || mode
== DImode
|| mode
== DFmode
366 /* Implement TARGET_CLASS_MAX_NREGS.
368 Return the number of hard registers needed to hold a value of MODE in
369 a register of class RCLASS. */
372 gcn_class_max_nregs (reg_class_t rclass
, machine_mode mode
)
374 /* Scalar registers are 32bit, vector registers are in fact tuples of
376 if (rclass
== VGPR_REGS
)
378 if (vgpr_1reg_mode_p (mode
))
380 if (vgpr_2reg_mode_p (mode
))
382 /* TImode is used by DImode compare_and_swap. */
386 else if (rclass
== VCC_CONDITIONAL_REG
&& mode
== BImode
)
388 return CEIL (GET_MODE_SIZE (mode
), 4);
391 /* Implement TARGET_HARD_REGNO_NREGS.
393 Return the number of hard registers needed to hold a value of MODE in
397 gcn_hard_regno_nregs (unsigned int regno
, machine_mode mode
)
399 return gcn_class_max_nregs (REGNO_REG_CLASS (regno
), mode
);
402 /* Implement TARGET_HARD_REGNO_MODE_OK.
404 Return true if REGNO can hold value in MODE. */
407 gcn_hard_regno_mode_ok (unsigned int regno
, machine_mode mode
)
409 /* Treat a complex mode as if it were a scalar mode of the same overall
410 size for the purposes of allocating hard registers. */
411 if (COMPLEX_MODE_P (mode
))
437 case FLAT_SCRATCH_LO_REG
:
438 case XNACK_MASK_LO_REG
:
441 return (mode
== SImode
|| mode
== DImode
);
444 return (mode
== BImode
|| mode
== SImode
|| mode
== DImode
);
446 case FLAT_SCRATCH_HI_REG
:
447 case XNACK_MASK_HI_REG
:
450 return mode
== SImode
;
454 return mode
== SImode
/*|| mode == V32BImode */ ;
458 return mode
== BImode
;
460 if (regno
== ARG_POINTER_REGNUM
|| regno
== FRAME_POINTER_REGNUM
)
462 if (SGPR_REGNO_P (regno
))
463 /* We restrict double register values to aligned registers. */
464 return (sgpr_1reg_mode_p (mode
)
465 || (!((regno
- FIRST_SGPR_REG
) & 1) && sgpr_2reg_mode_p (mode
))
466 || (((regno
- FIRST_SGPR_REG
) & 3) == 0 && mode
== TImode
));
467 if (VGPR_REGNO_P (regno
))
468 /* Vector instructions do not care about the alignment of register
469 pairs, but where there is no 64-bit instruction, many of the
470 define_split do not work if the input and output registers partially
471 overlap. We tried to fix this with early clobber and match
472 constraints, but it was bug prone, added complexity, and conflicts
473 with the 'U0' constraints on vec_merge.
474 Therefore, we restrict ourselved to aligned registers. */
475 return (vgpr_1reg_mode_p (mode
)
476 || (!((regno
- FIRST_VGPR_REG
) & 1) && vgpr_2reg_mode_p (mode
))
477 /* TImode is used by DImode compare_and_swap. */
479 && !((regno
- FIRST_VGPR_REG
) & 3)));
483 /* Implement REGNO_REG_CLASS via gcn.h.
485 Return smallest class containing REGNO. */
488 gcn_regno_reg_class (int regno
)
493 return SCC_CONDITIONAL_REG
;
496 return VCC_CONDITIONAL_REG
;
498 return VCCZ_CONDITIONAL_REG
;
500 return EXECZ_CONDITIONAL_REG
;
503 return EXEC_MASK_REG
;
505 if (VGPR_REGNO_P (regno
))
507 if (SGPR_REGNO_P (regno
))
509 if (regno
< FIRST_VGPR_REG
)
511 if (regno
== ARG_POINTER_REGNUM
|| regno
== FRAME_POINTER_REGNUM
)
516 /* Implement TARGET_CAN_CHANGE_MODE_CLASS.
518 GCC assumes that lowpart contains first part of value as stored in memory.
519 This is not the case for vector registers. */
522 gcn_can_change_mode_class (machine_mode from
, machine_mode to
,
523 reg_class_t regclass
)
525 if (!vgpr_vector_mode_p (from
) && !vgpr_vector_mode_p (to
))
527 return (gcn_class_max_nregs (regclass
, from
)
528 == gcn_class_max_nregs (regclass
, to
));
531 /* Implement TARGET_SMALL_REGISTER_CLASSES_FOR_MODE_P.
533 When this hook returns true for MODE, the compiler allows
534 registers explicitly used in the rtl to be used as spill registers
535 but prevents the compiler from extending the lifetime of these
539 gcn_small_register_classes_for_mode_p (machine_mode mode
)
541 /* We allocate into exec and vcc regs. Those make small register class. */
542 return mode
== DImode
|| mode
== SImode
;
545 /* Implement TARGET_CLASS_LIKELY_SPILLED_P.
547 Returns true if pseudos that have been assigned to registers of class RCLASS
548 would likely be spilled because registers of RCLASS are needed for spill
552 gcn_class_likely_spilled_p (reg_class_t rclass
)
554 return (rclass
== EXEC_MASK_REG
555 || reg_classes_intersect_p (ALL_CONDITIONAL_REGS
, rclass
));
558 /* Implement TARGET_MODES_TIEABLE_P.
560 Returns true if a value of MODE1 is accessible in MODE2 without
564 gcn_modes_tieable_p (machine_mode mode1
, machine_mode mode2
)
566 return (GET_MODE_BITSIZE (mode1
) <= MAX_FIXED_MODE_SIZE
567 && GET_MODE_BITSIZE (mode2
) <= MAX_FIXED_MODE_SIZE
);
570 /* Implement TARGET_TRULY_NOOP_TRUNCATION.
572 Returns true if it is safe to “convert” a value of INPREC bits to one of
573 OUTPREC bits (where OUTPREC is smaller than INPREC) by merely operating on
574 it as if it had only OUTPREC bits. */
577 gcn_truly_noop_truncation (poly_uint64 outprec
, poly_uint64 inprec
)
579 return ((inprec
<= 32) && (outprec
<= inprec
));
582 /* Return N-th part of value occupying multiple registers. */
585 gcn_operand_part (machine_mode mode
, rtx op
, int n
)
587 if (GET_MODE_SIZE (mode
) >= 256)
589 /*gcc_assert (GET_MODE_SIZE (mode) == 256 || n == 0); */
593 gcc_assert (REGNO (op
) + n
< FIRST_PSEUDO_REGISTER
);
594 return gen_rtx_REG (V64SImode
, REGNO (op
) + n
);
596 if (GET_CODE (op
) == CONST_VECTOR
)
598 int units
= GET_MODE_NUNITS (mode
);
599 rtvec v
= rtvec_alloc (units
);
601 for (int i
= 0; i
< units
; ++i
)
602 RTVEC_ELT (v
, i
) = gcn_operand_part (GET_MODE_INNER (mode
),
603 CONST_VECTOR_ELT (op
, i
), n
);
605 return gen_rtx_CONST_VECTOR (V64SImode
, v
);
607 if (GET_CODE (op
) == UNSPEC
&& XINT (op
, 1) == UNSPEC_VECTOR
)
608 return gcn_gen_undef (V64SImode
);
611 else if (GET_MODE_SIZE (mode
) == 8 && REG_P (op
))
613 gcc_assert (REGNO (op
) + n
< FIRST_PSEUDO_REGISTER
);
614 return gen_rtx_REG (SImode
, REGNO (op
) + n
);
618 if (GET_CODE (op
) == UNSPEC
&& XINT (op
, 1) == UNSPEC_VECTOR
)
619 return gcn_gen_undef (SImode
);
621 /* If it's a constant then let's assume it is of the largest mode
622 available, otherwise simplify_gen_subreg will fail. */
623 if (mode
== VOIDmode
&& CONST_INT_P (op
))
625 return simplify_gen_subreg (SImode
, op
, mode
, n
* 4);
629 /* Return N-th part of value occupying multiple registers. */
632 gcn_operand_doublepart (machine_mode mode
, rtx op
, int n
)
634 return simplify_gen_subreg (DImode
, op
, mode
, n
* 8);
637 /* Return true if OP can be split into subregs or high/low parts.
638 This is always true for scalars, but not normally true for vectors.
639 However, for vectors in hardregs we can use the low and high registers. */
642 gcn_can_split_p (machine_mode
, rtx op
)
644 if (vgpr_vector_mode_p (GET_MODE (op
)))
646 if (GET_CODE (op
) == SUBREG
)
647 op
= SUBREG_REG (op
);
650 return REGNO (op
) <= FIRST_PSEUDO_REGISTER
;
655 /* Implement TARGET_SPILL_CLASS.
657 Return class of registers which could be used for pseudo of MODE
658 and of class RCLASS for spilling instead of memory. Return NO_REGS
659 if it is not possible or non-profitable. */
662 gcn_spill_class (reg_class_t c
, machine_mode
/*mode */ )
664 if (reg_classes_intersect_p (ALL_CONDITIONAL_REGS
, c
)
665 || c
== VCC_CONDITIONAL_REG
)
671 /* Implement TARGET_IRA_CHANGE_PSEUDO_ALLOCNO_CLASS.
673 Change allocno class for given pseudo from allocno and best class
674 calculated by IRA. */
677 gcn_ira_change_pseudo_allocno_class (int regno
, reg_class_t cl
,
680 /* Avoid returning classes that contain both vgpr and sgpr registers. */
681 if (cl
!= ALL_REGS
&& cl
!= SRCDST_REGS
&& cl
!= ALL_GPR_REGS
)
683 if (best_cl
!= ALL_REGS
&& best_cl
!= SRCDST_REGS
684 && best_cl
!= ALL_GPR_REGS
)
687 machine_mode mode
= PSEUDO_REGNO_MODE (regno
);
688 if (vgpr_vector_mode_p (mode
))
694 /* Create a new DImode pseudo reg and emit an instruction to initialize
698 get_exec (int64_t val
)
700 rtx reg
= gen_reg_rtx (DImode
);
701 emit_insn (gen_rtx_SET (reg
, gen_int_mode (val
, DImode
)));
705 /* Return value of scalar exec register. */
713 /* Return pseudo holding scalar exec register. */
716 gcn_scalar_exec_reg ()
721 /* Return value of full exec register. */
729 /* Return pseudo holding full exec register. */
734 return get_exec (-1);
738 /* {{{ Immediate constants. */
740 /* Initialize shared numeric constants. */
743 init_ext_gcn_constants (void)
745 real_from_integer (&dconst4
, DFmode
, 4, SIGNED
);
747 /* FIXME: this constant probably does not match what hardware really loads.
748 Reality check it eventually. */
749 real_from_string (&dconst1over2pi
,
750 "0.1591549430918953357663423455968866839");
751 real_convert (&dconst1over2pi
, SFmode
, &dconst1over2pi
);
753 ext_gcn_constants_init
= 1;
756 /* Return non-zero if X is a constant that can appear as an inline operand.
757 This is 0, 0.5, -0.5, 1, -1, 2, -2, 4,-4, 1/(2*pi)
758 Or a vector of those.
759 The value returned should be the encoding of this constant. */
762 gcn_inline_fp_constant_p (rtx x
, bool allow_vector
)
764 machine_mode mode
= GET_MODE (x
);
766 if ((mode
== V64HFmode
|| mode
== V64SFmode
|| mode
== V64DFmode
)
770 if (GET_CODE (x
) != CONST_VECTOR
)
772 n
= gcn_inline_fp_constant_p (CONST_VECTOR_ELT (x
, 0), false);
775 for (int i
= 1; i
< 64; i
++)
776 if (CONST_VECTOR_ELT (x
, i
) != CONST_VECTOR_ELT (x
, 0))
781 if (mode
!= HFmode
&& mode
!= SFmode
&& mode
!= DFmode
)
784 const REAL_VALUE_TYPE
*r
;
786 if (x
== CONST0_RTX (mode
))
788 if (x
== CONST1_RTX (mode
))
791 r
= CONST_DOUBLE_REAL_VALUE (x
);
793 if (real_identical (r
, &dconstm1
))
796 if (real_identical (r
, &dconsthalf
))
798 if (real_identical (r
, &dconstm1
))
800 if (real_identical (r
, &dconst2
))
802 if (real_identical (r
, &dconst4
))
804 if (real_identical (r
, &dconst1over2pi
))
806 if (!ext_gcn_constants_init
)
807 init_ext_gcn_constants ();
808 real_value_negate (r
);
809 if (real_identical (r
, &dconsthalf
))
811 if (real_identical (r
, &dconst2
))
813 if (real_identical (r
, &dconst4
))
816 /* FIXME: add 4, -4 and 1/(2*PI). */
821 /* Return non-zero if X is a constant that can appear as an immediate operand.
822 This is 0, 0.5, -0.5, 1, -1, 2, -2, 4,-4, 1/(2*pi)
823 Or a vector of those.
824 The value returned should be the encoding of this constant. */
827 gcn_fp_constant_p (rtx x
, bool allow_vector
)
829 machine_mode mode
= GET_MODE (x
);
831 if ((mode
== V64HFmode
|| mode
== V64SFmode
|| mode
== V64DFmode
)
835 if (GET_CODE (x
) != CONST_VECTOR
)
837 n
= gcn_fp_constant_p (CONST_VECTOR_ELT (x
, 0), false);
840 for (int i
= 1; i
< 64; i
++)
841 if (CONST_VECTOR_ELT (x
, i
) != CONST_VECTOR_ELT (x
, 0))
845 if (mode
!= HFmode
&& mode
!= SFmode
&& mode
!= DFmode
)
848 if (gcn_inline_fp_constant_p (x
, false))
850 /* FIXME: It is not clear how 32bit immediates are interpreted here. */
851 return (mode
!= DFmode
);
854 /* Return true if X is a constant representable as an inline immediate
855 constant in a 32-bit instruction encoding. */
858 gcn_inline_constant_p (rtx x
)
860 if (GET_CODE (x
) == CONST_INT
)
861 return INTVAL (x
) >= -16 && INTVAL (x
) <= 64;
862 if (GET_CODE (x
) == CONST_DOUBLE
)
863 return gcn_inline_fp_constant_p (x
, false);
864 if (GET_CODE (x
) == CONST_VECTOR
)
867 if (!vgpr_vector_mode_p (GET_MODE (x
)))
869 n
= gcn_inline_constant_p (CONST_VECTOR_ELT (x
, 0));
872 for (int i
= 1; i
< 64; i
++)
873 if (CONST_VECTOR_ELT (x
, i
) != CONST_VECTOR_ELT (x
, 0))
880 /* Return true if X is a constant representable as an immediate constant
881 in a 32 or 64-bit instruction encoding. */
884 gcn_constant_p (rtx x
)
886 switch (GET_CODE (x
))
892 return gcn_fp_constant_p (x
, false);
897 if (!vgpr_vector_mode_p (GET_MODE (x
)))
899 n
= gcn_constant_p (CONST_VECTOR_ELT (x
, 0));
902 for (int i
= 1; i
< 64; i
++)
903 if (CONST_VECTOR_ELT (x
, i
) != CONST_VECTOR_ELT (x
, 0))
919 /* Return true if X is a constant representable as two inline immediate
920 constants in a 64-bit instruction that is split into two 32-bit
922 When MIXED is set, the low-part is permitted to use the full 32-bits. */
925 gcn_inline_constant64_p (rtx x
, bool mixed
)
927 if (GET_CODE (x
) == CONST_VECTOR
)
929 if (!vgpr_vector_mode_p (GET_MODE (x
)))
931 if (!gcn_inline_constant64_p (CONST_VECTOR_ELT (x
, 0), mixed
))
933 for (int i
= 1; i
< 64; i
++)
934 if (CONST_VECTOR_ELT (x
, i
) != CONST_VECTOR_ELT (x
, 0))
940 if (GET_CODE (x
) != CONST_INT
)
943 rtx val_lo
= gcn_operand_part (DImode
, x
, 0);
944 rtx val_hi
= gcn_operand_part (DImode
, x
, 1);
945 return ((mixed
|| gcn_inline_constant_p (val_lo
))
946 && gcn_inline_constant_p (val_hi
));
949 /* Return true if X is a constant representable as an immediate constant
950 in a 32 or 64-bit instruction encoding where the hardware will
951 extend the immediate to 64-bits. */
954 gcn_constant64_p (rtx x
)
956 if (!gcn_constant_p (x
))
959 if (GET_CODE (x
) != CONST_INT
)
962 /* Negative numbers are only allowed if they can be encoded within src0,
963 because the 32-bit immediates do not get sign-extended.
964 Unsigned numbers must not be encodable as 32-bit -1..-16, because the
965 assembler will use a src0 inline immediate and that will get
967 HOST_WIDE_INT val
= INTVAL (x
);
968 return (((val
& 0xffffffff) == val
/* Positive 32-bit. */
969 && (val
& 0xfffffff0) != 0xfffffff0) /* Not -1..-16. */
970 || gcn_inline_constant_p (x
)); /* Src0. */
973 /* Implement TARGET_LEGITIMATE_CONSTANT_P.
975 Returns true if X is a legitimate constant for a MODE immediate operand. */
978 gcn_legitimate_constant_p (machine_mode
, rtx x
)
980 return gcn_constant_p (x
);
983 /* Return true if X is a CONST_VECTOR of single constant. */
986 single_cst_vector_p (rtx x
)
988 if (GET_CODE (x
) != CONST_VECTOR
)
990 for (int i
= 1; i
< 64; i
++)
991 if (CONST_VECTOR_ELT (x
, i
) != CONST_VECTOR_ELT (x
, 0))
996 /* Create a CONST_VECTOR of duplicated value A. */
999 gcn_vec_constant (machine_mode mode
, int a
)
1002 return CONST0_RTX (mode);
1004 return CONSTM1_RTX (mode);
1006 return CONST1_RTX (mode);
1008 return CONST2_RTX (mode);*/
1010 int units
= GET_MODE_NUNITS (mode
);
1011 machine_mode innermode
= GET_MODE_INNER (mode
);
1014 if (FLOAT_MODE_P (innermode
))
1017 real_from_integer (&rv
, NULL
, a
, SIGNED
);
1018 tem
= const_double_from_real_value (rv
, innermode
);
1021 tem
= gen_int_mode (a
, innermode
);
1023 rtvec v
= rtvec_alloc (units
);
1024 for (int i
= 0; i
< units
; ++i
)
1025 RTVEC_ELT (v
, i
) = tem
;
1027 return gen_rtx_CONST_VECTOR (mode
, v
);
1030 /* Create a CONST_VECTOR of duplicated value A. */
1033 gcn_vec_constant (machine_mode mode
, rtx a
)
1035 int units
= GET_MODE_NUNITS (mode
);
1036 rtvec v
= rtvec_alloc (units
);
1038 for (int i
= 0; i
< units
; ++i
)
1039 RTVEC_ELT (v
, i
) = a
;
1041 return gen_rtx_CONST_VECTOR (mode
, v
);
1044 /* Create an undefined vector value, used where an insn operand is
1048 gcn_gen_undef (machine_mode mode
)
1050 return gen_rtx_UNSPEC (mode
, gen_rtvec (1, const0_rtx
), UNSPEC_VECTOR
);
1054 /* {{{ Addresses, pointers and moves. */
1056 /* Return true is REG is a valid place to store a pointer,
1057 for instructions that require an SGPR.
1061 gcn_address_register_p (rtx reg
, machine_mode mode
, bool strict
)
1063 if (GET_CODE (reg
) == SUBREG
)
1064 reg
= SUBREG_REG (reg
);
1069 if (GET_MODE (reg
) != mode
)
1072 int regno
= REGNO (reg
);
1074 if (regno
>= FIRST_PSEUDO_REGISTER
)
1082 regno
= reg_renumber
[regno
];
1085 return (SGPR_REGNO_P (regno
) || regno
== M0_REG
1086 || regno
== ARG_POINTER_REGNUM
|| regno
== FRAME_POINTER_REGNUM
);
1089 /* Return true is REG is a valid place to store a pointer,
1090 for instructions that require a VGPR. */
1093 gcn_vec_address_register_p (rtx reg
, machine_mode mode
, bool strict
)
1095 if (GET_CODE (reg
) == SUBREG
)
1096 reg
= SUBREG_REG (reg
);
1101 if (GET_MODE (reg
) != mode
)
1104 int regno
= REGNO (reg
);
1106 if (regno
>= FIRST_PSEUDO_REGISTER
)
1114 regno
= reg_renumber
[regno
];
1117 return VGPR_REGNO_P (regno
);
1120 /* Return true if X would be valid inside a MEM using the Flat address
1124 gcn_flat_address_p (rtx x
, machine_mode mode
)
1126 bool vec_mode
= (GET_MODE_CLASS (mode
) == MODE_VECTOR_INT
1127 || GET_MODE_CLASS (mode
) == MODE_VECTOR_FLOAT
);
1129 if (vec_mode
&& gcn_address_register_p (x
, DImode
, false))
1132 if (!vec_mode
&& gcn_vec_address_register_p (x
, DImode
, false))
1135 if (TARGET_GCN5_PLUS
1136 && GET_CODE (x
) == PLUS
1137 && gcn_vec_address_register_p (XEXP (x
, 0), DImode
, false)
1138 && CONST_INT_P (XEXP (x
, 1)))
1144 /* Return true if X would be valid inside a MEM using the Scalar Flat
1148 gcn_scalar_flat_address_p (rtx x
)
1150 if (gcn_address_register_p (x
, DImode
, false))
1153 if (GET_CODE (x
) == PLUS
1154 && gcn_address_register_p (XEXP (x
, 0), DImode
, false)
1155 && CONST_INT_P (XEXP (x
, 1)))
1161 /* Return true if MEM X would be valid for the Scalar Flat address space. */
1164 gcn_scalar_flat_mem_p (rtx x
)
1169 if (GET_MODE_SIZE (GET_MODE (x
)) < 4)
1172 return gcn_scalar_flat_address_p (XEXP (x
, 0));
1175 /* Return true if X would be valid inside a MEM using the LDS or GDS
1179 gcn_ds_address_p (rtx x
)
1181 if (gcn_vec_address_register_p (x
, SImode
, false))
1184 if (GET_CODE (x
) == PLUS
1185 && gcn_vec_address_register_p (XEXP (x
, 0), SImode
, false)
1186 && CONST_INT_P (XEXP (x
, 1)))
1192 /* Return true if ADDR would be valid inside a MEM using the Global
1196 gcn_global_address_p (rtx addr
)
1198 if (gcn_address_register_p (addr
, DImode
, false)
1199 || gcn_vec_address_register_p (addr
, DImode
, false))
1202 if (GET_CODE (addr
) == PLUS
)
1204 rtx base
= XEXP (addr
, 0);
1205 rtx offset
= XEXP (addr
, 1);
1206 bool immediate_p
= (CONST_INT_P (offset
)
1207 && INTVAL (offset
) >= -(1 << 12)
1208 && INTVAL (offset
) < (1 << 12));
1210 if ((gcn_address_register_p (base
, DImode
, false)
1211 || gcn_vec_address_register_p (base
, DImode
, false))
1213 /* SGPR + CONST or VGPR + CONST */
1216 if (gcn_address_register_p (base
, DImode
, false)
1217 && gcn_vgpr_register_operand (offset
, SImode
))
1221 if (GET_CODE (base
) == PLUS
1222 && gcn_address_register_p (XEXP (base
, 0), DImode
, false)
1223 && gcn_vgpr_register_operand (XEXP (base
, 1), SImode
)
1225 /* (SGPR + VGPR) + CONST */
1232 /* Implement TARGET_ADDR_SPACE_LEGITIMATE_ADDRESS_P.
1234 Recognizes RTL expressions that are valid memory addresses for an
1235 instruction. The MODE argument is the machine mode for the MEM
1236 expression that wants to use this address.
1238 It only recognizes address in canonical form. LEGITIMIZE_ADDRESS should
1239 convert common non-canonical forms to canonical form so that they will
1243 gcn_addr_space_legitimate_address_p (machine_mode mode
, rtx x
, bool strict
,
1246 /* All vector instructions need to work on addresses in registers. */
1247 if (!TARGET_GCN5_PLUS
&& (vgpr_vector_mode_p (mode
) && !REG_P (x
)))
1250 if (AS_SCALAR_FLAT_P (as
))
1252 if (mode
== QImode
|| mode
== HImode
)
1255 switch (GET_CODE (x
))
1258 return gcn_address_register_p (x
, DImode
, strict
);
1259 /* Addresses are in the form BASE+OFFSET
1260 OFFSET is either 20bit unsigned immediate, SGPR or M0.
1261 Writes and atomics do not accept SGPR. */
1264 rtx x0
= XEXP (x
, 0);
1265 rtx x1
= XEXP (x
, 1);
1266 if (!gcn_address_register_p (x0
, DImode
, strict
))
1268 /* FIXME: This is disabled because of the mode mismatch between
1269 SImode (for the address or m0 register) and the DImode PLUS.
1270 We'll need a zero_extend or similar.
1272 if (gcn_m0_register_p (x1, SImode, strict)
1273 || gcn_address_register_p (x1, SImode, strict))
1276 if (GET_CODE (x1
) == CONST_INT
)
1278 if (INTVAL (x1
) >= 0 && INTVAL (x1
) < (1 << 20)
1279 /* The low bits of the offset are ignored, even when
1280 they're meant to realign the pointer. */
1281 && !(INTVAL (x1
) & 0x3))
1291 else if (AS_SCRATCH_P (as
))
1292 return gcn_address_register_p (x
, SImode
, strict
);
1293 else if (AS_FLAT_P (as
) || AS_FLAT_SCRATCH_P (as
))
1295 if (TARGET_GCN3
|| GET_CODE (x
) == REG
)
1296 return ((GET_MODE_CLASS (mode
) == MODE_VECTOR_INT
1297 || GET_MODE_CLASS (mode
) == MODE_VECTOR_FLOAT
)
1298 ? gcn_address_register_p (x
, DImode
, strict
)
1299 : gcn_vec_address_register_p (x
, DImode
, strict
));
1302 gcc_assert (TARGET_GCN5_PLUS
);
1304 if (GET_CODE (x
) == PLUS
)
1306 rtx x1
= XEXP (x
, 1);
1308 if (VECTOR_MODE_P (mode
)
1309 ? !gcn_address_register_p (x
, DImode
, strict
)
1310 : !gcn_vec_address_register_p (x
, DImode
, strict
))
1313 if (GET_CODE (x1
) == CONST_INT
)
1315 if (INTVAL (x1
) >= 0 && INTVAL (x1
) < (1 << 12)
1316 /* The low bits of the offset are ignored, even when
1317 they're meant to realign the pointer. */
1318 && !(INTVAL (x1
) & 0x3))
1325 else if (AS_GLOBAL_P (as
))
1327 gcc_assert (TARGET_GCN5_PLUS
);
1329 if (GET_CODE (x
) == REG
)
1330 return (gcn_address_register_p (x
, DImode
, strict
)
1331 || (!VECTOR_MODE_P (mode
)
1332 && gcn_vec_address_register_p (x
, DImode
, strict
)));
1333 else if (GET_CODE (x
) == PLUS
)
1335 rtx base
= XEXP (x
, 0);
1336 rtx offset
= XEXP (x
, 1);
1338 bool immediate_p
= (GET_CODE (offset
) == CONST_INT
1339 /* Signed 13-bit immediate. */
1340 && INTVAL (offset
) >= -(1 << 12)
1341 && INTVAL (offset
) < (1 << 12)
1342 /* The low bits of the offset are ignored, even
1343 when they're meant to realign the pointer. */
1344 && !(INTVAL (offset
) & 0x3));
1346 if (!VECTOR_MODE_P (mode
))
1348 if ((gcn_address_register_p (base
, DImode
, strict
)
1349 || gcn_vec_address_register_p (base
, DImode
, strict
))
1351 /* SGPR + CONST or VGPR + CONST */
1354 if (gcn_address_register_p (base
, DImode
, strict
)
1355 && gcn_vgpr_register_operand (offset
, SImode
))
1359 if (GET_CODE (base
) == PLUS
1360 && gcn_address_register_p (XEXP (base
, 0), DImode
, strict
)
1361 && gcn_vgpr_register_operand (XEXP (base
, 1), SImode
)
1363 /* (SGPR + VGPR) + CONST */
1368 if (gcn_address_register_p (base
, DImode
, strict
)
1377 else if (AS_ANY_DS_P (as
))
1378 switch (GET_CODE (x
))
1381 return (VECTOR_MODE_P (mode
)
1382 ? gcn_address_register_p (x
, SImode
, strict
)
1383 : gcn_vec_address_register_p (x
, SImode
, strict
));
1384 /* Addresses are in the form BASE+OFFSET
1385 OFFSET is either 20bit unsigned immediate, SGPR or M0.
1386 Writes and atomics do not accept SGPR. */
1389 rtx x0
= XEXP (x
, 0);
1390 rtx x1
= XEXP (x
, 1);
1391 if (!gcn_vec_address_register_p (x0
, DImode
, strict
))
1393 if (GET_CODE (x1
) == REG
)
1395 if (GET_CODE (x1
) != REG
1396 || (REGNO (x1
) <= FIRST_PSEUDO_REGISTER
1397 && !gcn_ssrc_register_operand (x1
, DImode
)))
1400 else if (GET_CODE (x1
) == CONST_VECTOR
1401 && GET_CODE (CONST_VECTOR_ELT (x1
, 0)) == CONST_INT
1402 && single_cst_vector_p (x1
))
1404 x1
= CONST_VECTOR_ELT (x1
, 0);
1405 if (INTVAL (x1
) >= 0 && INTVAL (x1
) < (1 << 20))
1419 /* Implement TARGET_ADDR_SPACE_POINTER_MODE.
1421 Return the appropriate mode for a named address pointer. */
1423 static scalar_int_mode
1424 gcn_addr_space_pointer_mode (addr_space_t addrspace
)
1428 case ADDR_SPACE_SCRATCH
:
1429 case ADDR_SPACE_LDS
:
1430 case ADDR_SPACE_GDS
:
1432 case ADDR_SPACE_DEFAULT
:
1433 case ADDR_SPACE_FLAT
:
1434 case ADDR_SPACE_FLAT_SCRATCH
:
1435 case ADDR_SPACE_SCALAR_FLAT
:
1442 /* Implement TARGET_ADDR_SPACE_ADDRESS_MODE.
1444 Return the appropriate mode for a named address space address. */
1446 static scalar_int_mode
1447 gcn_addr_space_address_mode (addr_space_t addrspace
)
1449 return gcn_addr_space_pointer_mode (addrspace
);
1452 /* Implement TARGET_ADDR_SPACE_SUBSET_P.
1454 Determine if one named address space is a subset of another. */
1457 gcn_addr_space_subset_p (addr_space_t subset
, addr_space_t superset
)
1459 if (subset
== superset
)
1461 /* FIXME is this true? */
1462 if (AS_FLAT_P (superset
) || AS_SCALAR_FLAT_P (superset
))
1467 /* Convert from one address space to another. */
1470 gcn_addr_space_convert (rtx op
, tree from_type
, tree to_type
)
1472 gcc_assert (POINTER_TYPE_P (from_type
));
1473 gcc_assert (POINTER_TYPE_P (to_type
));
1475 addr_space_t as_from
= TYPE_ADDR_SPACE (TREE_TYPE (from_type
));
1476 addr_space_t as_to
= TYPE_ADDR_SPACE (TREE_TYPE (to_type
));
1478 if (AS_LDS_P (as_from
) && AS_FLAT_P (as_to
))
1480 rtx queue
= gen_rtx_REG (DImode
,
1481 cfun
->machine
->args
.reg
[QUEUE_PTR_ARG
]);
1482 rtx group_seg_aperture_hi
= gen_rtx_MEM (SImode
,
1483 gen_rtx_PLUS (DImode
, queue
,
1484 gen_int_mode (64, SImode
)));
1485 rtx tmp
= gen_reg_rtx (DImode
);
1487 emit_move_insn (gen_lowpart (SImode
, tmp
), op
);
1488 emit_move_insn (gen_highpart_mode (SImode
, DImode
, tmp
),
1489 group_seg_aperture_hi
);
1493 else if (as_from
== as_to
)
1500 /* Implement REGNO_MODE_CODE_OK_FOR_BASE_P via gcn.h
1502 Retun true if REGNO is OK for memory adressing. */
1505 gcn_regno_mode_code_ok_for_base_p (int regno
,
1506 machine_mode
, addr_space_t as
, int, int)
1508 if (regno
>= FIRST_PSEUDO_REGISTER
)
1511 regno
= reg_renumber
[regno
];
1516 return (VGPR_REGNO_P (regno
)
1517 || regno
== ARG_POINTER_REGNUM
|| regno
== FRAME_POINTER_REGNUM
);
1518 else if (AS_SCALAR_FLAT_P (as
))
1519 return (SGPR_REGNO_P (regno
)
1520 || regno
== ARG_POINTER_REGNUM
|| regno
== FRAME_POINTER_REGNUM
);
1521 else if (AS_GLOBAL_P (as
))
1523 return (SGPR_REGNO_P (regno
)
1524 || VGPR_REGNO_P (regno
)
1525 || regno
== ARG_POINTER_REGNUM
1526 || regno
== FRAME_POINTER_REGNUM
);
1533 /* Implement MODE_CODE_BASE_REG_CLASS via gcn.h.
1535 Return a suitable register class for memory addressing. */
1538 gcn_mode_code_base_reg_class (machine_mode mode
, addr_space_t as
, int oc
,
1543 case ADDR_SPACE_DEFAULT
:
1544 return gcn_mode_code_base_reg_class (mode
, DEFAULT_ADDR_SPACE
, oc
, ic
);
1545 case ADDR_SPACE_SCALAR_FLAT
:
1546 case ADDR_SPACE_SCRATCH
:
1549 case ADDR_SPACE_FLAT
:
1550 case ADDR_SPACE_FLAT_SCRATCH
:
1551 case ADDR_SPACE_LDS
:
1552 case ADDR_SPACE_GDS
:
1553 return ((GET_MODE_CLASS (mode
) == MODE_VECTOR_INT
1554 || GET_MODE_CLASS (mode
) == MODE_VECTOR_FLOAT
)
1555 ? SGPR_REGS
: VGPR_REGS
);
1556 case ADDR_SPACE_GLOBAL
:
1557 return ((GET_MODE_CLASS (mode
) == MODE_VECTOR_INT
1558 || GET_MODE_CLASS (mode
) == MODE_VECTOR_FLOAT
)
1559 ? SGPR_REGS
: ALL_GPR_REGS
);
1564 /* Implement REGNO_OK_FOR_INDEX_P via gcn.h.
1566 Return true if REGNO is OK for index of memory addressing. */
1569 regno_ok_for_index_p (int regno
)
1571 if (regno
>= FIRST_PSEUDO_REGISTER
)
1574 regno
= reg_renumber
[regno
];
1578 return regno
== M0_REG
|| VGPR_REGNO_P (regno
);
1581 /* Generate move which uses the exec flags. If EXEC is NULL, then it is
1582 assumed that all lanes normally relevant to the mode of the move are
1583 affected. If PREV is NULL, then a sensible default is supplied for
1584 the inactive lanes. */
1587 gen_mov_with_exec (rtx op0
, rtx op1
, rtx exec
= NULL
, rtx prev
= NULL
)
1589 machine_mode mode
= GET_MODE (op0
);
1591 if (vgpr_vector_mode_p (mode
))
1593 if (exec
&& exec
!= CONSTM1_RTX (DImode
))
1601 prev
= gcn_gen_undef (mode
);
1602 exec
= gcn_full_exec_reg ();
1605 rtx set
= gen_rtx_SET (op0
, gen_rtx_VEC_MERGE (mode
, op1
, prev
, exec
));
1607 return gen_rtx_PARALLEL (VOIDmode
,
1609 gen_rtx_CLOBBER (VOIDmode
,
1610 gen_rtx_SCRATCH (V64DImode
))));
1613 return (gen_rtx_PARALLEL
1615 gen_rtvec (2, gen_rtx_SET (op0
, op1
),
1616 gen_rtx_USE (VOIDmode
,
1617 exec
? exec
: gcn_scalar_exec ()))));
1620 /* Generate masked move. */
1623 gen_duplicate_load (rtx op0
, rtx op1
, rtx op2
= NULL
, rtx exec
= NULL
)
1626 return (gen_rtx_SET (op0
,
1627 gen_rtx_VEC_MERGE (GET_MODE (op0
),
1628 gen_rtx_VEC_DUPLICATE (GET_MODE
1632 return (gen_rtx_SET (op0
, gen_rtx_VEC_DUPLICATE (GET_MODE (op0
), op1
)));
1635 /* Expand vector init of OP0 by VEC.
1636 Implements vec_init instruction pattern. */
1639 gcn_expand_vector_init (rtx op0
, rtx vec
)
1641 int64_t initialized_mask
= 0;
1642 int64_t curr_mask
= 1;
1643 machine_mode mode
= GET_MODE (op0
);
1645 rtx val
= XVECEXP (vec
, 0, 0);
1647 for (int i
= 1; i
< 64; i
++)
1648 if (rtx_equal_p (val
, XVECEXP (vec
, 0, i
)))
1649 curr_mask
|= (int64_t) 1 << i
;
1651 if (gcn_constant_p (val
))
1652 emit_move_insn (op0
, gcn_vec_constant (mode
, val
));
1655 val
= force_reg (GET_MODE_INNER (mode
), val
);
1656 emit_insn (gen_duplicate_load (op0
, val
));
1658 initialized_mask
|= curr_mask
;
1659 for (int i
= 1; i
< 64; i
++)
1660 if (!(initialized_mask
& ((int64_t) 1 << i
)))
1662 curr_mask
= (int64_t) 1 << i
;
1663 rtx val
= XVECEXP (vec
, 0, i
);
1665 for (int j
= i
+ 1; j
< 64; j
++)
1666 if (rtx_equal_p (val
, XVECEXP (vec
, 0, j
)))
1667 curr_mask
|= (int64_t) 1 << j
;
1668 if (gcn_constant_p (val
))
1669 emit_insn (gen_mov_with_exec (op0
, gcn_vec_constant (mode
, val
),
1670 get_exec (curr_mask
)));
1673 val
= force_reg (GET_MODE_INNER (mode
), val
);
1674 emit_insn (gen_duplicate_load (op0
, val
, op0
,
1675 get_exec (curr_mask
)));
1677 initialized_mask
|= curr_mask
;
1681 /* Load vector constant where n-th lane contains BASE+n*VAL. */
1684 strided_constant (machine_mode mode
, int base
, int val
)
1686 rtx x
= gen_reg_rtx (mode
);
1687 emit_move_insn (x
, gcn_vec_constant (mode
, base
));
1688 emit_insn (gen_addv64si3_exec (x
, x
, gcn_vec_constant (mode
, val
* 32),
1689 x
, get_exec (0xffffffff00000000)));
1690 emit_insn (gen_addv64si3_exec (x
, x
, gcn_vec_constant (mode
, val
* 16),
1691 x
, get_exec (0xffff0000ffff0000)));
1692 emit_insn (gen_addv64si3_exec (x
, x
, gcn_vec_constant (mode
, val
* 8),
1693 x
, get_exec (0xff00ff00ff00ff00)));
1694 emit_insn (gen_addv64si3_exec (x
, x
, gcn_vec_constant (mode
, val
* 4),
1695 x
, get_exec (0xf0f0f0f0f0f0f0f0)));
1696 emit_insn (gen_addv64si3_exec (x
, x
, gcn_vec_constant (mode
, val
* 2),
1697 x
, get_exec (0xcccccccccccccccc)));
1698 emit_insn (gen_addv64si3_exec (x
, x
, gcn_vec_constant (mode
, val
* 1),
1699 x
, get_exec (0xaaaaaaaaaaaaaaaa)));
1703 /* Implement TARGET_ADDR_SPACE_LEGITIMIZE_ADDRESS. */
1706 gcn_addr_space_legitimize_address (rtx x
, rtx old
, machine_mode mode
,
1711 case ADDR_SPACE_DEFAULT
:
1712 return gcn_addr_space_legitimize_address (x
, old
, mode
,
1713 DEFAULT_ADDR_SPACE
);
1714 case ADDR_SPACE_SCALAR_FLAT
:
1715 case ADDR_SPACE_SCRATCH
:
1716 /* Instructions working on vectors need the address to be in
1718 if (vgpr_vector_mode_p (mode
))
1719 return force_reg (GET_MODE (x
), x
);
1722 case ADDR_SPACE_FLAT
:
1723 case ADDR_SPACE_FLAT_SCRATCH
:
1724 case ADDR_SPACE_GLOBAL
:
1725 return TARGET_GCN3
? force_reg (DImode
, x
) : x
;
1726 case ADDR_SPACE_LDS
:
1727 case ADDR_SPACE_GDS
:
1728 /* FIXME: LDS support offsets, handle them!. */
1729 if (vgpr_vector_mode_p (mode
) && GET_MODE (x
) != V64SImode
)
1731 rtx addrs
= gen_reg_rtx (V64SImode
);
1732 rtx base
= force_reg (SImode
, x
);
1733 rtx offsets
= strided_constant (V64SImode
, 0,
1734 GET_MODE_UNIT_SIZE (mode
));
1736 emit_insn (gen_vec_duplicatev64si (addrs
, base
));
1737 emit_insn (gen_addv64si3 (addrs
, offsets
, addrs
));
1745 /* Convert a (mem:<MODE> (reg:DI)) to (mem:<MODE> (reg:V64DI)) with the
1746 proper vector of stepped addresses.
1748 MEM will be a DImode address of a vector in an SGPR.
1749 TMP will be a V64DImode VGPR pair or (scratch:V64DI). */
1752 gcn_expand_scalar_to_vector_address (machine_mode mode
, rtx exec
, rtx mem
,
1755 gcc_assert (MEM_P (mem
));
1756 rtx mem_base
= XEXP (mem
, 0);
1757 rtx mem_index
= NULL_RTX
;
1759 if (!TARGET_GCN5_PLUS
)
1761 /* gcn_addr_space_legitimize_address should have put the address in a
1762 register. If not, it is too late to do anything about it. */
1763 gcc_assert (REG_P (mem_base
));
1766 if (GET_CODE (mem_base
) == PLUS
)
1768 mem_index
= XEXP (mem_base
, 1);
1769 mem_base
= XEXP (mem_base
, 0);
1772 /* RF and RM base registers for vector modes should be always an SGPR. */
1773 gcc_assert (SGPR_REGNO_P (REGNO (mem_base
))
1774 || REGNO (mem_base
) >= FIRST_PSEUDO_REGISTER
);
1776 machine_mode inner
= GET_MODE_INNER (mode
);
1777 int shift
= exact_log2 (GET_MODE_SIZE (inner
));
1778 rtx ramp
= gen_rtx_REG (V64SImode
, VGPR_REGNO (1));
1779 rtx undef_v64si
= gcn_gen_undef (V64SImode
);
1780 rtx new_base
= NULL_RTX
;
1781 addr_space_t as
= MEM_ADDR_SPACE (mem
);
1783 rtx tmplo
= (REG_P (tmp
)
1784 ? gcn_operand_part (V64DImode
, tmp
, 0)
1785 : gen_reg_rtx (V64SImode
));
1787 /* tmplo[:] = ramp[:] << shift */
1789 emit_insn (gen_ashlv64si3_exec (tmplo
, ramp
,
1790 gen_int_mode (shift
, SImode
),
1791 undef_v64si
, exec
));
1793 emit_insn (gen_ashlv64si3 (tmplo
, ramp
, gen_int_mode (shift
, SImode
)));
1797 rtx vcc
= gen_rtx_REG (DImode
, CC_SAVE_REG
);
1801 rtx mem_base_lo
= gcn_operand_part (DImode
, mem_base
, 0);
1802 rtx mem_base_hi
= gcn_operand_part (DImode
, mem_base
, 1);
1803 rtx tmphi
= gcn_operand_part (V64DImode
, tmp
, 1);
1805 /* tmphi[:] = mem_base_hi */
1807 emit_insn (gen_vec_duplicatev64si_exec (tmphi
, mem_base_hi
,
1808 undef_v64si
, exec
));
1810 emit_insn (gen_vec_duplicatev64si (tmphi
, mem_base_hi
));
1812 /* tmp[:] += zext (mem_base) */
1815 emit_insn (gen_addv64si3_vcc_dup_exec (tmplo
, mem_base_lo
, tmplo
,
1816 vcc
, undef_v64si
, exec
));
1817 emit_insn (gen_addcv64si3_exec (tmphi
, tmphi
, const0_rtx
,
1818 vcc
, vcc
, undef_v64si
, exec
));
1821 emit_insn (gen_addv64di3_vcc_zext_dup (tmp
, mem_base_lo
, tmp
, vcc
));
1825 tmp
= gen_reg_rtx (V64DImode
);
1827 emit_insn (gen_addv64di3_vcc_zext_dup2_exec
1828 (tmp
, tmplo
, mem_base
, vcc
, gcn_gen_undef (V64DImode
),
1831 emit_insn (gen_addv64di3_vcc_zext_dup2 (tmp
, tmplo
, mem_base
, vcc
));
1836 else if (AS_ANY_DS_P (as
))
1839 emit_insn (gen_addv64si3_dup (tmplo
, tmplo
, mem_base
));
1841 emit_insn (gen_addv64si3_dup_exec (tmplo
, tmplo
, mem_base
,
1842 gcn_gen_undef (V64SImode
), exec
));
1847 mem_base
= gen_rtx_VEC_DUPLICATE (V64DImode
, mem_base
);
1848 new_base
= gen_rtx_PLUS (V64DImode
, mem_base
,
1849 gen_rtx_SIGN_EXTEND (V64DImode
, tmplo
));
1852 return gen_rtx_PLUS (GET_MODE (new_base
), new_base
,
1853 gen_rtx_VEC_DUPLICATE (GET_MODE (new_base
),
1854 (mem_index
? mem_index
1858 /* Convert a BASE address, a vector of OFFSETS, and a SCALE, to addresses
1859 suitable for the given address space. This is indented for use in
1860 gather/scatter patterns.
1862 The offsets may be signed or unsigned, according to UNSIGNED_P.
1863 If EXEC is set then _exec patterns will be used, otherwise plain.
1866 ADDR_SPACE_FLAT - return V64DImode vector of absolute addresses.
1867 ADDR_SPACE_GLOBAL - return V64SImode vector of offsets. */
1870 gcn_expand_scaled_offsets (addr_space_t as
, rtx base
, rtx offsets
, rtx scale
,
1871 bool unsigned_p
, rtx exec
)
1873 rtx tmpsi
= gen_reg_rtx (V64SImode
);
1874 rtx tmpdi
= gen_reg_rtx (V64DImode
);
1875 rtx undefsi
= exec
? gcn_gen_undef (V64SImode
) : NULL
;
1876 rtx undefdi
= exec
? gcn_gen_undef (V64DImode
) : NULL
;
1878 if (CONST_INT_P (scale
)
1879 && INTVAL (scale
) > 0
1880 && exact_log2 (INTVAL (scale
)) >= 0)
1881 emit_insn (gen_ashlv64si3 (tmpsi
, offsets
,
1882 GEN_INT (exact_log2 (INTVAL (scale
)))));
1885 ? emit_insn (gen_mulv64si3_dup_exec (tmpsi
, offsets
, scale
, undefsi
,
1887 : emit_insn (gen_mulv64si3_dup (tmpsi
, offsets
, scale
)));
1889 /* "Global" instructions do not support negative register offsets. */
1890 if (as
== ADDR_SPACE_FLAT
|| !unsigned_p
)
1894 ? emit_insn (gen_addv64di3_zext_dup2_exec (tmpdi
, tmpsi
, base
,
1896 : emit_insn (gen_addv64di3_zext_dup2 (tmpdi
, tmpsi
, base
)));
1899 ? emit_insn (gen_addv64di3_sext_dup2_exec (tmpdi
, tmpsi
, base
,
1901 : emit_insn (gen_addv64di3_sext_dup2 (tmpdi
, tmpsi
, base
)));
1904 else if (as
== ADDR_SPACE_GLOBAL
)
1910 /* Return true if move from OP0 to OP1 is known to be executed in vector
1914 gcn_vgpr_move_p (rtx op0
, rtx op1
)
1916 if (MEM_P (op0
) && AS_SCALAR_FLAT_P (MEM_ADDR_SPACE (op0
)))
1918 if (MEM_P (op1
) && AS_SCALAR_FLAT_P (MEM_ADDR_SPACE (op1
)))
1920 return ((REG_P (op0
) && VGPR_REGNO_P (REGNO (op0
)))
1921 || (REG_P (op1
) && VGPR_REGNO_P (REGNO (op1
)))
1922 || vgpr_vector_mode_p (GET_MODE (op0
)));
1925 /* Return true if move from OP0 to OP1 is known to be executed in scalar
1926 unit. Used in the machine description. */
1929 gcn_sgpr_move_p (rtx op0
, rtx op1
)
1931 if (MEM_P (op0
) && AS_SCALAR_FLAT_P (MEM_ADDR_SPACE (op0
)))
1933 if (MEM_P (op1
) && AS_SCALAR_FLAT_P (MEM_ADDR_SPACE (op1
)))
1935 if (!REG_P (op0
) || REGNO (op0
) >= FIRST_PSEUDO_REGISTER
1936 || VGPR_REGNO_P (REGNO (op0
)))
1939 && REGNO (op1
) < FIRST_PSEUDO_REGISTER
1940 && !VGPR_REGNO_P (REGNO (op1
)))
1942 return immediate_operand (op1
, VOIDmode
) || memory_operand (op1
, VOIDmode
);
1945 /* Implement TARGET_SECONDARY_RELOAD.
1947 The address space determines which registers can be used for loads and
1951 gcn_secondary_reload (bool in_p
, rtx x
, reg_class_t rclass
,
1952 machine_mode reload_mode
, secondary_reload_info
*sri
)
1954 reg_class_t result
= NO_REGS
;
1955 bool spilled_pseudo
=
1956 (REG_P (x
) || GET_CODE (x
) == SUBREG
) && true_regnum (x
) == -1;
1958 if (dump_file
&& (dump_flags
& TDF_DETAILS
))
1960 fprintf (dump_file
, "gcn_secondary_reload: ");
1961 dump_value_slim (dump_file
, x
, 1);
1962 fprintf (dump_file
, " %s %s:%s", (in_p
? "->" : "<-"),
1963 reg_class_names
[rclass
], GET_MODE_NAME (reload_mode
));
1964 if (REG_P (x
) || GET_CODE (x
) == SUBREG
)
1965 fprintf (dump_file
, " (true regnum: %d \"%s\")", true_regnum (x
),
1966 (true_regnum (x
) >= 0
1967 && true_regnum (x
) < FIRST_PSEUDO_REGISTER
1968 ? reg_names
[true_regnum (x
)]
1969 : (spilled_pseudo
? "stack spill" : "??")));
1970 fprintf (dump_file
, "\n");
1973 /* Some callers don't use or initialize icode. */
1974 sri
->icode
= CODE_FOR_nothing
;
1976 if (MEM_P (x
) || spilled_pseudo
)
1978 addr_space_t as
= DEFAULT_ADDR_SPACE
;
1980 /* If we have a spilled pseudo, we can't find the address space
1981 directly, but we know it's in ADDR_SPACE_FLAT space for GCN3 or
1982 ADDR_SPACE_GLOBAL for GCN5. */
1984 as
= MEM_ADDR_SPACE (x
);
1986 if (as
== ADDR_SPACE_DEFAULT
)
1987 as
= DEFAULT_ADDR_SPACE
;
1991 case ADDR_SPACE_SCALAR_FLAT
:
1993 ((!MEM_P (x
) || rclass
== SGPR_REGS
) ? NO_REGS
: SGPR_REGS
);
1995 case ADDR_SPACE_FLAT
:
1996 case ADDR_SPACE_FLAT_SCRATCH
:
1997 case ADDR_SPACE_GLOBAL
:
1998 if (GET_MODE_CLASS (reload_mode
) == MODE_VECTOR_INT
1999 || GET_MODE_CLASS (reload_mode
) == MODE_VECTOR_FLOAT
)
2002 switch (reload_mode
)
2005 sri
->icode
= CODE_FOR_reload_inv64si
;
2008 sri
->icode
= CODE_FOR_reload_inv64sf
;
2011 sri
->icode
= CODE_FOR_reload_inv64hi
;
2014 sri
->icode
= CODE_FOR_reload_inv64hf
;
2017 sri
->icode
= CODE_FOR_reload_inv64qi
;
2020 sri
->icode
= CODE_FOR_reload_inv64di
;
2023 sri
->icode
= CODE_FOR_reload_inv64df
;
2029 switch (reload_mode
)
2032 sri
->icode
= CODE_FOR_reload_outv64si
;
2035 sri
->icode
= CODE_FOR_reload_outv64sf
;
2038 sri
->icode
= CODE_FOR_reload_outv64hi
;
2041 sri
->icode
= CODE_FOR_reload_outv64hf
;
2044 sri
->icode
= CODE_FOR_reload_outv64qi
;
2047 sri
->icode
= CODE_FOR_reload_outv64di
;
2050 sri
->icode
= CODE_FOR_reload_outv64df
;
2058 case ADDR_SPACE_LDS
:
2059 case ADDR_SPACE_GDS
:
2060 case ADDR_SPACE_SCRATCH
:
2061 result
= (rclass
== VGPR_REGS
? NO_REGS
: VGPR_REGS
);
2066 if (dump_file
&& (dump_flags
& TDF_DETAILS
))
2067 fprintf (dump_file
, " <= %s (icode: %s)\n", reg_class_names
[result
],
2068 get_insn_name (sri
->icode
));
2073 /* Update register usage after having seen the compiler flags and kernel
2074 attributes. We typically want to fix registers that contain values
2075 set by the HSA runtime. */
2078 gcn_conditional_register_usage (void)
2080 if (!cfun
|| !cfun
->machine
)
2083 if (cfun
->machine
->normal_function
)
2085 /* Restrict the set of SGPRs and VGPRs used by non-kernel functions. */
2086 for (int i
= SGPR_REGNO (MAX_NORMAL_SGPR_COUNT
);
2087 i
<= LAST_SGPR_REG
; i
++)
2088 fixed_regs
[i
] = 1, call_used_regs
[i
] = 1;
2090 for (int i
= VGPR_REGNO (MAX_NORMAL_VGPR_COUNT
);
2091 i
<= LAST_VGPR_REG
; i
++)
2092 fixed_regs
[i
] = 1, call_used_regs
[i
] = 1;
2097 /* If the set of requested args is the default set, nothing more needs to
2099 if (cfun
->machine
->args
.requested
== default_requested_args
)
2102 /* Requesting a set of args different from the default violates the ABI. */
2103 if (!leaf_function_p ())
2104 warning (0, "A non-default set of initial values has been requested, "
2105 "which violates the ABI!");
2107 for (int i
= SGPR_REGNO (0); i
< SGPR_REGNO (14); i
++)
2110 /* Fix the runtime argument register containing values that may be
2111 needed later. DISPATCH_PTR_ARG and FLAT_SCRATCH_* should not be
2112 needed after the prologue so there's no need to fix them. */
2113 if (cfun
->machine
->args
.reg
[PRIVATE_SEGMENT_WAVE_OFFSET_ARG
] >= 0)
2114 fixed_regs
[cfun
->machine
->args
.reg
[PRIVATE_SEGMENT_WAVE_OFFSET_ARG
]] = 1;
2115 if (cfun
->machine
->args
.reg
[PRIVATE_SEGMENT_BUFFER_ARG
] >= 0)
2117 /* The upper 32-bits of the 64-bit descriptor are not used, so allow
2118 the containing registers to be used for other purposes. */
2119 fixed_regs
[cfun
->machine
->args
.reg
[PRIVATE_SEGMENT_BUFFER_ARG
]] = 1;
2120 fixed_regs
[cfun
->machine
->args
.reg
[PRIVATE_SEGMENT_BUFFER_ARG
] + 1] = 1;
2122 if (cfun
->machine
->args
.reg
[KERNARG_SEGMENT_PTR_ARG
] >= 0)
2124 fixed_regs
[cfun
->machine
->args
.reg
[KERNARG_SEGMENT_PTR_ARG
]] = 1;
2125 fixed_regs
[cfun
->machine
->args
.reg
[KERNARG_SEGMENT_PTR_ARG
] + 1] = 1;
2127 if (cfun
->machine
->args
.reg
[DISPATCH_PTR_ARG
] >= 0)
2129 fixed_regs
[cfun
->machine
->args
.reg
[DISPATCH_PTR_ARG
]] = 1;
2130 fixed_regs
[cfun
->machine
->args
.reg
[DISPATCH_PTR_ARG
] + 1] = 1;
2132 if (cfun
->machine
->args
.reg
[WORKGROUP_ID_X_ARG
] >= 0)
2133 fixed_regs
[cfun
->machine
->args
.reg
[WORKGROUP_ID_X_ARG
]] = 1;
2134 if (cfun
->machine
->args
.reg
[WORK_ITEM_ID_X_ARG
] >= 0)
2135 fixed_regs
[cfun
->machine
->args
.reg
[WORK_ITEM_ID_X_ARG
]] = 1;
2136 if (cfun
->machine
->args
.reg
[WORK_ITEM_ID_Y_ARG
] >= 0)
2137 fixed_regs
[cfun
->machine
->args
.reg
[WORK_ITEM_ID_Y_ARG
]] = 1;
2138 if (cfun
->machine
->args
.reg
[WORK_ITEM_ID_Z_ARG
] >= 0)
2139 fixed_regs
[cfun
->machine
->args
.reg
[WORK_ITEM_ID_Z_ARG
]] = 1;
2141 if (TARGET_GCN5_PLUS
)
2142 /* v0 is always zero, for global nul-offsets. */
2143 fixed_regs
[VGPR_REGNO (0)] = 1;
2146 /* Determine if a load or store is valid, according to the register classes
2147 and address space. Used primarily by the machine description to decide
2148 when to split a move into two steps. */
2151 gcn_valid_move_p (machine_mode mode
, rtx dest
, rtx src
)
2153 if (!MEM_P (dest
) && !MEM_P (src
))
2157 && AS_FLAT_P (MEM_ADDR_SPACE (dest
))
2158 && (gcn_flat_address_p (XEXP (dest
, 0), mode
)
2159 || GET_CODE (XEXP (dest
, 0)) == SYMBOL_REF
2160 || GET_CODE (XEXP (dest
, 0)) == LABEL_REF
)
2161 && gcn_vgpr_register_operand (src
, mode
))
2163 else if (MEM_P (src
)
2164 && AS_FLAT_P (MEM_ADDR_SPACE (src
))
2165 && (gcn_flat_address_p (XEXP (src
, 0), mode
)
2166 || GET_CODE (XEXP (src
, 0)) == SYMBOL_REF
2167 || GET_CODE (XEXP (src
, 0)) == LABEL_REF
)
2168 && gcn_vgpr_register_operand (dest
, mode
))
2172 && AS_GLOBAL_P (MEM_ADDR_SPACE (dest
))
2173 && (gcn_global_address_p (XEXP (dest
, 0))
2174 || GET_CODE (XEXP (dest
, 0)) == SYMBOL_REF
2175 || GET_CODE (XEXP (dest
, 0)) == LABEL_REF
)
2176 && gcn_vgpr_register_operand (src
, mode
))
2178 else if (MEM_P (src
)
2179 && AS_GLOBAL_P (MEM_ADDR_SPACE (src
))
2180 && (gcn_global_address_p (XEXP (src
, 0))
2181 || GET_CODE (XEXP (src
, 0)) == SYMBOL_REF
2182 || GET_CODE (XEXP (src
, 0)) == LABEL_REF
)
2183 && gcn_vgpr_register_operand (dest
, mode
))
2187 && MEM_ADDR_SPACE (dest
) == ADDR_SPACE_SCALAR_FLAT
2188 && (gcn_scalar_flat_address_p (XEXP (dest
, 0))
2189 || GET_CODE (XEXP (dest
, 0)) == SYMBOL_REF
2190 || GET_CODE (XEXP (dest
, 0)) == LABEL_REF
)
2191 && gcn_ssrc_register_operand (src
, mode
))
2193 else if (MEM_P (src
)
2194 && MEM_ADDR_SPACE (src
) == ADDR_SPACE_SCALAR_FLAT
2195 && (gcn_scalar_flat_address_p (XEXP (src
, 0))
2196 || GET_CODE (XEXP (src
, 0)) == SYMBOL_REF
2197 || GET_CODE (XEXP (src
, 0)) == LABEL_REF
)
2198 && gcn_sdst_register_operand (dest
, mode
))
2202 && AS_ANY_DS_P (MEM_ADDR_SPACE (dest
))
2203 && gcn_ds_address_p (XEXP (dest
, 0))
2204 && gcn_vgpr_register_operand (src
, mode
))
2206 else if (MEM_P (src
)
2207 && AS_ANY_DS_P (MEM_ADDR_SPACE (src
))
2208 && gcn_ds_address_p (XEXP (src
, 0))
2209 && gcn_vgpr_register_operand (dest
, mode
))
2216 /* {{{ Functions and ABI. */
2218 /* Implement TARGET_FUNCTION_VALUE.
2220 Define how to find the value returned by a function.
2221 The register location is always the same, but the mode depends on
2225 gcn_function_value (const_tree valtype
, const_tree
, bool)
2227 machine_mode mode
= TYPE_MODE (valtype
);
2229 if (INTEGRAL_TYPE_P (valtype
)
2230 && GET_MODE_CLASS (mode
) == MODE_INT
2231 && GET_MODE_SIZE (mode
) < 4)
2234 return gen_rtx_REG (mode
, SGPR_REGNO (RETURN_VALUE_REG
));
2237 /* Implement TARGET_FUNCTION_VALUE_REGNO_P.
2239 Return true if N is a possible register number for the function return
2243 gcn_function_value_regno_p (const unsigned int n
)
2245 return n
== RETURN_VALUE_REG
;
2248 /* Calculate the number of registers required to hold function argument
2252 num_arg_regs (const function_arg_info
&arg
)
2254 if (targetm
.calls
.must_pass_in_stack (arg
))
2257 int size
= arg
.promoted_size_in_bytes ();
2258 return (size
+ UNITS_PER_WORD
- 1) / UNITS_PER_WORD
;
2261 /* Implement TARGET_STRICT_ARGUMENT_NAMING.
2263 Return true if the location where a function argument is passed
2264 depends on whether or not it is a named argument
2266 For gcn, we know how to handle functions declared as stdarg: by
2267 passing an extra pointer to the unnamed arguments. However, the
2268 Fortran frontend can produce a different situation, where a
2269 function pointer is declared with no arguments, but the actual
2270 function and calls to it take more arguments. In that case, we
2271 want to ensure the call matches the definition of the function. */
2274 gcn_strict_argument_naming (cumulative_args_t cum_v
)
2276 CUMULATIVE_ARGS
*cum
= get_cumulative_args (cum_v
);
2278 return cum
->fntype
== NULL_TREE
|| stdarg_p (cum
->fntype
);
2281 /* Implement TARGET_PRETEND_OUTGOING_VARARGS_NAMED.
2283 See comment on gcn_strict_argument_naming. */
2286 gcn_pretend_outgoing_varargs_named (cumulative_args_t cum_v
)
2288 return !gcn_strict_argument_naming (cum_v
);
2291 /* Implement TARGET_FUNCTION_ARG.
2293 Return an RTX indicating whether a function argument is passed in a register
2294 and if so, which register. */
2297 gcn_function_arg (cumulative_args_t cum_v
, const function_arg_info
&arg
)
2299 CUMULATIVE_ARGS
*cum
= get_cumulative_args (cum_v
);
2300 if (cum
->normal_function
)
2302 if (!arg
.named
|| arg
.end_marker_p ())
2305 if (targetm
.calls
.must_pass_in_stack (arg
))
2308 /* Vector parameters are not supported yet. */
2309 if (VECTOR_MODE_P (arg
.mode
))
2312 int reg_num
= FIRST_PARM_REG
+ cum
->num
;
2313 int num_regs
= num_arg_regs (arg
);
2315 while (reg_num
% num_regs
!= 0)
2317 if (reg_num
+ num_regs
<= FIRST_PARM_REG
+ NUM_PARM_REGS
)
2318 return gen_rtx_REG (arg
.mode
, reg_num
);
2322 if (cum
->num
>= cum
->args
.nargs
)
2324 cum
->offset
= (cum
->offset
+ TYPE_ALIGN (arg
.type
) / 8 - 1)
2325 & -(TYPE_ALIGN (arg
.type
) / 8);
2326 cfun
->machine
->kernarg_segment_alignment
2327 = MAX ((unsigned) cfun
->machine
->kernarg_segment_alignment
,
2328 TYPE_ALIGN (arg
.type
) / 8);
2329 rtx addr
= gen_rtx_REG (DImode
,
2330 cum
->args
.reg
[KERNARG_SEGMENT_PTR_ARG
]);
2332 addr
= gen_rtx_PLUS (DImode
, addr
,
2333 gen_int_mode (cum
->offset
, DImode
));
2334 rtx mem
= gen_rtx_MEM (arg
.mode
, addr
);
2335 set_mem_attributes (mem
, arg
.type
, 1);
2336 set_mem_addr_space (mem
, ADDR_SPACE_SCALAR_FLAT
);
2337 MEM_READONLY_P (mem
) = 1;
2341 int a
= cum
->args
.order
[cum
->num
];
2342 if (arg
.mode
!= gcn_kernel_arg_types
[a
].mode
)
2344 error ("wrong type of argument %s", gcn_kernel_arg_types
[a
].name
);
2347 return gen_rtx_REG ((machine_mode
) gcn_kernel_arg_types
[a
].mode
,
2353 /* Implement TARGET_FUNCTION_ARG_ADVANCE.
2355 Updates the summarizer variable pointed to by CUM_V to advance past an
2356 argument in the argument list. */
2359 gcn_function_arg_advance (cumulative_args_t cum_v
,
2360 const function_arg_info
&arg
)
2362 CUMULATIVE_ARGS
*cum
= get_cumulative_args (cum_v
);
2364 if (cum
->normal_function
)
2369 int num_regs
= num_arg_regs (arg
);
2371 while ((FIRST_PARM_REG
+ cum
->num
) % num_regs
!= 0)
2373 cum
->num
+= num_regs
;
2377 if (cum
->num
< cum
->args
.nargs
)
2381 cum
->offset
+= tree_to_uhwi (TYPE_SIZE_UNIT (arg
.type
));
2382 cfun
->machine
->kernarg_segment_byte_size
= cum
->offset
;
2387 /* Implement TARGET_ARG_PARTIAL_BYTES.
2389 Returns the number of bytes at the beginning of an argument that must be put
2390 in registers. The value must be zero for arguments that are passed entirely
2391 in registers or that are entirely pushed on the stack. */
2394 gcn_arg_partial_bytes (cumulative_args_t cum_v
, const function_arg_info
&arg
)
2396 CUMULATIVE_ARGS
*cum
= get_cumulative_args (cum_v
);
2401 if (targetm
.calls
.must_pass_in_stack (arg
))
2404 if (cum
->num
>= NUM_PARM_REGS
)
2407 /* If the argument fits entirely in registers, return 0. */
2408 if (cum
->num
+ num_arg_regs (arg
) <= NUM_PARM_REGS
)
2411 return (NUM_PARM_REGS
- cum
->num
) * UNITS_PER_WORD
;
2414 /* A normal function which takes a pointer argument (to a scalar) may be
2415 passed a pointer to LDS space (via a high-bits-set aperture), and that only
2416 works with FLAT addressing, not GLOBAL. Force FLAT addressing if the
2417 function has an incoming pointer-to-scalar parameter. */
2420 gcn_detect_incoming_pointer_arg (tree fndecl
)
2422 gcc_assert (cfun
&& cfun
->machine
);
2424 for (tree arg
= TYPE_ARG_TYPES (TREE_TYPE (fndecl
));
2426 arg
= TREE_CHAIN (arg
))
2427 if (POINTER_TYPE_P (TREE_VALUE (arg
))
2428 && !AGGREGATE_TYPE_P (TREE_TYPE (TREE_VALUE (arg
))))
2429 cfun
->machine
->use_flat_addressing
= true;
2432 /* Implement INIT_CUMULATIVE_ARGS, via gcn.h.
2434 Initialize a variable CUM of type CUMULATIVE_ARGS for a call to a function
2435 whose data type is FNTYPE. For a library call, FNTYPE is 0. */
2438 gcn_init_cumulative_args (CUMULATIVE_ARGS
*cum
/* Argument info to init */ ,
2439 tree fntype
/* tree ptr for function decl */ ,
2440 rtx libname
/* SYMBOL_REF of library name or 0 */ ,
2441 tree fndecl
, int caller
)
2443 memset (cum
, 0, sizeof (*cum
));
2444 cum
->fntype
= fntype
;
2447 gcc_assert (cfun
&& cfun
->machine
);
2448 cum
->normal_function
= true;
2451 cfun
->machine
->normal_function
= true;
2452 gcn_detect_incoming_pointer_arg (fndecl
);
2458 attr
= lookup_attribute ("amdgpu_hsa_kernel", DECL_ATTRIBUTES (fndecl
));
2459 if (fndecl
&& !attr
)
2460 attr
= lookup_attribute ("amdgpu_hsa_kernel",
2461 TYPE_ATTRIBUTES (TREE_TYPE (fndecl
)));
2462 if (!attr
&& fntype
)
2463 attr
= lookup_attribute ("amdgpu_hsa_kernel", TYPE_ATTRIBUTES (fntype
));
2464 /* Handle main () as kernel, so we can run testsuite.
2465 Handle OpenACC kernels similarly to main. */
2466 if (!attr
&& !caller
&& fndecl
2467 && (MAIN_NAME_P (DECL_NAME (fndecl
))
2468 || lookup_attribute ("omp target entrypoint",
2469 DECL_ATTRIBUTES (fndecl
)) != NULL_TREE
))
2470 gcn_parse_amdgpu_hsa_kernel_attribute (&cum
->args
, NULL_TREE
);
2473 if (!attr
|| caller
)
2475 gcc_assert (cfun
&& cfun
->machine
);
2476 cum
->normal_function
= true;
2478 cfun
->machine
->normal_function
= true;
2480 gcn_parse_amdgpu_hsa_kernel_attribute
2481 (&cum
->args
, attr
? TREE_VALUE (attr
) : NULL_TREE
);
2483 cfun
->machine
->args
= cum
->args
;
2484 if (!caller
&& cfun
->machine
->normal_function
)
2485 gcn_detect_incoming_pointer_arg (fndecl
);
2491 gcn_return_in_memory (const_tree type
, const_tree
ARG_UNUSED (fntype
))
2493 machine_mode mode
= TYPE_MODE (type
);
2494 HOST_WIDE_INT size
= int_size_in_bytes (type
);
2496 if (AGGREGATE_TYPE_P (type
))
2499 /* Vector return values are not supported yet. */
2500 if (VECTOR_TYPE_P (type
))
2503 if (mode
== BLKmode
)
2506 if (size
> 2 * UNITS_PER_WORD
)
2512 /* Implement TARGET_PROMOTE_FUNCTION_MODE.
2514 Return the mode to use for outgoing function arguments. */
2517 gcn_promote_function_mode (const_tree
ARG_UNUSED (type
), machine_mode mode
,
2518 int *ARG_UNUSED (punsignedp
),
2519 const_tree
ARG_UNUSED (funtype
),
2520 int ARG_UNUSED (for_return
))
2522 if (GET_MODE_CLASS (mode
) == MODE_INT
&& GET_MODE_SIZE (mode
) < 4)
2528 /* Implement TARGET_GIMPLIFY_VA_ARG_EXPR.
2530 Derived from hppa_gimplify_va_arg_expr. The generic routine doesn't handle
2531 ARGS_GROW_DOWNWARDS. */
2534 gcn_gimplify_va_arg_expr (tree valist
, tree type
,
2535 gimple_seq
*ARG_UNUSED (pre_p
),
2536 gimple_seq
*ARG_UNUSED (post_p
))
2538 tree ptr
= build_pointer_type (type
);
2543 indirect
= pass_va_arg_by_reference (type
);
2547 ptr
= build_pointer_type (type
);
2549 valist_type
= TREE_TYPE (valist
);
2551 /* Args grow down. Not handled by generic routines. */
2553 u
= fold_convert (sizetype
, size_in_bytes (type
));
2554 u
= fold_build1 (NEGATE_EXPR
, sizetype
, u
);
2555 t
= fold_build_pointer_plus (valist
, u
);
2557 /* Align to 8 byte boundary. */
2559 u
= build_int_cst (TREE_TYPE (t
), -8);
2560 t
= build2 (BIT_AND_EXPR
, TREE_TYPE (t
), t
, u
);
2561 t
= fold_convert (valist_type
, t
);
2563 t
= build2 (MODIFY_EXPR
, valist_type
, valist
, t
);
2565 t
= fold_convert (ptr
, t
);
2566 t
= build_va_arg_indirect_ref (t
);
2569 t
= build_va_arg_indirect_ref (t
);
2574 /* Return 1 if TRAIT NAME is present in the OpenMP context's
2575 device trait set, return 0 if not present in any OpenMP context in the
2576 whole translation unit, or -1 if not present in the current OpenMP context
2577 but might be present in another OpenMP context in the same TU. */
2580 gcn_omp_device_kind_arch_isa (enum omp_device_kind_arch_isa trait
,
2585 case omp_device_kind
:
2586 return strcmp (name
, "gpu") == 0;
2587 case omp_device_arch
:
2588 return strcmp (name
, "gcn") == 0;
2589 case omp_device_isa
:
2590 if (strcmp (name
, "fiji") == 0)
2591 return gcn_arch
== PROCESSOR_FIJI
;
2592 if (strcmp (name
, "gfx900") == 0)
2593 return gcn_arch
== PROCESSOR_VEGA10
;
2594 if (strcmp (name
, "gfx906") == 0)
2595 return gcn_arch
== PROCESSOR_VEGA20
;
2602 /* Calculate stack offsets needed to create prologues and epilogues. */
2604 static struct machine_function
*
2605 gcn_compute_frame_offsets (void)
2607 machine_function
*offsets
= cfun
->machine
;
2609 if (reload_completed
)
2612 offsets
->need_frame_pointer
= frame_pointer_needed
;
2614 offsets
->outgoing_args_size
= crtl
->outgoing_args_size
;
2615 offsets
->pretend_size
= crtl
->args
.pretend_args_size
;
2617 offsets
->local_vars
= get_frame_size ();
2619 offsets
->lr_needs_saving
= (!leaf_function_p ()
2620 || df_regs_ever_live_p (LR_REGNUM
)
2621 || df_regs_ever_live_p (LR_REGNUM
+ 1));
2623 offsets
->callee_saves
= offsets
->lr_needs_saving
? 8 : 0;
2625 for (int regno
= 0; regno
< FIRST_PSEUDO_REGISTER
; regno
++)
2626 if ((df_regs_ever_live_p (regno
) && !call_used_or_fixed_reg_p (regno
))
2627 || ((regno
& ~1) == HARD_FRAME_POINTER_REGNUM
2628 && frame_pointer_needed
))
2629 offsets
->callee_saves
+= (VGPR_REGNO_P (regno
) ? 256 : 4);
2631 /* Round up to 64-bit boundary to maintain stack alignment. */
2632 offsets
->callee_saves
= (offsets
->callee_saves
+ 7) & ~7;
2637 /* Insert code into the prologue or epilogue to store or load any
2638 callee-save register to/from the stack.
2640 Helper function for gcn_expand_prologue and gcn_expand_epilogue. */
2643 move_callee_saved_registers (rtx sp
, machine_function
*offsets
,
2646 int regno
, offset
, saved_scalars
;
2647 rtx exec
= gen_rtx_REG (DImode
, EXEC_REG
);
2648 rtx vcc
= gen_rtx_REG (DImode
, VCC_LO_REG
);
2649 rtx offreg
= gen_rtx_REG (SImode
, SGPR_REGNO (22));
2650 rtx as
= gen_rtx_CONST_INT (VOIDmode
, STACK_ADDR_SPACE
);
2651 HOST_WIDE_INT exec_set
= 0;
2656 /* Move scalars into two vector registers. */
2657 for (regno
= 0, saved_scalars
= 0; regno
< FIRST_VGPR_REG
; regno
++)
2658 if ((df_regs_ever_live_p (regno
) && !call_used_or_fixed_reg_p (regno
))
2659 || ((regno
& ~1) == LINK_REGNUM
&& offsets
->lr_needs_saving
)
2660 || ((regno
& ~1) == HARD_FRAME_POINTER_REGNUM
2661 && offsets
->need_frame_pointer
))
2663 rtx reg
= gen_rtx_REG (SImode
, regno
);
2664 rtx vreg
= gen_rtx_REG (V64SImode
,
2665 VGPR_REGNO (6 + (saved_scalars
/ 64)));
2666 int lane
= saved_scalars
% 64;
2669 emit_insn (gen_vec_setv64si (vreg
, reg
, GEN_INT (lane
)));
2671 emit_insn (gen_vec_extractv64sisi (reg
, vreg
, GEN_INT (lane
)));
2676 rtx move_scalars
= get_insns ();
2680 /* Ensure that all vector lanes are moved. */
2682 emit_move_insn (exec
, GEN_INT (exec_set
));
2684 /* Set up a vector stack pointer. */
2685 rtx _0_1_2_3
= gen_rtx_REG (V64SImode
, VGPR_REGNO (1));
2686 rtx _0_4_8_12
= gen_rtx_REG (V64SImode
, VGPR_REGNO (3));
2687 emit_insn (gen_ashlv64si3_exec (_0_4_8_12
, _0_1_2_3
, GEN_INT (2),
2688 gcn_gen_undef (V64SImode
), exec
));
2689 rtx vsp
= gen_rtx_REG (V64DImode
, VGPR_REGNO (4));
2690 emit_insn (gen_vec_duplicatev64di_exec (vsp
, sp
, gcn_gen_undef (V64DImode
),
2692 emit_insn (gen_addv64si3_vcc_exec (gcn_operand_part (V64SImode
, vsp
, 0),
2693 gcn_operand_part (V64SImode
, vsp
, 0),
2694 _0_4_8_12
, vcc
, gcn_gen_undef (V64SImode
),
2696 emit_insn (gen_addcv64si3_exec (gcn_operand_part (V64SImode
, vsp
, 1),
2697 gcn_operand_part (V64SImode
, vsp
, 1),
2698 const0_rtx
, vcc
, vcc
,
2699 gcn_gen_undef (V64SImode
), exec
));
2702 for (regno
= FIRST_VGPR_REG
, offset
= offsets
->pretend_size
;
2703 regno
< FIRST_PSEUDO_REGISTER
; regno
++)
2704 if ((df_regs_ever_live_p (regno
) && !call_used_or_fixed_reg_p (regno
))
2705 || (regno
== VGPR_REGNO (6) && saved_scalars
> 0)
2706 || (regno
== VGPR_REGNO (7) && saved_scalars
> 63))
2708 rtx reg
= gen_rtx_REG (V64SImode
, regno
);
2711 if (regno
== VGPR_REGNO (6) && saved_scalars
< 64)
2712 size
= saved_scalars
* 4;
2713 else if (regno
== VGPR_REGNO (7) && saved_scalars
< 128)
2714 size
= (saved_scalars
- 64) * 4;
2716 if (size
!= 256 || exec_set
!= -1)
2718 exec_set
= ((unsigned HOST_WIDE_INT
) 1 << (size
/ 4)) - 1;
2719 emit_move_insn (exec
, gen_int_mode (exec_set
, DImode
));
2723 emit_insn (gen_scatterv64si_insn_1offset_exec (vsp
, const0_rtx
, reg
,
2724 as
, const0_rtx
, exec
));
2726 emit_insn (gen_gatherv64si_insn_1offset_exec
2727 (reg
, vsp
, const0_rtx
, as
, const0_rtx
,
2728 gcn_gen_undef (V64SImode
), exec
));
2730 /* Move our VSP to the next stack entry. */
2731 if (offreg_set
!= size
)
2734 emit_move_insn (offreg
, GEN_INT (size
));
2739 emit_move_insn (exec
, GEN_INT (exec_set
));
2741 emit_insn (gen_addv64si3_vcc_dup_exec
2742 (gcn_operand_part (V64SImode
, vsp
, 0),
2743 offreg
, gcn_operand_part (V64SImode
, vsp
, 0),
2744 vcc
, gcn_gen_undef (V64SImode
), exec
));
2745 emit_insn (gen_addcv64si3_exec
2746 (gcn_operand_part (V64SImode
, vsp
, 1),
2747 gcn_operand_part (V64SImode
, vsp
, 1),
2748 const0_rtx
, vcc
, vcc
, gcn_gen_undef (V64SImode
), exec
));
2753 rtx move_vectors
= get_insns ();
2758 emit_insn (move_scalars
);
2759 emit_insn (move_vectors
);
2763 emit_insn (move_vectors
);
2764 emit_insn (move_scalars
);
2768 /* Generate prologue. Called from gen_prologue during pro_and_epilogue pass.
2770 For a non-kernel function, the stack layout looks like this (interim),
2774 |__________________| <-- current SP
2776 |__________________|
2778 |__________________|
2780 |__________________| <-- FP/hard FP
2781 | callee-save regs |
2782 |__________________| <-- soft arg pointer
2784 |__________________| <-- incoming SP
2786 lo |..................|
2788 This implies arguments (beyond the first N in registers) must grow
2789 downwards (as, apparently, PA has them do).
2791 For a kernel function we have the simpler:
2794 |__________________| <-- current SP
2796 |__________________|
2798 |__________________|
2800 lo |__________________| <-- FP/hard FP
2805 gcn_expand_prologue ()
2807 machine_function
*offsets
= gcn_compute_frame_offsets ();
2809 if (!cfun
|| !cfun
->machine
|| cfun
->machine
->normal_function
)
2811 rtx sp
= gen_rtx_REG (Pmode
, STACK_POINTER_REGNUM
);
2812 rtx sp_hi
= gcn_operand_part (Pmode
, sp
, 1);
2813 rtx sp_lo
= gcn_operand_part (Pmode
, sp
, 0);
2814 rtx fp
= gen_rtx_REG (Pmode
, HARD_FRAME_POINTER_REGNUM
);
2815 rtx fp_hi
= gcn_operand_part (Pmode
, fp
, 1);
2816 rtx fp_lo
= gcn_operand_part (Pmode
, fp
, 0);
2820 if (offsets
->pretend_size
> 0)
2822 /* FIXME: Do the actual saving of register pretend args to the stack.
2823 Register order needs consideration. */
2826 /* Save callee-save regs. */
2827 move_callee_saved_registers (sp
, offsets
, true);
2829 HOST_WIDE_INT sp_adjust
= offsets
->pretend_size
2830 + offsets
->callee_saves
2831 + offsets
->local_vars
+ offsets
->outgoing_args_size
;
2834 /* Adding RTX_FRAME_RELATED_P effectively disables spliting, so
2835 we use split add explictly, and specify the DImode add in
2837 rtx scc
= gen_rtx_REG (BImode
, SCC_REG
);
2838 rtx adjustment
= gen_int_mode (sp_adjust
, SImode
);
2839 rtx insn
= emit_insn (gen_addsi3_scalar_carry (sp_lo
, sp_lo
,
2841 RTX_FRAME_RELATED_P (insn
) = 1;
2842 add_reg_note (insn
, REG_FRAME_RELATED_EXPR
,
2844 gen_rtx_PLUS (DImode
, sp
, adjustment
)));
2845 emit_insn (gen_addcsi3_scalar_zero (sp_hi
, sp_hi
, scc
));
2848 if (offsets
->need_frame_pointer
)
2850 /* Adding RTX_FRAME_RELATED_P effectively disables spliting, so
2851 we use split add explictly, and specify the DImode add in
2853 rtx scc
= gen_rtx_REG (BImode
, SCC_REG
);
2854 int fp_adjust
= -(offsets
->local_vars
+ offsets
->outgoing_args_size
);
2855 rtx adjustment
= gen_int_mode (fp_adjust
, SImode
);
2856 rtx insn
= emit_insn (gen_addsi3_scalar_carry(fp_lo
, sp_lo
,
2858 RTX_FRAME_RELATED_P (insn
) = 1;
2859 add_reg_note (insn
, REG_FRAME_RELATED_EXPR
,
2861 gen_rtx_PLUS (DImode
, sp
, adjustment
)));
2862 emit_insn (gen_addcsi3_scalar (fp_hi
, sp_hi
,
2863 (fp_adjust
< 0 ? GEN_INT (-1)
2868 rtx_insn
*seq
= get_insns ();
2871 /* FIXME: Prologue insns should have this flag set for debug output, etc.
2872 but it causes issues for now.
2873 for (insn = seq; insn; insn = NEXT_INSN (insn))
2875 RTX_FRAME_RELATED_P (insn) = 1;*/
2881 rtx wave_offset
= gen_rtx_REG (SImode
,
2882 cfun
->machine
->args
.
2883 reg
[PRIVATE_SEGMENT_WAVE_OFFSET_ARG
]);
2885 if (cfun
->machine
->args
.requested
& (1 << FLAT_SCRATCH_INIT_ARG
))
2888 gen_rtx_REG (SImode
,
2889 cfun
->machine
->args
.reg
[FLAT_SCRATCH_INIT_ARG
]);
2891 gen_rtx_REG (SImode
,
2892 cfun
->machine
->args
.reg
[FLAT_SCRATCH_INIT_ARG
] + 1);
2893 rtx fs_reg_lo
= gen_rtx_REG (SImode
, FLAT_SCRATCH_REG
);
2894 rtx fs_reg_hi
= gen_rtx_REG (SImode
, FLAT_SCRATCH_REG
+ 1);
2896 /*rtx queue = gen_rtx_REG(DImode,
2897 cfun->machine->args.reg[QUEUE_PTR_ARG]);
2898 rtx aperture = gen_rtx_MEM (SImode,
2899 gen_rtx_PLUS (DImode, queue,
2900 gen_int_mode (68, SImode)));
2901 set_mem_addr_space (aperture, ADDR_SPACE_SCALAR_FLAT);*/
2903 /* Set up flat_scratch. */
2904 emit_insn (gen_addsi3_scc (fs_reg_hi
, fs_init_lo
, wave_offset
));
2905 emit_insn (gen_lshrsi3_scc (fs_reg_hi
, fs_reg_hi
,
2906 gen_int_mode (8, SImode
)));
2907 emit_move_insn (fs_reg_lo
, fs_init_hi
);
2910 /* Set up frame pointer and stack pointer. */
2911 rtx sp
= gen_rtx_REG (DImode
, STACK_POINTER_REGNUM
);
2912 rtx sp_hi
= simplify_gen_subreg (SImode
, sp
, DImode
, 4);
2913 rtx sp_lo
= simplify_gen_subreg (SImode
, sp
, DImode
, 0);
2914 rtx fp
= gen_rtx_REG (DImode
, HARD_FRAME_POINTER_REGNUM
);
2915 rtx fp_hi
= simplify_gen_subreg (SImode
, fp
, DImode
, 4);
2916 rtx fp_lo
= simplify_gen_subreg (SImode
, fp
, DImode
, 0);
2918 HOST_WIDE_INT sp_adjust
= (offsets
->local_vars
2919 + offsets
->outgoing_args_size
);
2921 /* Initialise FP and SP from the buffer descriptor in s[0:3]. */
2922 emit_move_insn (fp_lo
, gen_rtx_REG (SImode
, 0));
2923 emit_insn (gen_andsi3_scc (fp_hi
, gen_rtx_REG (SImode
, 1),
2924 gen_int_mode (0xffff, SImode
)));
2925 rtx scc
= gen_rtx_REG (BImode
, SCC_REG
);
2926 emit_insn (gen_addsi3_scalar_carry (fp_lo
, fp_lo
, wave_offset
, scc
));
2927 emit_insn (gen_addcsi3_scalar_zero (fp_hi
, fp_hi
, scc
));
2929 /* Adding RTX_FRAME_RELATED_P effectively disables spliting, so we use
2930 split add explictly, and specify the DImode add in the note.
2931 The DWARF info expects that the callee-save data is in the frame,
2932 even though it isn't (because this is the entry point), so we
2933 make a notional adjustment to the DWARF frame offset here. */
2934 rtx dbg_adjustment
= gen_int_mode (sp_adjust
+ offsets
->callee_saves
,
2939 rtx scc
= gen_rtx_REG (BImode
, SCC_REG
);
2940 rtx adjustment
= gen_int_mode (sp_adjust
, DImode
);
2941 insn
= emit_insn (gen_addsi3_scalar_carry(sp_lo
, fp_lo
, adjustment
,
2943 emit_insn (gen_addcsi3_scalar_zero (sp_hi
, fp_hi
, scc
));
2946 insn
= emit_move_insn (sp
, fp
);
2947 RTX_FRAME_RELATED_P (insn
) = 1;
2948 add_reg_note (insn
, REG_FRAME_RELATED_EXPR
,
2949 gen_rtx_SET (sp
, gen_rtx_PLUS (DImode
, sp
,
2952 /* Make sure the flat scratch reg doesn't get optimised away. */
2953 emit_insn (gen_prologue_use (gen_rtx_REG (DImode
, FLAT_SCRATCH_REG
)));
2956 /* Ensure that the scheduler doesn't do anything unexpected. */
2957 emit_insn (gen_blockage ());
2959 /* m0 is initialized for the usual LDS DS and FLAT memory case.
2960 The low-part is the address of the topmost addressable byte, which is
2961 size-1. The high-part is an offset and should be zero. */
2962 emit_move_insn (gen_rtx_REG (SImode
, M0_REG
),
2963 gen_int_mode (LDS_SIZE
-1, SImode
));
2965 emit_insn (gen_prologue_use (gen_rtx_REG (SImode
, M0_REG
)));
2967 if (cfun
&& cfun
->machine
&& !cfun
->machine
->normal_function
&& flag_openmp
)
2969 /* OpenMP kernels have an implicit call to gomp_gcn_enter_kernel. */
2970 rtx fn_reg
= gen_rtx_REG (Pmode
, FIRST_PARM_REG
);
2971 emit_move_insn (fn_reg
, gen_rtx_SYMBOL_REF (Pmode
,
2972 "gomp_gcn_enter_kernel"));
2973 emit_call_insn (gen_gcn_indirect_call (fn_reg
, const0_rtx
));
2977 /* Generate epilogue. Called from gen_epilogue during pro_and_epilogue pass.
2979 See gcn_expand_prologue for stack details. */
2982 gcn_expand_epilogue (void)
2984 /* Ensure that the scheduler doesn't do anything unexpected. */
2985 emit_insn (gen_blockage ());
2987 if (!cfun
|| !cfun
->machine
|| cfun
->machine
->normal_function
)
2989 machine_function
*offsets
= gcn_compute_frame_offsets ();
2990 rtx sp
= gen_rtx_REG (Pmode
, STACK_POINTER_REGNUM
);
2991 rtx fp
= gen_rtx_REG (Pmode
, HARD_FRAME_POINTER_REGNUM
);
2993 HOST_WIDE_INT sp_adjust
= offsets
->callee_saves
+ offsets
->pretend_size
;
2995 if (offsets
->need_frame_pointer
)
2997 /* Restore old SP from the frame pointer. */
2999 emit_insn (gen_subdi3 (sp
, fp
, gen_int_mode (sp_adjust
, DImode
)));
3001 emit_move_insn (sp
, fp
);
3005 /* Restore old SP from current SP. */
3006 sp_adjust
+= offsets
->outgoing_args_size
+ offsets
->local_vars
;
3009 emit_insn (gen_subdi3 (sp
, sp
, gen_int_mode (sp_adjust
, DImode
)));
3012 move_callee_saved_registers (sp
, offsets
, false);
3014 /* There's no explicit use of the link register on the return insn. Emit
3015 one here instead. */
3016 if (offsets
->lr_needs_saving
)
3017 emit_use (gen_rtx_REG (DImode
, LINK_REGNUM
));
3019 /* Similar for frame pointer. */
3020 if (offsets
->need_frame_pointer
)
3021 emit_use (gen_rtx_REG (DImode
, HARD_FRAME_POINTER_REGNUM
));
3023 else if (flag_openmp
)
3025 /* OpenMP kernels have an implicit call to gomp_gcn_exit_kernel. */
3026 rtx fn_reg
= gen_rtx_REG (Pmode
, FIRST_PARM_REG
);
3027 emit_move_insn (fn_reg
,
3028 gen_rtx_SYMBOL_REF (Pmode
, "gomp_gcn_exit_kernel"));
3029 emit_call_insn (gen_gcn_indirect_call (fn_reg
, const0_rtx
));
3031 else if (TREE_CODE (TREE_TYPE (DECL_RESULT (cfun
->decl
))) != VOID_TYPE
)
3033 /* Assume that an exit value compatible with gcn-run is expected.
3034 That is, the third input parameter is an int*.
3036 We can't allocate any new registers, but the kernarg_reg is
3037 dead after this, so we'll use that. */
3038 rtx kernarg_reg
= gen_rtx_REG (DImode
, cfun
->machine
->args
.reg
3039 [KERNARG_SEGMENT_PTR_ARG
]);
3040 rtx retptr_mem
= gen_rtx_MEM (DImode
,
3041 gen_rtx_PLUS (DImode
, kernarg_reg
,
3043 set_mem_addr_space (retptr_mem
, ADDR_SPACE_SCALAR_FLAT
);
3044 emit_move_insn (kernarg_reg
, retptr_mem
);
3046 rtx retval_mem
= gen_rtx_MEM (SImode
, kernarg_reg
);
3047 set_mem_addr_space (retval_mem
, ADDR_SPACE_SCALAR_FLAT
);
3048 emit_move_insn (retval_mem
,
3049 gen_rtx_REG (SImode
, SGPR_REGNO (RETURN_VALUE_REG
)));
3052 emit_jump_insn (gen_gcn_return ());
3055 /* Implement TARGET_CAN_ELIMINATE.
3057 Return true if the compiler is allowed to try to replace register number
3058 FROM_REG with register number TO_REG.
3060 FIXME: is the default "true" not enough? Should this be a negative set? */
3063 gcn_can_eliminate_p (int /*from_reg */ , int to_reg
)
3065 return (to_reg
== HARD_FRAME_POINTER_REGNUM
3066 || to_reg
== STACK_POINTER_REGNUM
);
3069 /* Implement INITIAL_ELIMINATION_OFFSET.
3071 Returns the initial difference between the specified pair of registers, in
3072 terms of stack position. */
3075 gcn_initial_elimination_offset (int from
, int to
)
3077 machine_function
*offsets
= gcn_compute_frame_offsets ();
3081 case ARG_POINTER_REGNUM
:
3082 if (to
== STACK_POINTER_REGNUM
)
3083 return -(offsets
->callee_saves
+ offsets
->local_vars
3084 + offsets
->outgoing_args_size
);
3085 else if (to
== FRAME_POINTER_REGNUM
|| to
== HARD_FRAME_POINTER_REGNUM
)
3086 return -offsets
->callee_saves
;
3091 case FRAME_POINTER_REGNUM
:
3092 if (to
== STACK_POINTER_REGNUM
)
3093 return -(offsets
->local_vars
+ offsets
->outgoing_args_size
);
3094 else if (to
== HARD_FRAME_POINTER_REGNUM
)
3105 /* Implement HARD_REGNO_RENAME_OK.
3107 Return true if it is permissible to rename a hard register from
3108 FROM_REG to TO_REG. */
3111 gcn_hard_regno_rename_ok (unsigned int from_reg
, unsigned int to_reg
)
3113 if (from_reg
== SCC_REG
3114 || from_reg
== VCC_LO_REG
|| from_reg
== VCC_HI_REG
3115 || from_reg
== EXEC_LO_REG
|| from_reg
== EXEC_HI_REG
3116 || to_reg
== SCC_REG
3117 || to_reg
== VCC_LO_REG
|| to_reg
== VCC_HI_REG
3118 || to_reg
== EXEC_LO_REG
|| to_reg
== EXEC_HI_REG
)
3121 /* Allow the link register to be used if it was saved. */
3122 if ((to_reg
& ~1) == LINK_REGNUM
)
3123 return !cfun
|| cfun
->machine
->lr_needs_saving
;
3125 /* Allow the registers used for the static chain to be used if the chain is
3126 not in active use. */
3127 if ((to_reg
& ~1) == STATIC_CHAIN_REGNUM
)
3129 || !(cfun
->static_chain_decl
3130 && df_regs_ever_live_p (STATIC_CHAIN_REGNUM
)
3131 && df_regs_ever_live_p (STATIC_CHAIN_REGNUM
+ 1));
3136 /* Implement HARD_REGNO_CALLER_SAVE_MODE.
3138 Which mode is required for saving NREGS of a pseudo-register in
3139 call-clobbered hard register REGNO. */
3142 gcn_hard_regno_caller_save_mode (unsigned int regno
, unsigned int nregs
,
3143 machine_mode regmode
)
3145 machine_mode result
= choose_hard_reg_mode (regno
, nregs
, NULL
);
3147 if (VECTOR_MODE_P (result
) && !VECTOR_MODE_P (regmode
))
3148 result
= (nregs
== 1 ? SImode
: DImode
);
3153 /* Implement TARGET_ASM_TRAMPOLINE_TEMPLATE.
3155 Output assembler code for a block containing the constant parts
3156 of a trampoline, leaving space for the variable parts. */
3159 gcn_asm_trampoline_template (FILE *f
)
3161 /* The source operand of the move instructions must be a 32-bit
3162 constant following the opcode. */
3163 asm_fprintf (f
, "\ts_mov_b32\ts%i, 0xffff\n", STATIC_CHAIN_REGNUM
);
3164 asm_fprintf (f
, "\ts_mov_b32\ts%i, 0xffff\n", STATIC_CHAIN_REGNUM
+ 1);
3165 asm_fprintf (f
, "\ts_mov_b32\ts%i, 0xffff\n", CC_SAVE_REG
);
3166 asm_fprintf (f
, "\ts_mov_b32\ts%i, 0xffff\n", CC_SAVE_REG
+ 1);
3167 asm_fprintf (f
, "\ts_setpc_b64\ts[%i:%i]\n", CC_SAVE_REG
, CC_SAVE_REG
+ 1);
3170 /* Implement TARGET_TRAMPOLINE_INIT.
3172 Emit RTL insns to initialize the variable parts of a trampoline.
3173 FNDECL is the decl of the target address, M_TRAMP is a MEM for
3174 the trampoline, and CHAIN_VALUE is an RTX for the static chain
3175 to be passed to the target function. */
3178 gcn_trampoline_init (rtx m_tramp
, tree fndecl
, rtx chain_value
)
3180 if (TARGET_GCN5_PLUS
)
3181 sorry ("nested function trampolines not supported on GCN5 due to"
3182 " non-executable stacks");
3184 emit_block_move (m_tramp
, assemble_trampoline_template (),
3185 GEN_INT (TRAMPOLINE_SIZE
), BLOCK_OP_NORMAL
);
3187 rtx fnaddr
= XEXP (DECL_RTL (fndecl
), 0);
3188 rtx chain_value_reg
= copy_to_reg (chain_value
);
3189 rtx fnaddr_reg
= copy_to_reg (fnaddr
);
3191 for (int i
= 0; i
< 4; i
++)
3193 rtx mem
= adjust_address (m_tramp
, SImode
, i
* 8 + 4);
3194 rtx reg
= i
< 2 ? chain_value_reg
: fnaddr_reg
;
3195 emit_move_insn (mem
, gen_rtx_SUBREG (SImode
, reg
, (i
% 2) * 4));
3198 rtx tramp_addr
= XEXP (m_tramp
, 0);
3199 emit_insn (gen_clear_icache (tramp_addr
,
3200 plus_constant (ptr_mode
, tramp_addr
,
3205 /* {{{ Miscellaneous. */
3207 /* Implement TARGET_CANNOT_COPY_INSN_P.
3209 Return true if INSN must not be duplicated. */
3212 gcn_cannot_copy_insn_p (rtx_insn
*insn
)
3214 if (recog_memoized (insn
) == CODE_FOR_gcn_wavefront_barrier
)
3220 /* Implement TARGET_DEBUG_UNWIND_INFO.
3222 Defines the mechanism that will be used for describing frame unwind
3223 information to the debugger. */
3225 static enum unwind_info_type
3226 gcn_debug_unwind_info ()
3228 /* No support for debug info, yet. */
3232 /* Determine if there is a suitable hardware conversion instruction.
3233 Used primarily by the machine description. */
3236 gcn_valid_cvt_p (machine_mode from
, machine_mode to
, enum gcn_cvt_t op
)
3238 if (VECTOR_MODE_P (from
) != VECTOR_MODE_P (to
))
3241 if (VECTOR_MODE_P (from
))
3243 from
= GET_MODE_INNER (from
);
3244 to
= GET_MODE_INNER (to
);
3250 case fixuns_trunc_cvt
:
3251 if (GET_MODE_CLASS (from
) != MODE_FLOAT
3252 || GET_MODE_CLASS (to
) != MODE_INT
)
3257 if (GET_MODE_CLASS (from
) != MODE_INT
3258 || GET_MODE_CLASS (to
) != MODE_FLOAT
)
3262 if (GET_MODE_CLASS (from
) != MODE_FLOAT
3263 || GET_MODE_CLASS (to
) != MODE_FLOAT
3264 || GET_MODE_SIZE (from
) >= GET_MODE_SIZE (to
))
3268 if (GET_MODE_CLASS (from
) != MODE_FLOAT
3269 || GET_MODE_CLASS (to
) != MODE_FLOAT
3270 || GET_MODE_SIZE (from
) <= GET_MODE_SIZE (to
))
3275 return ((to
== HImode
&& from
== HFmode
)
3276 || (to
== SImode
&& (from
== SFmode
|| from
== DFmode
))
3277 || (to
== HFmode
&& (from
== HImode
|| from
== SFmode
))
3278 || (to
== SFmode
&& (from
== SImode
|| from
== HFmode
3280 || (to
== DFmode
&& (from
== SImode
|| from
== SFmode
)));
3283 /* Implement TARGET_EMUTLS_VAR_INIT.
3285 Disable emutls (gthr-gcn.h does not support it, yet). */
3288 gcn_emutls_var_init (tree
, tree decl
, tree
)
3290 sorry_at (DECL_SOURCE_LOCATION (decl
), "TLS is not implemented for GCN.");
3297 /* Implement TARGET_RTX_COSTS.
3299 Compute a (partial) cost for rtx X. Return true if the complete
3300 cost has been computed, and false if subexpressions should be
3301 scanned. In either case, *TOTAL contains the cost result. */
3304 gcn_rtx_costs (rtx x
, machine_mode
, int, int, int *total
, bool)
3306 enum rtx_code code
= GET_CODE (x
);
3313 if (gcn_inline_constant_p (x
))
3315 else if (code
== CONST_INT
3316 && ((unsigned HOST_WIDE_INT
) INTVAL (x
) + 0x8000) < 0x10000)
3318 else if (gcn_constant_p (x
))
3321 *total
= vgpr_vector_mode_p (GET_MODE (x
)) ? 64 : 4;
3334 /* Implement TARGET_MEMORY_MOVE_COST.
3336 Return the cost of moving data of mode M between a
3337 register and memory. A value of 2 is the default; this cost is
3338 relative to those in `REGISTER_MOVE_COST'.
3340 This function is used extensively by register_move_cost that is used to
3341 build tables at startup. Make it inline in this case.
3342 When IN is 2, return maximum of in and out move cost.
3344 If moving between registers and memory is more expensive than
3345 between two registers, you should define this macro to express the
3348 Model also increased moving costs of QImode registers in non
3351 #define LOAD_COST 32
3352 #define STORE_COST 32
3354 gcn_memory_move_cost (machine_mode mode
, reg_class_t regclass
, bool in
)
3356 int nregs
= CEIL (GET_MODE_SIZE (mode
), 4);
3359 case SCC_CONDITIONAL_REG
:
3360 case VCCZ_CONDITIONAL_REG
:
3361 case VCC_CONDITIONAL_REG
:
3362 case EXECZ_CONDITIONAL_REG
:
3363 case ALL_CONDITIONAL_REGS
:
3365 case SGPR_EXEC_REGS
:
3367 case SGPR_VOP_SRC_REGS
:
3368 case SGPR_MEM_SRC_REGS
:
3374 return (STORE_COST
+ 2) * nregs
;
3375 return LOAD_COST
* nregs
;
3378 return (LOAD_COST
+ 2) * nregs
;
3379 return STORE_COST
* nregs
;
3384 return (LOAD_COST
+ 2) * nregs
;
3385 return (STORE_COST
+ 2) * nregs
;
3391 /* Implement TARGET_REGISTER_MOVE_COST.
3393 Return the cost of moving data from a register in class CLASS1 to
3394 one in class CLASS2. Base value is 2. */
3397 gcn_register_move_cost (machine_mode
, reg_class_t dst
, reg_class_t src
)
3399 /* Increase cost of moving from and to vector registers. While this is
3400 fast in hardware (I think), it has hidden cost of setting up the exec
3402 if ((src
< VGPR_REGS
) != (dst
< VGPR_REGS
))
3410 /* Type codes used by GCN built-in definitions. */
3412 enum gcn_builtin_type_index
3414 GCN_BTI_END_OF_PARAMS
,
3433 GCN_BTI_LDS_VOIDPTR
,
3438 static GTY(()) tree gcn_builtin_types
[GCN_BTI_MAX
];
3440 #define exec_type_node (gcn_builtin_types[GCN_BTI_EXEC])
3441 #define sf_type_node (gcn_builtin_types[GCN_BTI_SF])
3442 #define v64si_type_node (gcn_builtin_types[GCN_BTI_V64SI])
3443 #define v64sf_type_node (gcn_builtin_types[GCN_BTI_V64SF])
3444 #define v64ptr_type_node (gcn_builtin_types[GCN_BTI_V64PTR])
3445 #define siptr_type_node (gcn_builtin_types[GCN_BTI_SIPTR])
3446 #define sfptr_type_node (gcn_builtin_types[GCN_BTI_SFPTR])
3447 #define voidptr_type_node (gcn_builtin_types[GCN_BTI_VOIDPTR])
3448 #define size_t_type_node (gcn_builtin_types[GCN_BTI_SIZE_T])
3450 static rtx
gcn_expand_builtin_1 (tree
, rtx
, rtx
, machine_mode
, int,
3451 struct gcn_builtin_description
*);
3452 static rtx
gcn_expand_builtin_binop (tree
, rtx
, rtx
, machine_mode
, int,
3453 struct gcn_builtin_description
*);
3455 struct gcn_builtin_description
;
3456 typedef rtx (*gcn_builtin_expander
) (tree
, rtx
, rtx
, machine_mode
, int,
3457 struct gcn_builtin_description
*);
3459 enum gcn_builtin_type
3461 B_UNIMPLEMENTED
, /* Sorry out */
3462 B_INSN
, /* Emit a pattern */
3463 B_OVERLOAD
/* Placeholder for an overloaded function */
3466 struct gcn_builtin_description
3471 enum gcn_builtin_type type
;
3472 /* The first element of parm is always the return type. The rest
3473 are a zero terminated list of parameters. */
3475 gcn_builtin_expander expander
;
3478 /* Read in the GCN builtins from gcn-builtins.def. */
3480 extern GTY(()) struct gcn_builtin_description gcn_builtins
[GCN_BUILTIN_MAX
];
3482 struct gcn_builtin_description gcn_builtins
[] = {
3483 #define DEF_BUILTIN(fcode, icode, name, type, params, expander) \
3484 {GCN_BUILTIN_ ## fcode, icode, name, type, params, expander},
3486 #define DEF_BUILTIN_BINOP_INT_FP(fcode, ic, name) \
3487 {GCN_BUILTIN_ ## fcode ## _V64SI, \
3488 CODE_FOR_ ## ic ##v64si3_exec, name "_v64int", B_INSN, \
3489 {GCN_BTI_V64SI, GCN_BTI_EXEC, GCN_BTI_V64SI, GCN_BTI_V64SI, \
3490 GCN_BTI_V64SI, GCN_BTI_END_OF_PARAMS}, gcn_expand_builtin_binop}, \
3491 {GCN_BUILTIN_ ## fcode ## _V64SI_unspec, \
3492 CODE_FOR_ ## ic ##v64si3_exec, name "_v64int_unspec", B_INSN, \
3493 {GCN_BTI_V64SI, GCN_BTI_EXEC, GCN_BTI_V64SI, GCN_BTI_V64SI, \
3494 GCN_BTI_END_OF_PARAMS}, gcn_expand_builtin_binop},
3496 #include "gcn-builtins.def"
3497 #undef DEF_BUILTIN_BINOP_INT_FP
3501 static GTY(()) tree gcn_builtin_decls
[GCN_BUILTIN_MAX
];
3503 /* Implement TARGET_BUILTIN_DECL.
3505 Return the GCN builtin for CODE. */
3508 gcn_builtin_decl (unsigned code
, bool ARG_UNUSED (initialize_p
))
3510 if (code
>= GCN_BUILTIN_MAX
)
3511 return error_mark_node
;
3513 return gcn_builtin_decls
[code
];
3516 /* Helper function for gcn_init_builtins. */
3519 gcn_init_builtin_types (void)
3521 gcn_builtin_types
[GCN_BTI_VOID
] = void_type_node
;
3522 gcn_builtin_types
[GCN_BTI_BOOL
] = boolean_type_node
;
3523 gcn_builtin_types
[GCN_BTI_INT
] = intSI_type_node
;
3524 gcn_builtin_types
[GCN_BTI_UINT
] = unsigned_type_for (intSI_type_node
);
3525 gcn_builtin_types
[GCN_BTI_SIZE_T
] = size_type_node
;
3526 gcn_builtin_types
[GCN_BTI_LLINT
] = intDI_type_node
;
3527 gcn_builtin_types
[GCN_BTI_LLUINT
] = unsigned_type_for (intDI_type_node
);
3529 exec_type_node
= unsigned_intDI_type_node
;
3530 sf_type_node
= float32_type_node
;
3531 v64si_type_node
= build_vector_type (intSI_type_node
, 64);
3532 v64sf_type_node
= build_vector_type (float_type_node
, 64);
3533 v64ptr_type_node
= build_vector_type (unsigned_intDI_type_node
3534 /*build_pointer_type
3535 (integer_type_node) */
3537 tree tmp
= build_distinct_type_copy (intSI_type_node
);
3538 TYPE_ADDR_SPACE (tmp
) = ADDR_SPACE_FLAT
;
3539 siptr_type_node
= build_pointer_type (tmp
);
3541 tmp
= build_distinct_type_copy (float_type_node
);
3542 TYPE_ADDR_SPACE (tmp
) = ADDR_SPACE_FLAT
;
3543 sfptr_type_node
= build_pointer_type (tmp
);
3545 tmp
= build_distinct_type_copy (void_type_node
);
3546 TYPE_ADDR_SPACE (tmp
) = ADDR_SPACE_FLAT
;
3547 voidptr_type_node
= build_pointer_type (tmp
);
3549 tmp
= build_distinct_type_copy (void_type_node
);
3550 TYPE_ADDR_SPACE (tmp
) = ADDR_SPACE_LDS
;
3551 gcn_builtin_types
[GCN_BTI_LDS_VOIDPTR
] = build_pointer_type (tmp
);
3554 /* Implement TARGET_INIT_BUILTINS.
3556 Set up all builtin functions for this target. */
3559 gcn_init_builtins (void)
3561 gcn_init_builtin_types ();
3563 struct gcn_builtin_description
*d
;
3565 for (i
= 0, d
= gcn_builtins
; i
< GCN_BUILTIN_MAX
; i
++, d
++)
3568 char name
[64]; /* build_function will make a copy. */
3571 /* FIXME: Is this necessary/useful? */
3575 /* Find last parm. */
3576 for (parm
= 1; d
->parm
[parm
] != GCN_BTI_END_OF_PARAMS
; parm
++)
3581 p
= tree_cons (NULL_TREE
, gcn_builtin_types
[d
->parm
[--parm
]], p
);
3583 p
= build_function_type (gcn_builtin_types
[d
->parm
[0]], p
);
3585 sprintf (name
, "__builtin_gcn_%s", d
->name
);
3586 gcn_builtin_decls
[i
]
3587 = add_builtin_function (name
, p
, i
, BUILT_IN_MD
, NULL
, NULL_TREE
);
3589 /* These builtins don't throw. */
3590 TREE_NOTHROW (gcn_builtin_decls
[i
]) = 1;
3593 /* FIXME: remove the ifdef once OpenACC support is merged upstream. */
3594 #ifdef BUILT_IN_GOACC_SINGLE_START
3595 /* These builtins need to take/return an LDS pointer: override the generic
3598 set_builtin_decl (BUILT_IN_GOACC_SINGLE_START
,
3599 gcn_builtin_decls
[GCN_BUILTIN_ACC_SINGLE_START
], false);
3601 set_builtin_decl (BUILT_IN_GOACC_SINGLE_COPY_START
,
3602 gcn_builtin_decls
[GCN_BUILTIN_ACC_SINGLE_COPY_START
],
3605 set_builtin_decl (BUILT_IN_GOACC_SINGLE_COPY_END
,
3606 gcn_builtin_decls
[GCN_BUILTIN_ACC_SINGLE_COPY_END
],
3609 set_builtin_decl (BUILT_IN_GOACC_BARRIER
,
3610 gcn_builtin_decls
[GCN_BUILTIN_ACC_BARRIER
], false);
3614 /* Expand the CMP_SWAP GCN builtins. We have our own versions that do
3615 not require taking the address of any object, other than the memory
3616 cell being operated on.
3618 Helper function for gcn_expand_builtin_1. */
3621 gcn_expand_cmp_swap (tree exp
, rtx target
)
3623 machine_mode mode
= TYPE_MODE (TREE_TYPE (exp
));
3625 = TYPE_ADDR_SPACE (TREE_TYPE (TREE_TYPE (CALL_EXPR_ARG (exp
, 0))));
3626 machine_mode as_mode
= gcn_addr_space_address_mode (as
);
3629 target
= gen_reg_rtx (mode
);
3631 rtx addr
= expand_expr (CALL_EXPR_ARG (exp
, 0),
3632 NULL_RTX
, as_mode
, EXPAND_NORMAL
);
3633 rtx cmp
= expand_expr (CALL_EXPR_ARG (exp
, 1),
3634 NULL_RTX
, mode
, EXPAND_NORMAL
);
3635 rtx src
= expand_expr (CALL_EXPR_ARG (exp
, 2),
3636 NULL_RTX
, mode
, EXPAND_NORMAL
);
3639 rtx mem
= gen_rtx_MEM (mode
, force_reg (as_mode
, addr
));
3640 set_mem_addr_space (mem
, as
);
3643 cmp
= copy_to_mode_reg (mode
, cmp
);
3645 src
= copy_to_mode_reg (mode
, src
);
3648 pat
= gen_sync_compare_and_swapsi (target
, mem
, cmp
, src
);
3650 pat
= gen_sync_compare_and_swapdi (target
, mem
, cmp
, src
);
3657 /* Expand many different builtins.
3659 Intended for use in gcn-builtins.def. */
3662 gcn_expand_builtin_1 (tree exp
, rtx target
, rtx
/*subtarget */ ,
3663 machine_mode
/*mode */ , int ignore
,
3664 struct gcn_builtin_description
*)
3666 tree fndecl
= TREE_OPERAND (CALL_EXPR_FN (exp
), 0);
3667 switch (DECL_MD_FUNCTION_CODE (fndecl
))
3669 case GCN_BUILTIN_FLAT_LOAD_INT32
:
3675 expand_expr (CALL_EXPR_ARG (exp
, 0), NULL_RTX
, DImode
,
3678 force_reg (V64DImode
,
3679 expand_expr (CALL_EXPR_ARG (exp
, 1), NULL_RTX
, V64DImode
,
3681 /*emit_insn (gen_vector_flat_loadv64si
3682 (target, gcn_gen_undef (V64SImode), ptr, exec)); */
3685 case GCN_BUILTIN_FLAT_LOAD_PTR_INT32
:
3686 case GCN_BUILTIN_FLAT_LOAD_PTR_FLOAT
:
3690 rtx exec
= force_reg (DImode
,
3691 expand_expr (CALL_EXPR_ARG (exp
, 0), NULL_RTX
,
3694 rtx ptr
= force_reg (DImode
,
3695 expand_expr (CALL_EXPR_ARG (exp
, 1), NULL_RTX
,
3698 rtx offsets
= force_reg (V64SImode
,
3699 expand_expr (CALL_EXPR_ARG (exp
, 2),
3700 NULL_RTX
, V64DImode
,
3702 rtx addrs
= gen_reg_rtx (V64DImode
);
3703 rtx tmp
= gen_reg_rtx (V64SImode
);
3704 emit_insn (gen_ashlv64si3_exec (tmp
, offsets
,
3706 gcn_gen_undef (V64SImode
), exec
));
3707 emit_insn (gen_addv64di3_zext_dup2_exec (addrs
, tmp
, ptr
,
3708 gcn_gen_undef (V64DImode
),
3710 rtx mem
= gen_rtx_MEM (GET_MODE (target
), addrs
);
3711 /*set_mem_addr_space (mem, ADDR_SPACE_FLAT); */
3712 /* FIXME: set attributes. */
3713 emit_insn (gen_mov_with_exec (target
, mem
, exec
));
3716 case GCN_BUILTIN_FLAT_STORE_PTR_INT32
:
3717 case GCN_BUILTIN_FLAT_STORE_PTR_FLOAT
:
3719 rtx exec
= force_reg (DImode
,
3720 expand_expr (CALL_EXPR_ARG (exp
, 0), NULL_RTX
,
3723 rtx ptr
= force_reg (DImode
,
3724 expand_expr (CALL_EXPR_ARG (exp
, 1), NULL_RTX
,
3727 rtx offsets
= force_reg (V64SImode
,
3728 expand_expr (CALL_EXPR_ARG (exp
, 2),
3729 NULL_RTX
, V64DImode
,
3731 machine_mode vmode
= TYPE_MODE (TREE_TYPE (CALL_EXPR_ARG (exp
,
3733 rtx val
= force_reg (vmode
,
3734 expand_expr (CALL_EXPR_ARG (exp
, 3), NULL_RTX
,
3737 rtx addrs
= gen_reg_rtx (V64DImode
);
3738 rtx tmp
= gen_reg_rtx (V64SImode
);
3739 emit_insn (gen_ashlv64si3_exec (tmp
, offsets
,
3741 gcn_gen_undef (V64SImode
), exec
));
3742 emit_insn (gen_addv64di3_zext_dup2_exec (addrs
, tmp
, ptr
,
3743 gcn_gen_undef (V64DImode
),
3745 rtx mem
= gen_rtx_MEM (vmode
, addrs
);
3746 /*set_mem_addr_space (mem, ADDR_SPACE_FLAT); */
3747 /* FIXME: set attributes. */
3748 emit_insn (gen_mov_with_exec (mem
, val
, exec
));
3751 case GCN_BUILTIN_SQRTVF
:
3755 rtx exec
= gcn_full_exec_reg ();
3756 rtx arg
= force_reg (V64SFmode
,
3757 expand_expr (CALL_EXPR_ARG (exp
, 0), NULL_RTX
,
3760 emit_insn (gen_sqrtv64sf2_exec
3761 (target
, arg
, gcn_gen_undef (V64SFmode
), exec
));
3764 case GCN_BUILTIN_SQRTF
:
3768 rtx arg
= force_reg (SFmode
,
3769 expand_expr (CALL_EXPR_ARG (exp
, 0), NULL_RTX
,
3772 emit_insn (gen_sqrtsf2 (target
, arg
));
3775 case GCN_BUILTIN_OMP_DIM_SIZE
:
3779 emit_insn (gen_oacc_dim_size (target
,
3780 expand_expr (CALL_EXPR_ARG (exp
, 0),
3785 case GCN_BUILTIN_OMP_DIM_POS
:
3789 emit_insn (gen_oacc_dim_pos (target
,
3790 expand_expr (CALL_EXPR_ARG (exp
, 0),
3795 case GCN_BUILTIN_CMP_SWAP
:
3796 case GCN_BUILTIN_CMP_SWAPLL
:
3797 return gcn_expand_cmp_swap (exp
, target
);
3799 case GCN_BUILTIN_ACC_SINGLE_START
:
3804 rtx wavefront
= gcn_oacc_dim_pos (1);
3805 rtx cond
= gen_rtx_EQ (VOIDmode
, wavefront
, const0_rtx
);
3806 rtx cc
= (target
&& REG_P (target
)) ? target
: gen_reg_rtx (BImode
);
3807 emit_insn (gen_cstoresi4 (cc
, cond
, wavefront
, const0_rtx
));
3811 case GCN_BUILTIN_ACC_SINGLE_COPY_START
:
3813 rtx blk
= force_reg (SImode
,
3814 expand_expr (CALL_EXPR_ARG (exp
, 0), NULL_RTX
,
3815 SImode
, EXPAND_NORMAL
));
3816 rtx wavefront
= gcn_oacc_dim_pos (1);
3817 rtx cond
= gen_rtx_NE (VOIDmode
, wavefront
, const0_rtx
);
3818 rtx not_zero
= gen_label_rtx ();
3819 emit_insn (gen_cbranchsi4 (cond
, wavefront
, const0_rtx
, not_zero
));
3820 emit_move_insn (blk
, const0_rtx
);
3821 emit_label (not_zero
);
3825 case GCN_BUILTIN_ACC_SINGLE_COPY_END
:
3828 case GCN_BUILTIN_ACC_BARRIER
:
3829 emit_insn (gen_gcn_wavefront_barrier ());
3837 /* Expansion of simple arithmetic and bit binary operation builtins.
3839 Intended for use with gcn_builtins table. */
3842 gcn_expand_builtin_binop (tree exp
, rtx target
, rtx
/*subtarget */ ,
3843 machine_mode
/*mode */ , int ignore
,
3844 struct gcn_builtin_description
*d
)
3846 int icode
= d
->icode
;
3850 rtx exec
= force_reg (DImode
,
3851 expand_expr (CALL_EXPR_ARG (exp
, 0), NULL_RTX
, DImode
,
3854 machine_mode m1
= insn_data
[icode
].operand
[1].mode
;
3855 rtx arg1
= expand_expr (CALL_EXPR_ARG (exp
, 1), NULL_RTX
, m1
,
3857 if (!insn_data
[icode
].operand
[1].predicate (arg1
, m1
))
3858 arg1
= force_reg (m1
, arg1
);
3860 machine_mode m2
= insn_data
[icode
].operand
[2].mode
;
3861 rtx arg2
= expand_expr (CALL_EXPR_ARG (exp
, 2), NULL_RTX
, m2
,
3863 if (!insn_data
[icode
].operand
[2].predicate (arg2
, m2
))
3864 arg2
= force_reg (m2
, arg2
);
3867 if (call_expr_nargs (exp
) == 4)
3869 machine_mode m_prev
= insn_data
[icode
].operand
[4].mode
;
3870 arg_prev
= force_reg (m_prev
,
3871 expand_expr (CALL_EXPR_ARG (exp
, 3), NULL_RTX
,
3872 m_prev
, EXPAND_NORMAL
));
3875 arg_prev
= gcn_gen_undef (GET_MODE (target
));
3877 rtx pat
= GEN_FCN (icode
) (target
, arg1
, arg2
, exec
, arg_prev
);
3882 /* Implement TARGET_EXPAND_BUILTIN.
3884 Expand an expression EXP that calls a built-in function, with result going
3885 to TARGET if that's convenient (and in mode MODE if that's convenient).
3886 SUBTARGET may be used as the target for computing one of EXP's operands.
3887 IGNORE is nonzero if the value is to be ignored. */
3890 gcn_expand_builtin (tree exp
, rtx target
, rtx subtarget
, machine_mode mode
,
3893 tree fndecl
= TREE_OPERAND (CALL_EXPR_FN (exp
), 0);
3894 unsigned int fcode
= DECL_MD_FUNCTION_CODE (fndecl
);
3895 struct gcn_builtin_description
*d
;
3897 gcc_assert (fcode
< GCN_BUILTIN_MAX
);
3898 d
= &gcn_builtins
[fcode
];
3900 if (d
->type
== B_UNIMPLEMENTED
)
3901 sorry ("Builtin not implemented");
3903 return d
->expander (exp
, target
, subtarget
, mode
, ignore
, d
);
3907 /* {{{ Vectorization. */
3909 /* Implement TARGET_VECTORIZE_GET_MASK_MODE.
3911 A vector mask is a value that holds one boolean result for every element in
3915 gcn_vectorize_get_mask_mode (machine_mode
)
3917 /* GCN uses a DImode bit-mask. */
3921 /* Return an RTX that references a vector with the i-th lane containing
3924 Helper function for gcn_vectorize_vec_perm_const. */
3927 gcn_make_vec_perm_address (unsigned int *perm
)
3929 rtx x
= gen_reg_rtx (V64SImode
);
3930 emit_move_insn (x
, gcn_vec_constant (V64SImode
, 0));
3932 /* Permutation addresses use byte addressing. With each vector lane being
3933 4 bytes wide, and with 64 lanes in total, only bits 2..7 are significant,
3936 The permutation given to the vec_perm* patterns range from 0 to 2N-1 to
3937 select between lanes in two vectors, but as the DS_BPERMUTE* instructions
3938 only take one source vector, the most-significant bit can be ignored
3939 here. Instead, we can use EXEC masking to select the relevant part of
3940 each source vector after they are permuted separately. */
3941 uint64_t bit_mask
= 1 << 2;
3942 for (int i
= 2; i
< 8; i
++, bit_mask
<<= 1)
3944 uint64_t exec_mask
= 0;
3945 uint64_t lane_mask
= 1;
3946 for (int j
= 0; j
< 64; j
++, lane_mask
<<= 1)
3947 if ((perm
[j
] * 4) & bit_mask
)
3948 exec_mask
|= lane_mask
;
3951 emit_insn (gen_addv64si3_exec (x
, x
,
3952 gcn_vec_constant (V64SImode
,
3954 x
, get_exec (exec_mask
)));
3960 /* Implement TARGET_VECTORIZE_VEC_PERM_CONST.
3962 Return true if permutation with SEL is possible.
3964 If DST/SRC0/SRC1 are non-null, emit the instructions to perform the
3968 gcn_vectorize_vec_perm_const (machine_mode vmode
, rtx dst
,
3970 const vec_perm_indices
& sel
)
3972 unsigned int nelt
= GET_MODE_NUNITS (vmode
);
3974 gcc_assert (VECTOR_MODE_P (vmode
));
3975 gcc_assert (nelt
<= 64);
3976 gcc_assert (sel
.length () == nelt
);
3980 /* All vector permutations are possible on this architecture,
3981 with varying degrees of efficiency depending on the permutation. */
3985 unsigned int perm
[64];
3986 for (unsigned int i
= 0; i
< nelt
; ++i
)
3987 perm
[i
] = sel
[i
] & (2 * nelt
- 1);
3989 /* Make life a bit easier by swapping operands if necessary so that
3990 the first element always comes from src0. */
3991 if (perm
[0] >= nelt
)
3997 for (unsigned int i
= 0; i
< nelt
; ++i
)
4004 /* TODO: There are more efficient ways to implement certain permutations
4005 using ds_swizzle_b32 and/or DPP. Test for and expand them here, before
4006 this more inefficient generic approach is used. */
4008 int64_t src1_lanes
= 0;
4009 int64_t lane_bit
= 1;
4011 for (unsigned int i
= 0; i
< nelt
; ++i
, lane_bit
<<= 1)
4013 /* Set the bits for lanes from src1. */
4014 if (perm
[i
] >= nelt
)
4015 src1_lanes
|= lane_bit
;
4018 rtx addr
= gcn_make_vec_perm_address (perm
);
4019 rtx (*ds_bpermute
) (rtx
, rtx
, rtx
, rtx
);
4024 ds_bpermute
= gen_ds_bpermutev64qi
;
4027 ds_bpermute
= gen_ds_bpermutev64hi
;
4030 ds_bpermute
= gen_ds_bpermutev64si
;
4033 ds_bpermute
= gen_ds_bpermutev64hf
;
4036 ds_bpermute
= gen_ds_bpermutev64sf
;
4039 ds_bpermute
= gen_ds_bpermutev64di
;
4042 ds_bpermute
= gen_ds_bpermutev64df
;
4048 /* Load elements from src0 to dst. */
4049 gcc_assert (~src1_lanes
);
4050 emit_insn (ds_bpermute (dst
, addr
, src0
, gcn_full_exec_reg ()));
4052 /* Load elements from src1 to dst. */
4055 /* Masking a lane masks both the destination and source lanes for
4056 DS_BPERMUTE, so we need to have all lanes enabled for the permute,
4057 then add an extra masked move to merge the results of permuting
4058 the two source vectors together.
4060 rtx tmp
= gen_reg_rtx (vmode
);
4061 emit_insn (ds_bpermute (tmp
, addr
, src1
, gcn_full_exec_reg ()));
4062 emit_insn (gen_mov_with_exec (dst
, tmp
, get_exec (src1_lanes
)));
4068 /* Implements TARGET_VECTOR_MODE_SUPPORTED_P.
4070 Return nonzero if vector MODE is supported with at least move
4074 gcn_vector_mode_supported_p (machine_mode mode
)
4076 return (mode
== V64QImode
|| mode
== V64HImode
4077 || mode
== V64SImode
|| mode
== V64DImode
4078 || mode
== V64SFmode
|| mode
== V64DFmode
);
4081 /* Implement TARGET_VECTORIZE_PREFERRED_SIMD_MODE.
4083 Enables autovectorization for all supported modes. */
4086 gcn_vectorize_preferred_simd_mode (scalar_mode mode
)
4107 /* Implement TARGET_VECTORIZE_RELATED_MODE.
4109 All GCN vectors are 64-lane, so this is simpler than other architectures.
4110 In particular, we do *not* want to match vector bit-size. */
4112 static opt_machine_mode
4113 gcn_related_vector_mode (machine_mode
ARG_UNUSED (vector_mode
),
4114 scalar_mode element_mode
, poly_uint64 nunits
)
4116 if (known_ne (nunits
, 0U) && known_ne (nunits
, 64U))
4119 machine_mode pref_mode
= gcn_vectorize_preferred_simd_mode (element_mode
);
4120 if (!VECTOR_MODE_P (pref_mode
))
4126 /* Implement TARGET_VECTORIZE_PREFERRED_VECTOR_ALIGNMENT.
4128 Returns the preferred alignment in bits for accesses to vectors of type type
4129 in vectorized code. This might be less than or greater than the ABI-defined
4130 value returned by TARGET_VECTOR_ALIGNMENT. It can be equal to the alignment
4131 of a single element, in which case the vectorizer will not try to optimize
4135 gcn_preferred_vector_alignment (const_tree type
)
4137 return TYPE_ALIGN (TREE_TYPE (type
));
4140 /* Implement TARGET_VECTORIZE_SUPPORT_VECTOR_MISALIGNMENT.
4142 Return true if the target supports misaligned vector store/load of a
4143 specific factor denoted in the misalignment parameter. */
4146 gcn_vectorize_support_vector_misalignment (machine_mode
ARG_UNUSED (mode
),
4147 const_tree type
, int misalignment
,
4153 /* If the misalignment is unknown, we should be able to handle the access
4154 so long as it is not to a member of a packed data structure. */
4155 if (misalignment
== -1)
4158 /* Return true if the misalignment is a multiple of the natural alignment
4159 of the vector's element type. This is probably always going to be
4160 true in practice, since we've already established that this isn't a
4162 return misalignment
% TYPE_ALIGN_UNIT (type
) == 0;
4165 /* Implement TARGET_VECTORIZE_VECTOR_ALIGNMENT_REACHABLE.
4167 Return true if vector alignment is reachable (by peeling N iterations) for
4168 the given scalar type TYPE. */
4171 gcn_vector_alignment_reachable (const_tree
ARG_UNUSED (type
), bool is_packed
)
4173 /* Vectors which aren't in packed structures will not be less aligned than
4174 the natural alignment of their element type, so this is safe. */
4178 /* Generate DPP instructions used for vector reductions.
4180 The opcode is given by INSN.
4181 The first operand of the operation is shifted right by SHIFT vector lanes.
4182 SHIFT must be a power of 2. If SHIFT is 16, the 15th lane of each row is
4183 broadcast the next row (thereby acting like a shift of 16 for the end of
4184 each row). If SHIFT is 32, lane 31 is broadcast to all the
4185 following lanes (thereby acting like a shift of 32 for lane 63). */
4188 gcn_expand_dpp_shr_insn (machine_mode mode
, const char *insn
,
4189 int unspec
, int shift
)
4191 static char buf
[128];
4193 const char *vcc_in
= "";
4194 const char *vcc_out
= "";
4196 /* Add the vcc operand if needed. */
4197 if (GET_MODE_CLASS (mode
) == MODE_VECTOR_INT
)
4199 if (unspec
== UNSPEC_PLUS_CARRY_IN_DPP_SHR
)
4202 if (unspec
== UNSPEC_PLUS_CARRY_DPP_SHR
4203 || unspec
== UNSPEC_PLUS_CARRY_IN_DPP_SHR
)
4207 /* Add the DPP modifiers. */
4211 dpp
= "row_shr:1 bound_ctrl:0";
4214 dpp
= "row_shr:2 bound_ctrl:0";
4217 dpp
= "row_shr:4 bank_mask:0xe";
4220 dpp
= "row_shr:8 bank_mask:0xc";
4223 dpp
= "row_bcast:15 row_mask:0xa";
4226 dpp
= "row_bcast:31 row_mask:0xc";
4232 if (unspec
== UNSPEC_MOV_DPP_SHR
&& vgpr_2reg_mode_p (mode
))
4233 sprintf (buf
, "%s\t%%L0, %%L1 %s\n\t%s\t%%H0, %%H1 %s",
4234 insn
, dpp
, insn
, dpp
);
4235 else if (unspec
== UNSPEC_MOV_DPP_SHR
)
4236 sprintf (buf
, "%s\t%%0, %%1 %s", insn
, dpp
);
4238 sprintf (buf
, "%s\t%%0%s, %%1, %%2%s %s", insn
, vcc_out
, vcc_in
, dpp
);
4243 /* Generate vector reductions in terms of DPP instructions.
4245 The vector register SRC of mode MODE is reduced using the operation given
4246 by UNSPEC, and the scalar result is returned in lane 63 of a vector
4250 gcn_expand_reduc_scalar (machine_mode mode
, rtx src
, int unspec
)
4252 machine_mode orig_mode
= mode
;
4253 bool use_moves
= (((unspec
== UNSPEC_SMIN_DPP_SHR
4254 || unspec
== UNSPEC_SMAX_DPP_SHR
4255 || unspec
== UNSPEC_UMIN_DPP_SHR
4256 || unspec
== UNSPEC_UMAX_DPP_SHR
)
4257 && mode
== V64DImode
)
4258 || (unspec
== UNSPEC_PLUS_DPP_SHR
4259 && mode
== V64DFmode
));
4260 rtx_code code
= (unspec
== UNSPEC_SMIN_DPP_SHR
? SMIN
4261 : unspec
== UNSPEC_SMAX_DPP_SHR
? SMAX
4262 : unspec
== UNSPEC_UMIN_DPP_SHR
? UMIN
4263 : unspec
== UNSPEC_UMAX_DPP_SHR
? UMAX
4264 : unspec
== UNSPEC_PLUS_DPP_SHR
? PLUS
4266 bool use_extends
= ((unspec
== UNSPEC_SMIN_DPP_SHR
4267 || unspec
== UNSPEC_SMAX_DPP_SHR
4268 || unspec
== UNSPEC_UMIN_DPP_SHR
4269 || unspec
== UNSPEC_UMAX_DPP_SHR
)
4270 && (mode
== V64QImode
4271 || mode
== V64HImode
));
4272 bool unsignedp
= (unspec
== UNSPEC_UMIN_DPP_SHR
4273 || unspec
== UNSPEC_UMAX_DPP_SHR
);
4274 bool use_plus_carry
= unspec
== UNSPEC_PLUS_DPP_SHR
4275 && GET_MODE_CLASS (mode
) == MODE_VECTOR_INT
4276 && (TARGET_GCN3
|| mode
== V64DImode
);
4279 unspec
= UNSPEC_PLUS_CARRY_DPP_SHR
;
4283 rtx tmp
= gen_reg_rtx (V64SImode
);
4284 convert_move (tmp
, src
, unsignedp
);
4289 /* Perform reduction by first performing the reduction operation on every
4290 pair of lanes, then on every pair of results from the previous
4291 iteration (thereby effectively reducing every 4 lanes) and so on until
4292 all lanes are reduced. */
4294 for (int i
= 0, shift
= 1; i
< 6; i
++, shift
<<= 1)
4296 rtx shift_val
= gen_rtx_CONST_INT (VOIDmode
, shift
);
4298 out
= gen_reg_rtx (mode
);
4302 rtx tmp
= gen_reg_rtx (mode
);
4303 emit_insn (gen_dpp_move (mode
, tmp
, in
, shift_val
));
4304 emit_insn (gen_rtx_SET (out
, gen_rtx_fmt_ee (code
, mode
, tmp
, in
)));
4308 rtx insn
= gen_rtx_SET (out
,
4309 gen_rtx_UNSPEC (mode
,
4310 gen_rtvec (3, in
, in
,
4314 /* Add clobber for instructions that set the carry flags. */
4317 rtx clobber
= gen_rtx_CLOBBER (VOIDmode
,
4318 gen_rtx_REG (DImode
, VCC_REG
));
4319 insn
= gen_rtx_PARALLEL (VOIDmode
,
4320 gen_rtvec (2, insn
, clobber
));
4329 rtx tmp
= gen_reg_rtx (orig_mode
);
4330 convert_move (tmp
, out
, unsignedp
);
4337 /* Implement TARGET_VECTORIZE_BUILTIN_VECTORIZATION_COST. */
4340 gcn_vectorization_cost (enum vect_cost_for_stmt
ARG_UNUSED (type_of_cost
),
4341 tree
ARG_UNUSED (vectype
), int ARG_UNUSED (misalign
))
4343 /* Always vectorize. */
4348 /* {{{ md_reorg pass. */
4350 /* Identify VMEM instructions from their "type" attribute. */
4353 gcn_vmem_insn_p (attr_type type
)
4384 /* If INSN sets the EXEC register to a constant value, return the value,
4385 otherwise return zero. */
4388 gcn_insn_exec_value (rtx_insn
*insn
)
4390 if (!NONDEBUG_INSN_P (insn
))
4393 rtx pattern
= PATTERN (insn
);
4395 if (GET_CODE (pattern
) == SET
)
4397 rtx dest
= XEXP (pattern
, 0);
4398 rtx src
= XEXP (pattern
, 1);
4400 if (GET_MODE (dest
) == DImode
4401 && REG_P (dest
) && REGNO (dest
) == EXEC_REG
4402 && CONST_INT_P (src
))
4403 return INTVAL (src
);
4409 /* Sets the EXEC register before INSN to the value that it had after
4410 LAST_EXEC_DEF. The constant value of the EXEC register is returned if
4411 known, otherwise it returns zero. */
4414 gcn_restore_exec (rtx_insn
*insn
, rtx_insn
*last_exec_def
, int64_t curr_exec
,
4415 bool curr_exec_known
, bool &last_exec_def_saved
)
4417 rtx exec_reg
= gen_rtx_REG (DImode
, EXEC_REG
);
4420 int64_t exec_value
= gcn_insn_exec_value (last_exec_def
);
4424 /* If the EXEC value is a constant and it happens to be the same as the
4425 current EXEC value, the restore can be skipped. */
4426 if (curr_exec_known
&& exec_value
== curr_exec
)
4429 exec
= GEN_INT (exec_value
);
4433 /* If the EXEC value is not a constant, save it in a register after the
4434 point of definition. */
4435 rtx exec_save_reg
= gen_rtx_REG (DImode
, EXEC_SAVE_REG
);
4437 if (!last_exec_def_saved
)
4440 emit_move_insn (exec_save_reg
, exec_reg
);
4441 rtx_insn
*seq
= get_insns ();
4444 emit_insn_after (seq
, last_exec_def
);
4445 if (dump_file
&& (dump_flags
& TDF_DETAILS
))
4446 fprintf (dump_file
, "Saving EXEC after insn %d.\n",
4447 INSN_UID (last_exec_def
));
4449 last_exec_def_saved
= true;
4452 exec
= exec_save_reg
;
4455 /* Restore EXEC register before the usage. */
4457 emit_move_insn (exec_reg
, exec
);
4458 rtx_insn
*seq
= get_insns ();
4460 emit_insn_before (seq
, insn
);
4462 if (dump_file
&& (dump_flags
& TDF_DETAILS
))
4465 fprintf (dump_file
, "Restoring EXEC to %ld before insn %d.\n",
4466 exec_value
, INSN_UID (insn
));
4469 "Restoring EXEC from saved value before insn %d.\n",
4476 /* Implement TARGET_MACHINE_DEPENDENT_REORG.
4478 Ensure that pipeline dependencies and lane masking are set correctly. */
4484 rtx exec_reg
= gen_rtx_REG (DImode
, EXEC_REG
);
4487 INIT_REG_SET (&live
);
4489 compute_bb_for_insn ();
4494 if (dump_file
&& (dump_flags
& TDF_DETAILS
))
4496 fprintf (dump_file
, "After split:\n");
4497 print_rtl_with_bb (dump_file
, get_insns (), dump_flags
);
4500 /* Update data-flow information for split instructions. */
4501 df_insn_rescan_all ();
4506 /* This pass ensures that the EXEC register is set correctly, according
4507 to the "exec" attribute. However, care must be taken so that the
4508 value that reaches explicit uses of the EXEC register remains the
4512 FOR_EACH_BB_FN (bb
, cfun
)
4514 if (dump_file
&& (dump_flags
& TDF_DETAILS
))
4515 fprintf (dump_file
, "BB %d:\n", bb
->index
);
4517 rtx_insn
*insn
, *curr
;
4518 rtx_insn
*last_exec_def
= BB_HEAD (bb
);
4519 bool last_exec_def_saved
= false;
4520 bool curr_exec_explicit
= true;
4521 bool curr_exec_known
= true;
4522 int64_t curr_exec
= 0; /* 0 here means 'the value is that of EXEC
4523 after last_exec_def is executed'. */
4525 FOR_BB_INSNS_SAFE (bb
, insn
, curr
)
4527 if (!NONDEBUG_INSN_P (insn
))
4530 if (GET_CODE (PATTERN (insn
)) == USE
4531 || GET_CODE (PATTERN (insn
)) == CLOBBER
)
4534 HARD_REG_SET defs
, uses
;
4535 CLEAR_HARD_REG_SET (defs
);
4536 CLEAR_HARD_REG_SET (uses
);
4537 note_stores (insn
, record_hard_reg_sets
, &defs
);
4538 note_uses (&PATTERN (insn
), record_hard_reg_uses
, &uses
);
4540 bool exec_lo_def_p
= TEST_HARD_REG_BIT (defs
, EXEC_LO_REG
);
4541 bool exec_hi_def_p
= TEST_HARD_REG_BIT (defs
, EXEC_HI_REG
);
4542 bool exec_used
= (hard_reg_set_intersect_p
4543 (uses
, reg_class_contents
[(int) EXEC_MASK_REG
])
4544 || TEST_HARD_REG_BIT (uses
, EXECZ_REG
));
4546 /* Check the instruction for implicit setting of EXEC via an
4548 attr_exec exec_attr
= get_attr_exec (insn
);
4558 /* Instructions that do not involve memory accesses only require
4559 bit 0 of EXEC to be set. */
4560 if (gcn_vmem_insn_p (get_attr_type (insn
))
4561 || get_attr_type (insn
) == TYPE_DS
)
4564 new_exec
= curr_exec
| 1;
4571 default: /* Auto-detect what setting is appropriate. */
4575 /* If EXEC is referenced explicitly then we don't need to do
4576 anything to set it, so we're done. */
4580 /* Scan the insn for VGPRs defs or uses. The mode determines
4581 what kind of exec is needed. */
4582 subrtx_iterator::array_type array
;
4583 FOR_EACH_SUBRTX (iter
, array
, PATTERN (insn
), NONCONST
)
4585 const_rtx x
= *iter
;
4586 if (REG_P (x
) && VGPR_REGNO_P (REGNO (x
)))
4588 if (VECTOR_MODE_P (GET_MODE (x
)))
4601 if (new_exec
&& (!curr_exec_known
|| new_exec
!= curr_exec
))
4604 emit_move_insn (exec_reg
, GEN_INT (new_exec
));
4605 rtx_insn
*seq
= get_insns ();
4607 emit_insn_before (seq
, insn
);
4609 if (dump_file
&& (dump_flags
& TDF_DETAILS
))
4610 fprintf (dump_file
, "Setting EXEC to %ld before insn %d.\n",
4611 new_exec
, INSN_UID (insn
));
4613 curr_exec
= new_exec
;
4614 curr_exec_explicit
= false;
4615 curr_exec_known
= true;
4617 else if (new_exec
&& dump_file
&& (dump_flags
& TDF_DETAILS
))
4619 fprintf (dump_file
, "Exec already is %ld before insn %d.\n",
4620 new_exec
, INSN_UID (insn
));
4623 /* The state of the EXEC register is unknown after a
4626 curr_exec_known
= false;
4628 /* Handle explicit uses of EXEC. If the instruction is a partial
4629 explicit definition of EXEC, then treat it as an explicit use of
4631 if (exec_used
|| exec_lo_def_p
!= exec_hi_def_p
)
4633 /* An instruction that explicitly uses EXEC should not also
4634 implicitly define it. */
4635 gcc_assert (!exec_used
|| !new_exec
);
4637 if (!curr_exec_known
|| !curr_exec_explicit
)
4639 /* Restore the previous explicitly defined value. */
4640 curr_exec
= gcn_restore_exec (insn
, last_exec_def
,
4641 curr_exec
, curr_exec_known
,
4642 last_exec_def_saved
);
4643 curr_exec_explicit
= true;
4644 curr_exec_known
= true;
4648 /* Handle explicit definitions of EXEC. */
4649 if (exec_lo_def_p
|| exec_hi_def_p
)
4651 last_exec_def
= insn
;
4652 last_exec_def_saved
= false;
4653 curr_exec
= gcn_insn_exec_value (insn
);
4654 curr_exec_explicit
= true;
4655 curr_exec_known
= true;
4657 if (dump_file
&& (dump_flags
& TDF_DETAILS
))
4659 "Found %s definition of EXEC at insn %d.\n",
4660 exec_lo_def_p
== exec_hi_def_p
? "full" : "partial",
4665 COPY_REG_SET (&live
, DF_LR_OUT (bb
));
4666 df_simulate_initialize_backwards (bb
, &live
);
4668 /* If EXEC is live after the basic block, restore the value of EXEC
4669 at the end of the block. */
4670 if ((REGNO_REG_SET_P (&live
, EXEC_LO_REG
)
4671 || REGNO_REG_SET_P (&live
, EXEC_HI_REG
))
4672 && (!curr_exec_known
|| !curr_exec_explicit
))
4674 rtx_insn
*end_insn
= BB_END (bb
);
4676 /* If the instruction is not a jump instruction, do the restore
4677 after the last instruction in the basic block. */
4678 if (NONJUMP_INSN_P (end_insn
))
4679 end_insn
= NEXT_INSN (end_insn
);
4681 gcn_restore_exec (end_insn
, last_exec_def
, curr_exec
,
4682 curr_exec_known
, last_exec_def_saved
);
4686 CLEAR_REG_SET (&live
);
4688 /* "Manually Inserted Wait States (NOPs)."
4690 GCN hardware detects most kinds of register dependencies, but there
4691 are some exceptions documented in the ISA manual. This pass
4692 detects the missed cases, and inserts the documented number of NOPs
4693 required for correct execution. */
4695 const int max_waits
= 5;
4700 attr_delayeduse delayeduse
;
4701 HARD_REG_SET writes
;
4706 for (int i
= 0; i
< max_waits
; i
++)
4707 back
[i
].insn
= NULL
;
4709 rtx_insn
*insn
, *last_insn
= NULL
;
4710 for (insn
= get_insns (); insn
!= 0; insn
= NEXT_INSN (insn
))
4712 if (!NONDEBUG_INSN_P (insn
))
4715 if (GET_CODE (PATTERN (insn
)) == USE
4716 || GET_CODE (PATTERN (insn
)) == CLOBBER
)
4719 attr_type itype
= get_attr_type (insn
);
4720 attr_unit iunit
= get_attr_unit (insn
);
4721 attr_delayeduse idelayeduse
= get_attr_delayeduse (insn
);
4722 HARD_REG_SET ireads
, iwrites
;
4723 CLEAR_HARD_REG_SET (ireads
);
4724 CLEAR_HARD_REG_SET (iwrites
);
4725 note_stores (insn
, record_hard_reg_sets
, &iwrites
);
4726 note_uses (&PATTERN (insn
), record_hard_reg_uses
, &ireads
);
4728 /* Scan recent previous instructions for dependencies not handled in
4731 for (int i
= oldest
; i
< oldest
+ max_waits
; i
++)
4733 struct ilist
*prev_insn
= &back
[i
% max_waits
];
4735 if (!prev_insn
->insn
)
4738 /* VALU writes SGPR followed by VMEM reading the same SGPR
4739 requires 5 wait states. */
4740 if ((prev_insn
->age
+ nops_rqd
) < 5
4741 && prev_insn
->unit
== UNIT_VECTOR
4742 && gcn_vmem_insn_p (itype
))
4744 HARD_REG_SET regs
= prev_insn
->writes
& ireads
;
4745 if (hard_reg_set_intersect_p
4746 (regs
, reg_class_contents
[(int) SGPR_REGS
]))
4747 nops_rqd
= 5 - prev_insn
->age
;
4750 /* VALU sets VCC/EXEC followed by VALU uses VCCZ/EXECZ
4751 requires 5 wait states. */
4752 if ((prev_insn
->age
+ nops_rqd
) < 5
4753 && prev_insn
->unit
== UNIT_VECTOR
4754 && iunit
== UNIT_VECTOR
4755 && ((hard_reg_set_intersect_p
4757 reg_class_contents
[(int) EXEC_MASK_REG
])
4758 && TEST_HARD_REG_BIT (ireads
, EXECZ_REG
))
4760 (hard_reg_set_intersect_p
4762 reg_class_contents
[(int) VCC_CONDITIONAL_REG
])
4763 && TEST_HARD_REG_BIT (ireads
, VCCZ_REG
))))
4764 nops_rqd
= 5 - prev_insn
->age
;
4766 /* VALU writes SGPR/VCC followed by v_{read,write}lane using
4767 SGPR/VCC as lane select requires 4 wait states. */
4768 if ((prev_insn
->age
+ nops_rqd
) < 4
4769 && prev_insn
->unit
== UNIT_VECTOR
4770 && get_attr_laneselect (insn
) == LANESELECT_YES
)
4772 HARD_REG_SET regs
= prev_insn
->writes
& ireads
;
4773 if (hard_reg_set_intersect_p
4774 (regs
, reg_class_contents
[(int) SGPR_REGS
])
4775 || hard_reg_set_intersect_p
4776 (regs
, reg_class_contents
[(int) VCC_CONDITIONAL_REG
]))
4777 nops_rqd
= 4 - prev_insn
->age
;
4780 /* VALU writes VGPR followed by VALU_DPP reading that VGPR
4781 requires 2 wait states. */
4782 if ((prev_insn
->age
+ nops_rqd
) < 2
4783 && prev_insn
->unit
== UNIT_VECTOR
4784 && itype
== TYPE_VOP_DPP
)
4786 HARD_REG_SET regs
= prev_insn
->writes
& ireads
;
4787 if (hard_reg_set_intersect_p
4788 (regs
, reg_class_contents
[(int) VGPR_REGS
]))
4789 nops_rqd
= 2 - prev_insn
->age
;
4792 /* Store that requires input registers are not overwritten by
4793 following instruction. */
4794 if ((prev_insn
->age
+ nops_rqd
) < 1
4795 && prev_insn
->delayeduse
== DELAYEDUSE_YES
4796 && ((hard_reg_set_intersect_p
4797 (prev_insn
->reads
, iwrites
))))
4798 nops_rqd
= 1 - prev_insn
->age
;
4801 /* Insert the required number of NOPs. */
4802 for (int i
= nops_rqd
; i
> 0; i
--)
4803 emit_insn_after (gen_nop (), last_insn
);
4805 /* Age the previous instructions. We can also ignore writes to
4806 registers subsequently overwritten. */
4807 HARD_REG_SET written
;
4808 CLEAR_HARD_REG_SET (written
);
4809 for (int i
= oldest
+ max_waits
- 1; i
> oldest
; i
--)
4811 struct ilist
*prev_insn
= &back
[i
% max_waits
];
4813 /* Assume all instructions are equivalent to one "wait", the same
4814 as s_nop. This is probably true for SALU, but not VALU (which
4815 may take longer), so this is not optimal. However, AMD do
4816 not publish the cycle times for instructions. */
4817 prev_insn
->age
+= 1 + nops_rqd
;
4820 prev_insn
->writes
&= ~written
;
4823 /* Track the current instruction as a previous instruction. */
4824 back
[oldest
].insn
= insn
;
4825 back
[oldest
].unit
= iunit
;
4826 back
[oldest
].delayeduse
= idelayeduse
;
4827 back
[oldest
].writes
= iwrites
;
4828 back
[oldest
].reads
= ireads
;
4829 back
[oldest
].age
= 0;
4830 oldest
= (oldest
+ 1) % max_waits
;
4837 /* {{{ OpenACC / OpenMP. */
4839 #define GCN_DEFAULT_GANGS 0 /* Choose at runtime. */
4840 #define GCN_DEFAULT_WORKERS 0 /* Choose at runtime. */
4841 #define GCN_DEFAULT_VECTORS 1 /* Use autovectorization only, for now. */
4843 /* Implement TARGET_GOACC_VALIDATE_DIMS.
4845 Check the launch dimensions provided for an OpenACC compute
4846 region, or routine. */
4849 gcn_goacc_validate_dims (tree decl
, int dims
[], int fn_level
,
4852 bool changed
= false;
4854 /* FIXME: remove -facc-experimental-workers when they're ready. */
4855 int max_workers
= flag_worker_partitioning
? 16 : 1;
4857 gcc_assert (!flag_worker_partitioning
);
4859 /* The vector size must appear to be 64, to the user, unless this is a
4860 SEQ routine. The real, internal value is always 1, which means use
4861 autovectorization, but the user should not see that. */
4862 if (fn_level
<= GOMP_DIM_VECTOR
&& fn_level
>= -1
4863 && dims
[GOMP_DIM_VECTOR
] >= 0)
4865 if (fn_level
< 0 && dims
[GOMP_DIM_VECTOR
] >= 0
4866 && dims
[GOMP_DIM_VECTOR
] != 64)
4867 warning_at (decl
? DECL_SOURCE_LOCATION (decl
) : UNKNOWN_LOCATION
,
4869 (dims
[GOMP_DIM_VECTOR
]
4870 ? G_("using vector_length (64), ignoring %d")
4871 : G_("using vector_length (64), "
4872 "ignoring runtime setting")),
4873 dims
[GOMP_DIM_VECTOR
]);
4874 dims
[GOMP_DIM_VECTOR
] = 1;
4878 /* Check the num workers is not too large. */
4879 if (dims
[GOMP_DIM_WORKER
] > max_workers
)
4881 warning_at (decl
? DECL_SOURCE_LOCATION (decl
) : UNKNOWN_LOCATION
,
4883 "using num_workers (%d), ignoring %d",
4884 max_workers
, dims
[GOMP_DIM_WORKER
]);
4885 dims
[GOMP_DIM_WORKER
] = max_workers
;
4889 /* Set global defaults. */
4892 dims
[GOMP_DIM_VECTOR
] = GCN_DEFAULT_VECTORS
;
4893 if (dims
[GOMP_DIM_WORKER
] < 0)
4894 dims
[GOMP_DIM_WORKER
] = (flag_worker_partitioning
4895 ? GCN_DEFAULT_WORKERS
: 1);
4896 if (dims
[GOMP_DIM_GANG
] < 0)
4897 dims
[GOMP_DIM_GANG
] = GCN_DEFAULT_GANGS
;
4904 /* Helper function for oacc_dim_size instruction.
4905 Also used for OpenMP, via builtin_gcn_dim_size, and the omp_gcn pass. */
4908 gcn_oacc_dim_size (int dim
)
4910 if (dim
< 0 || dim
> 2)
4911 error ("offload dimension out of range (%d)", dim
);
4913 /* Vectors are a special case. */
4915 return const1_rtx
; /* Think of this as 1 times 64. */
4917 static int offset
[] = {
4918 /* Offsets into dispatch packet. */
4919 12, /* X dim = Gang / Team / Work-group. */
4920 20, /* Z dim = Worker / Thread / Wavefront. */
4921 16 /* Y dim = Vector / SIMD / Work-item. */
4923 rtx addr
= gen_rtx_PLUS (DImode
,
4924 gen_rtx_REG (DImode
,
4925 cfun
->machine
->args
.
4926 reg
[DISPATCH_PTR_ARG
]),
4927 GEN_INT (offset
[dim
]));
4928 return gen_rtx_MEM (SImode
, addr
);
4931 /* Helper function for oacc_dim_pos instruction.
4932 Also used for OpenMP, via builtin_gcn_dim_pos, and the omp_gcn pass. */
4935 gcn_oacc_dim_pos (int dim
)
4937 if (dim
< 0 || dim
> 2)
4938 error ("offload dimension out of range (%d)", dim
);
4940 static const int reg
[] = {
4941 WORKGROUP_ID_X_ARG
, /* Gang / Team / Work-group. */
4942 WORK_ITEM_ID_Z_ARG
, /* Worker / Thread / Wavefront. */
4943 WORK_ITEM_ID_Y_ARG
/* Vector / SIMD / Work-item. */
4946 int reg_num
= cfun
->machine
->args
.reg
[reg
[dim
]];
4948 /* The information must have been requested by the kernel. */
4949 gcc_assert (reg_num
>= 0);
4951 return gen_rtx_REG (SImode
, reg_num
);
4954 /* Implement TARGET_GOACC_FORK_JOIN. */
4957 gcn_fork_join (gcall
*ARG_UNUSED (call
), const int *ARG_UNUSED (dims
),
4958 bool ARG_UNUSED (is_fork
))
4960 /* GCN does not use the fork/join concept invented for NVPTX.
4961 Instead we use standard autovectorization. */
4965 /* Implement ???????
4966 FIXME make this a real hook.
4968 Adjust FNDECL such that options inherited from the host compiler
4969 are made appropriate for the accelerator compiler. */
4972 gcn_fixup_accel_lto_options (tree fndecl
)
4974 tree func_optimize
= DECL_FUNCTION_SPECIFIC_OPTIMIZATION (fndecl
);
4979 = build_optimization_node (&global_options
, &global_options_set
);
4982 /* If the function changed the optimization levels as well as
4983 setting target options, start with the optimizations
4985 if (func_optimize
!= old_optimize
)
4986 cl_optimization_restore (&global_options
, &global_options_set
,
4987 TREE_OPTIMIZATION (func_optimize
));
4989 gcn_option_override ();
4991 /* The target attributes may also change some optimization flags,
4992 so update the optimization options if necessary. */
4993 new_optimize
= build_optimization_node (&global_options
,
4994 &global_options_set
);
4996 if (old_optimize
!= new_optimize
)
4998 DECL_FUNCTION_SPECIFIC_OPTIMIZATION (fndecl
) = new_optimize
;
4999 cl_optimization_restore (&global_options
, &global_options_set
,
5000 TREE_OPTIMIZATION (old_optimize
));
5005 /* {{{ ASM Output. */
5007 /* Implement TARGET_ASM_FILE_START.
5009 Print assembler file header text. */
5012 output_file_start (void)
5017 case PROCESSOR_FIJI
: cpu
= "gfx803"; break;
5018 case PROCESSOR_VEGA10
: cpu
= "gfx900"; break;
5019 case PROCESSOR_VEGA20
: cpu
= "gfx906"; break;
5020 default: gcc_unreachable ();
5023 fprintf(asm_out_file
, "\t.amdgcn_target \"amdgcn-unknown-amdhsa--%s\"\n", cpu
);
5026 /* Implement ASM_DECLARE_FUNCTION_NAME via gcn-hsa.h.
5028 Print the initial definition of a function name.
5030 For GCN kernel entry points this includes all the HSA meta-data, special
5031 alignment constraints that don't apply to regular functions, and magic
5032 comments that pass information to mkoffload. */
5035 gcn_hsa_declare_function_name (FILE *file
, const char *name
, tree
)
5038 bool xnack_enabled
= false;
5040 fputs ("\n\n", file
);
5042 if (cfun
&& cfun
->machine
&& cfun
->machine
->normal_function
)
5044 fputs ("\t.type\t", file
);
5045 assemble_name (file
, name
);
5046 fputs (",@function\n", file
);
5047 assemble_name (file
, name
);
5048 fputs (":\n", file
);
5052 /* Determine count of sgpr/vgpr registers by looking for last
5054 for (sgpr
= 101; sgpr
>= 0; sgpr
--)
5055 if (df_regs_ever_live_p (FIRST_SGPR_REG
+ sgpr
))
5058 for (vgpr
= 255; vgpr
>= 0; vgpr
--)
5059 if (df_regs_ever_live_p (FIRST_VGPR_REG
+ vgpr
))
5063 if (!leaf_function_p ())
5065 /* We can't know how many registers function calls might use. */
5066 if (vgpr
< MAX_NORMAL_VGPR_COUNT
)
5067 vgpr
= MAX_NORMAL_VGPR_COUNT
;
5068 if (sgpr
< MAX_NORMAL_SGPR_COUNT
)
5069 sgpr
= MAX_NORMAL_SGPR_COUNT
;
5072 fputs ("\t.rodata\n"
5074 "\t.amdhsa_kernel\t", file
);
5075 assemble_name (file
, name
);
5077 int reg
= FIRST_SGPR_REG
;
5078 for (int a
= 0; a
< GCN_KERNEL_ARG_TYPES
; a
++)
5082 if ((cfun
->machine
->args
.requested
& (1 << a
))
5083 && (gcn_kernel_arg_types
[a
].fixed_regno
< 0))
5086 reg_last
= (reg_first
5087 + (GET_MODE_SIZE (gcn_kernel_arg_types
[a
].mode
)
5088 / UNITS_PER_WORD
) - 1);
5092 if (gcn_kernel_arg_types
[a
].header_pseudo
)
5094 fprintf (file
, "\t %s%s\t%i",
5095 (cfun
->machine
->args
.requested
& (1 << a
)) != 0 ? "" : ";",
5096 gcn_kernel_arg_types
[a
].header_pseudo
,
5097 (cfun
->machine
->args
.requested
& (1 << a
)) != 0);
5098 if (reg_first
!= -1)
5100 fprintf (file
, " ; (");
5101 for (int i
= reg_first
; i
<= reg_last
; ++i
)
5104 fprintf (file
, ", ");
5105 fprintf (file
, "%s", reg_names
[i
]);
5107 fprintf (file
, ")");
5109 fprintf (file
, "\n");
5111 else if (gcn_kernel_arg_types
[a
].fixed_regno
>= 0
5112 && cfun
->machine
->args
.requested
& (1 << a
))
5113 fprintf (file
, "\t ; %s\t%i (%s)\n",
5114 gcn_kernel_arg_types
[a
].name
,
5115 (cfun
->machine
->args
.requested
& (1 << a
)) != 0,
5116 reg_names
[gcn_kernel_arg_types
[a
].fixed_regno
]);
5118 fprintf (file
, "\t .amdhsa_system_vgpr_workitem_id\t%i\n",
5119 (cfun
->machine
->args
.requested
& (1 << WORK_ITEM_ID_Z_ARG
))
5121 : cfun
->machine
->args
.requested
& (1 << WORK_ITEM_ID_Y_ARG
)
5124 "\t .amdhsa_next_free_vgpr\t%i\n"
5125 "\t .amdhsa_next_free_sgpr\t%i\n"
5126 "\t .amdhsa_reserve_vcc\t1\n"
5127 "\t .amdhsa_reserve_flat_scratch\t0\n"
5128 "\t .amdhsa_reserve_xnack_mask\t%i\n"
5129 "\t .amdhsa_private_segment_fixed_size\t%i\n"
5130 "\t .amdhsa_group_segment_fixed_size\t%u\n"
5131 "\t .amdhsa_float_denorm_mode_32\t3\n"
5132 "\t .amdhsa_float_denorm_mode_16_64\t3\n",
5136 /* workitem_private_segment_bytes_size needs to be
5137 one 64th the wave-front stack size. */
5138 stack_size_opt
/ 64,
5140 fputs ("\t.end_amdhsa_kernel\n", file
);
5143 /* The following is YAML embedded in assembler; tabs are not allowed. */
5144 fputs (" .amdgpu_metadata\n"
5145 " amdhsa.version:\n"
5148 " amdhsa.kernels:\n"
5149 " - .name: ", file
);
5150 assemble_name (file
, name
);
5151 fputs ("\n .symbol: ", file
);
5152 assemble_name (file
, name
);
5155 " .kernarg_segment_size: %i\n"
5156 " .kernarg_segment_align: %i\n"
5157 " .group_segment_fixed_size: %u\n"
5158 " .private_segment_fixed_size: %i\n"
5159 " .wavefront_size: 64\n"
5160 " .sgpr_count: %i\n"
5161 " .vgpr_count: %i\n"
5162 " .max_flat_workgroup_size: 1024\n",
5163 cfun
->machine
->kernarg_segment_byte_size
,
5164 cfun
->machine
->kernarg_segment_alignment
,
5166 stack_size_opt
/ 64,
5168 fputs (" .end_amdgpu_metadata\n", file
);
5171 fputs ("\t.text\n", file
);
5172 fputs ("\t.align\t256\n", file
);
5173 fputs ("\t.type\t", file
);
5174 assemble_name (file
, name
);
5175 fputs (",@function\n", file
);
5176 assemble_name (file
, name
);
5177 fputs (":\n", file
);
5179 /* This comment is read by mkoffload. */
5181 fprintf (file
, "\t;; OPENACC-DIMS: %d, %d, %d : %s\n",
5182 oacc_get_fn_dim_size (cfun
->decl
, GOMP_DIM_GANG
),
5183 oacc_get_fn_dim_size (cfun
->decl
, GOMP_DIM_WORKER
),
5184 oacc_get_fn_dim_size (cfun
->decl
, GOMP_DIM_VECTOR
), name
);
5187 /* Implement TARGET_ASM_SELECT_SECTION.
5189 Return the section into which EXP should be placed. */
5192 gcn_asm_select_section (tree exp
, int reloc
, unsigned HOST_WIDE_INT align
)
5194 if (TREE_TYPE (exp
) != error_mark_node
5195 && TYPE_ADDR_SPACE (TREE_TYPE (exp
)) == ADDR_SPACE_LDS
)
5198 return get_section (".lds_bss",
5199 SECTION_WRITE
| SECTION_BSS
| SECTION_DEBUG
,
5202 return get_named_section (exp
, ".lds_bss", reloc
);
5205 return default_elf_select_section (exp
, reloc
, align
);
5208 /* Implement TARGET_ASM_FUNCTION_PROLOGUE.
5210 Emits custom text into the assembler file at the head of each function. */
5213 gcn_target_asm_function_prologue (FILE *file
)
5215 machine_function
*offsets
= gcn_compute_frame_offsets ();
5217 asm_fprintf (file
, "\t; using %s addressing in function\n",
5218 offsets
->use_flat_addressing
? "flat" : "global");
5220 if (offsets
->normal_function
)
5222 asm_fprintf (file
, "\t; frame pointer needed: %s\n",
5223 offsets
->need_frame_pointer
? "true" : "false");
5224 asm_fprintf (file
, "\t; lr needs saving: %s\n",
5225 offsets
->lr_needs_saving
? "true" : "false");
5226 asm_fprintf (file
, "\t; outgoing args size: %wd\n",
5227 offsets
->outgoing_args_size
);
5228 asm_fprintf (file
, "\t; pretend size: %wd\n", offsets
->pretend_size
);
5229 asm_fprintf (file
, "\t; local vars size: %wd\n", offsets
->local_vars
);
5230 asm_fprintf (file
, "\t; callee save size: %wd\n",
5231 offsets
->callee_saves
);
5235 asm_fprintf (file
, "\t; HSA kernel entry point\n");
5236 asm_fprintf (file
, "\t; local vars size: %wd\n", offsets
->local_vars
);
5237 asm_fprintf (file
, "\t; outgoing args size: %wd\n",
5238 offsets
->outgoing_args_size
);
5242 /* Helper function for print_operand and print_operand_address.
5244 Print a register as the assembler requires, according to mode and name. */
5247 print_reg (FILE *file
, rtx x
)
5249 machine_mode mode
= GET_MODE (x
);
5250 if (mode
== BImode
|| mode
== QImode
|| mode
== HImode
|| mode
== SImode
5251 || mode
== HFmode
|| mode
== SFmode
5252 || mode
== V64SFmode
|| mode
== V64SImode
5253 || mode
== V64QImode
|| mode
== V64HImode
)
5254 fprintf (file
, "%s", reg_names
[REGNO (x
)]);
5255 else if (mode
== DImode
|| mode
== V64DImode
5256 || mode
== DFmode
|| mode
== V64DFmode
)
5258 if (SGPR_REGNO_P (REGNO (x
)))
5259 fprintf (file
, "s[%i:%i]", REGNO (x
) - FIRST_SGPR_REG
,
5260 REGNO (x
) - FIRST_SGPR_REG
+ 1);
5261 else if (VGPR_REGNO_P (REGNO (x
)))
5262 fprintf (file
, "v[%i:%i]", REGNO (x
) - FIRST_VGPR_REG
,
5263 REGNO (x
) - FIRST_VGPR_REG
+ 1);
5264 else if (REGNO (x
) == FLAT_SCRATCH_REG
)
5265 fprintf (file
, "flat_scratch");
5266 else if (REGNO (x
) == EXEC_REG
)
5267 fprintf (file
, "exec");
5268 else if (REGNO (x
) == VCC_LO_REG
)
5269 fprintf (file
, "vcc");
5271 fprintf (file
, "[%s:%s]",
5272 reg_names
[REGNO (x
)], reg_names
[REGNO (x
) + 1]);
5274 else if (mode
== TImode
)
5276 if (SGPR_REGNO_P (REGNO (x
)))
5277 fprintf (file
, "s[%i:%i]", REGNO (x
) - FIRST_SGPR_REG
,
5278 REGNO (x
) - FIRST_SGPR_REG
+ 3);
5279 else if (VGPR_REGNO_P (REGNO (x
)))
5280 fprintf (file
, "v[%i:%i]", REGNO (x
) - FIRST_VGPR_REG
,
5281 REGNO (x
) - FIRST_VGPR_REG
+ 3);
5289 /* Implement TARGET_SECTION_TYPE_FLAGS.
5291 Return a set of section attributes for use by TARGET_ASM_NAMED_SECTION. */
5294 gcn_section_type_flags (tree decl
, const char *name
, int reloc
)
5296 if (strcmp (name
, ".lds_bss") == 0)
5297 return SECTION_WRITE
| SECTION_BSS
| SECTION_DEBUG
;
5299 return default_section_type_flags (decl
, name
, reloc
);
5302 /* Helper function for gcn_asm_output_symbol_ref.
5304 FIXME: If we want to have propagation blocks allocated separately and
5305 statically like this, it would be better done via symbol refs and the
5306 assembler/linker. This is a temporary hack. */
5309 gcn_print_lds_decl (FILE *f
, tree var
)
5312 machine_function
*machfun
= cfun
->machine
;
5314 if ((offset
= machfun
->lds_allocs
->get (var
)))
5315 fprintf (f
, "%u", (unsigned) *offset
);
5318 unsigned HOST_WIDE_INT align
= DECL_ALIGN_UNIT (var
);
5319 tree type
= TREE_TYPE (var
);
5320 unsigned HOST_WIDE_INT size
= tree_to_uhwi (TYPE_SIZE_UNIT (type
));
5321 if (size
> align
&& size
> 4 && align
< 8)
5324 machfun
->lds_allocated
= ((machfun
->lds_allocated
+ align
- 1)
5327 machfun
->lds_allocs
->put (var
, machfun
->lds_allocated
);
5328 fprintf (f
, "%u", machfun
->lds_allocated
);
5329 machfun
->lds_allocated
+= size
;
5330 if (machfun
->lds_allocated
> LDS_SIZE
)
5331 error ("local data-share memory exhausted");
5335 /* Implement ASM_OUTPUT_SYMBOL_REF via gcn-hsa.h. */
5338 gcn_asm_output_symbol_ref (FILE *file
, rtx x
)
5342 && (decl
= SYMBOL_REF_DECL (x
)) != 0
5343 && TREE_CODE (decl
) == VAR_DECL
5344 && AS_LDS_P (TYPE_ADDR_SPACE (TREE_TYPE (decl
))))
5346 /* LDS symbols (emitted using this hook) are only used at present
5347 to propagate worker values from an active thread to neutered
5348 threads. Use the same offset for each such block, but don't
5349 use zero because null pointers are used to identify the active
5350 thread in GOACC_single_copy_start calls. */
5351 gcn_print_lds_decl (file
, decl
);
5355 assemble_name (file
, XSTR (x
, 0));
5356 /* FIXME: See above -- this condition is unreachable. */
5358 && (decl
= SYMBOL_REF_DECL (x
)) != 0
5359 && TREE_CODE (decl
) == VAR_DECL
5360 && AS_LDS_P (TYPE_ADDR_SPACE (TREE_TYPE (decl
))))
5361 fputs ("@abs32", file
);
5365 /* Implement TARGET_CONSTANT_ALIGNMENT.
5367 Returns the alignment in bits of a constant that is being placed in memory.
5368 CONSTANT is the constant and BASIC_ALIGN is the alignment that the object
5369 would ordinarily have. */
5371 static HOST_WIDE_INT
5372 gcn_constant_alignment (const_tree
ARG_UNUSED (constant
),
5373 HOST_WIDE_INT basic_align
)
5375 return basic_align
> 128 ? basic_align
: 128;
5378 /* Implement PRINT_OPERAND_ADDRESS via gcn.h. */
5381 print_operand_address (FILE *file
, rtx mem
)
5383 gcc_assert (MEM_P (mem
));
5387 addr_space_t as
= MEM_ADDR_SPACE (mem
);
5388 rtx addr
= XEXP (mem
, 0);
5389 gcc_assert (REG_P (addr
) || GET_CODE (addr
) == PLUS
);
5391 if (AS_SCRATCH_P (as
))
5392 switch (GET_CODE (addr
))
5395 print_reg (file
, addr
);
5399 reg
= XEXP (addr
, 0);
5400 offset
= XEXP (addr
, 1);
5401 print_reg (file
, reg
);
5402 if (GET_CODE (offset
) == CONST_INT
)
5403 fprintf (file
, " offset:" HOST_WIDE_INT_PRINT_DEC
, INTVAL (offset
));
5412 else if (AS_ANY_FLAT_P (as
))
5414 if (GET_CODE (addr
) == REG
)
5415 print_reg (file
, addr
);
5418 gcc_assert (TARGET_GCN5_PLUS
);
5419 print_reg (file
, XEXP (addr
, 0));
5422 else if (AS_GLOBAL_P (as
))
5424 gcc_assert (TARGET_GCN5_PLUS
);
5427 rtx vgpr_offset
= NULL_RTX
;
5429 if (GET_CODE (addr
) == PLUS
)
5431 base
= XEXP (addr
, 0);
5433 if (GET_CODE (base
) == PLUS
)
5435 /* (SGPR + VGPR) + CONST */
5436 vgpr_offset
= XEXP (base
, 1);
5437 base
= XEXP (base
, 0);
5441 rtx offset
= XEXP (addr
, 1);
5445 vgpr_offset
= offset
;
5446 else if (CONST_INT_P (offset
))
5447 /* VGPR + CONST or SGPR + CONST */
5450 output_operand_lossage ("bad ADDR_SPACE_GLOBAL address");
5456 if (VGPR_REGNO_P (REGNO (base
)))
5457 print_reg (file
, base
);
5458 else if (SGPR_REGNO_P (REGNO (base
)))
5460 /* The assembler requires a 64-bit VGPR pair here, even though
5461 the offset should be only 32-bit. */
5462 if (vgpr_offset
== NULL_RTX
)
5463 /* In this case, the vector offset is zero, so we use the first
5464 lane of v1, which is initialized to zero. */
5465 fprintf (file
, "v[1:2]");
5466 else if (REG_P (vgpr_offset
)
5467 && VGPR_REGNO_P (REGNO (vgpr_offset
)))
5469 fprintf (file
, "v[%d:%d]",
5470 REGNO (vgpr_offset
) - FIRST_VGPR_REG
,
5471 REGNO (vgpr_offset
) - FIRST_VGPR_REG
+ 1);
5474 output_operand_lossage ("bad ADDR_SPACE_GLOBAL address");
5478 output_operand_lossage ("bad ADDR_SPACE_GLOBAL address");
5480 else if (AS_ANY_DS_P (as
))
5481 switch (GET_CODE (addr
))
5484 print_reg (file
, addr
);
5488 reg
= XEXP (addr
, 0);
5489 print_reg (file
, reg
);
5497 switch (GET_CODE (addr
))
5500 print_reg (file
, addr
);
5501 fprintf (file
, ", 0");
5505 reg
= XEXP (addr
, 0);
5506 offset
= XEXP (addr
, 1);
5507 print_reg (file
, reg
);
5508 fprintf (file
, ", ");
5509 if (GET_CODE (offset
) == REG
)
5510 print_reg (file
, reg
);
5511 else if (GET_CODE (offset
) == CONST_INT
)
5512 fprintf (file
, HOST_WIDE_INT_PRINT_DEC
, INTVAL (offset
));
5523 /* Implement PRINT_OPERAND via gcn.h.
5525 b - print operand size as untyped operand (b8/b16/b32/b64)
5526 B - print operand size as SI/DI untyped operand (b32/b32/b32/b64)
5527 i - print operand size as untyped operand (i16/b32/i64)
5528 I - print operand size as SI/DI untyped operand(i32/b32/i64)
5529 u - print operand size as untyped operand (u16/u32/u64)
5530 U - print operand size as SI/DI untyped operand(u32/u64)
5531 o - print operand size as memory access size for loads
5532 (ubyte/ushort/dword/dwordx2/wordx3/dwordx4)
5533 s - print operand size as memory access size for stores
5534 (byte/short/dword/dwordx2/wordx3/dwordx4)
5535 C - print conditional code for s_cbranch (_sccz/_sccnz/_vccz/_vccnz...)
5536 c - print inverse conditional code for s_cbranch
5537 D - print conditional code for s_cmp (eq_u64/lg_u64...)
5538 E - print conditional code for v_cmp (eq_u64/ne_u64...)
5539 A - print address in formatting suitable for given address space.
5540 O - print offset:n for data share operations.
5541 ^ - print "_co" suffix for GCN5 mnemonics
5542 g - print "glc", if appropriate for given MEM
5546 print_operand (FILE *file
, rtx x
, int code
)
5548 int xcode
= x
? GET_CODE (x
) : 0;
5549 bool invert
= false;
5552 /* Instructions have the following suffixes.
5553 If there are two suffixes, the first is the destination type,
5554 and the second is the source type.
5556 B32 Bitfield (untyped data) 32-bit
5557 B64 Bitfield (untyped data) 64-bit
5558 F16 floating-point 16-bit
5559 F32 floating-point 32-bit (IEEE 754 single-precision float)
5560 F64 floating-point 64-bit (IEEE 754 double-precision float)
5561 I16 signed 32-bit integer
5562 I32 signed 32-bit integer
5563 I64 signed 64-bit integer
5564 U16 unsigned 32-bit integer
5565 U32 unsigned 32-bit integer
5566 U64 unsigned 64-bit integer */
5568 /* Print operand size as untyped suffix. */
5572 machine_mode mode
= GET_MODE (x
);
5573 if (VECTOR_MODE_P (mode
))
5574 mode
= GET_MODE_INNER (mode
);
5575 switch (GET_MODE_SIZE (mode
))
5590 output_operand_lossage ("invalid operand %%xn code");
5599 machine_mode mode
= GET_MODE (x
);
5600 if (VECTOR_MODE_P (mode
))
5601 mode
= GET_MODE_INNER (mode
);
5602 switch (GET_MODE_SIZE (mode
))
5613 output_operand_lossage ("invalid operand %%xn code");
5620 fputs ("sext(", file
);
5621 print_operand (file
, x
, 0);
5629 bool signed_p
= code
== 'i';
5630 bool min32_p
= code
== 'I' || code
== 'U';
5632 machine_mode mode
= GET_MODE (x
);
5633 if (VECTOR_MODE_P (mode
))
5634 mode
= GET_MODE_INNER (mode
);
5635 if (mode
== VOIDmode
)
5636 switch (GET_CODE (x
))
5639 s
= signed_p
? "_i32" : "_u32";
5645 output_operand_lossage ("invalid operand %%xn code");
5648 else if (FLOAT_MODE_P (mode
))
5649 switch (GET_MODE_SIZE (mode
))
5661 output_operand_lossage ("invalid operand %%xn code");
5665 switch (GET_MODE_SIZE (mode
))
5670 s
= signed_p
? "_i32" : "_u32";
5673 s
= signed_p
? "_i64" : "_u64";
5676 output_operand_lossage ("invalid operand %%xn code");
5680 switch (GET_MODE_SIZE (mode
))
5683 s
= signed_p
? "_i8" : "_u8";
5686 s
= signed_p
? "_i16" : "_u16";
5689 s
= signed_p
? "_i32" : "_u32";
5692 s
= signed_p
? "_i64" : "_u64";
5695 output_operand_lossage ("invalid operand %%xn code");
5701 /* Print operand size as untyped suffix. */
5705 switch (GET_MODE_SIZE (GET_MODE (x
)))
5713 /* The following are full-vector variants. */
5728 /* Fall-through - the other cases for 'o' are the same as for 's'. */
5734 switch (GET_MODE_SIZE (GET_MODE (x
)))
5758 s
= VECTOR_MODE_P (GET_MODE (x
)) ? "_byte" : "_dwordx16";
5760 /* The following are full-vector variants. */
5771 output_operand_lossage ("invalid operand %%xn code");
5780 output_operand_lossage ("invalid %%xn code");
5783 print_operand_address (file
, x
);
5789 output_operand_lossage ("invalid %%xn code");
5792 if (AS_GDS_P (MEM_ADDR_SPACE (x
)))
5793 fprintf (file
, " gds");
5795 rtx x0
= XEXP (x
, 0);
5796 if (AS_GLOBAL_P (MEM_ADDR_SPACE (x
)))
5798 gcc_assert (TARGET_GCN5_PLUS
);
5800 fprintf (file
, ", ");
5803 rtx const_offset
= NULL_RTX
;
5805 if (GET_CODE (base
) == PLUS
)
5807 rtx offset
= XEXP (x0
, 1);
5808 base
= XEXP (x0
, 0);
5810 if (GET_CODE (base
) == PLUS
)
5811 /* (SGPR + VGPR) + CONST */
5812 /* Ignore the VGPR offset for this operand. */
5813 base
= XEXP (base
, 0);
5815 if (CONST_INT_P (offset
))
5816 const_offset
= XEXP (x0
, 1);
5817 else if (REG_P (offset
))
5819 /* Ignore the VGPR offset for this operand. */
5822 output_operand_lossage ("bad ADDR_SPACE_GLOBAL address");
5827 if (VGPR_REGNO_P (REGNO (base
)))
5828 /* The VGPR address is specified in the %A operand. */
5829 fprintf (file
, "off");
5830 else if (SGPR_REGNO_P (REGNO (base
)))
5831 print_reg (file
, base
);
5833 output_operand_lossage ("bad ADDR_SPACE_GLOBAL address");
5836 output_operand_lossage ("bad ADDR_SPACE_GLOBAL address");
5838 if (const_offset
!= NULL_RTX
)
5839 fprintf (file
, " offset:" HOST_WIDE_INT_PRINT_DEC
,
5840 INTVAL (const_offset
));
5845 if (GET_CODE (x0
) == REG
)
5847 if (GET_CODE (x0
) != PLUS
)
5849 output_operand_lossage ("invalid %%xn code");
5852 rtx val
= XEXP (x0
, 1);
5853 if (GET_CODE (val
) == CONST_VECTOR
)
5854 val
= CONST_VECTOR_ELT (val
, 0);
5855 if (GET_CODE (val
) != CONST_INT
)
5857 output_operand_lossage ("invalid %%xn code");
5860 fprintf (file
, " offset:" HOST_WIDE_INT_PRINT_DEC
, INTVAL (val
));
5871 if ((xcode
!= EQ
&& xcode
!= NE
) || !REG_P (XEXP (x
, 0)))
5873 output_operand_lossage ("invalid %%xn code");
5876 switch (REGNO (XEXP (x
, 0)))
5883 /* For some reason llvm-mc insists on scc0 instead of sccz. */
5891 output_operand_lossage ("invalid %%xn code");
5895 if (xcode
== (invert
? NE
: EQ
))
5896 fputc (num
? '0' : 'z', file
);
5898 fputs (num
? "1" : "nz", file
);
5904 bool cmp_signed
= false;
5942 output_operand_lossage ("invalid %%xn code");
5946 fputc (cmp_signed
? 'i' : 'u', file
);
5948 machine_mode mode
= GET_MODE (XEXP (x
, 0));
5950 if (mode
== VOIDmode
)
5951 mode
= GET_MODE (XEXP (x
, 1));
5953 /* If both sides are constants, then assume the instruction is in
5954 SImode since s_cmp can only do integer compares. */
5955 if (mode
== VOIDmode
)
5958 switch (GET_MODE_SIZE (mode
))
5967 output_operand_lossage ("invalid operand %%xn code");
5976 bool cmp_signed
= false;
5977 machine_mode mode
= GET_MODE (XEXP (x
, 0));
5979 if (mode
== VOIDmode
)
5980 mode
= GET_MODE (XEXP (x
, 1));
5982 /* If both sides are constants, assume the instruction is in SFmode
5983 if either operand is floating point, otherwise assume SImode. */
5984 if (mode
== VOIDmode
)
5986 if (GET_CODE (XEXP (x
, 0)) == CONST_DOUBLE
5987 || GET_CODE (XEXP (x
, 1)) == CONST_DOUBLE
)
5993 /* Use the same format code for vector comparisons. */
5994 if (GET_MODE_CLASS (mode
) == MODE_VECTOR_FLOAT
5995 || GET_MODE_CLASS (mode
) == MODE_VECTOR_INT
)
5996 mode
= GET_MODE_INNER (mode
);
5998 bool float_p
= GET_MODE_CLASS (mode
) == MODE_FLOAT
;
6006 s
= float_p
? "_neq_" : "_ne_";
6061 output_operand_lossage ("invalid %%xn code");
6065 fputc (float_p
? 'f' : cmp_signed
? 'i' : 'u', file
);
6067 switch (GET_MODE_SIZE (mode
))
6070 output_operand_lossage ("operand %%xn code invalid for QImode");
6082 output_operand_lossage ("invalid operand %%xn code");
6089 print_operand (file
, gcn_operand_part (GET_MODE (x
), x
, 0), 0);
6092 print_operand (file
, gcn_operand_part (GET_MODE (x
), x
, 1), 0);
6095 /* Print a scalar register number as an integer. Temporary hack. */
6096 gcc_assert (REG_P (x
));
6097 fprintf (file
, "%u", (int) REGNO (x
));
6100 /* Print a vector register number as an integer. Temporary hack. */
6101 gcc_assert (REG_P (x
));
6102 fprintf (file
, "%u", (int) REGNO (x
) - FIRST_VGPR_REG
);
6106 print_reg (file
, x
);
6107 else if (xcode
== MEM
)
6108 output_address (GET_MODE (x
), x
);
6109 else if (xcode
== CONST_INT
)
6110 fprintf (file
, "%i", (int) INTVAL (x
));
6111 else if (xcode
== CONST_VECTOR
)
6112 print_operand (file
, CONST_VECTOR_ELT (x
, 0), code
);
6113 else if (xcode
== CONST_DOUBLE
)
6116 switch (gcn_inline_fp_constant_p (x
, false))
6146 rtx ix
= simplify_gen_subreg (GET_MODE (x
) == DFmode
6148 x
, GET_MODE (x
), 0);
6150 print_operand (file
, ix
, code
);
6152 output_operand_lossage ("invalid fp constant");
6156 fprintf (file
, str
);
6160 output_addr_const (file
, x
);
6163 if (TARGET_GCN5_PLUS
)
6164 fputs ("_co", file
);
6167 gcc_assert (xcode
== MEM
);
6168 if (MEM_VOLATILE_P (x
))
6169 fputs (" glc", file
);
6172 output_operand_lossage ("invalid %%xn code");
6177 /* Implement DBX_REGISTER_NUMBER macro.
6179 Return the DWARF register number that corresponds to the GCC internal
6183 gcn_dwarf_register_number (unsigned int regno
)
6185 /* Registers defined in DWARF. */
6186 if (regno
== EXEC_LO_REG
)
6188 /* We need to use a more complex DWARF expression for this
6189 else if (regno == EXEC_HI_REG)
6191 else if (regno
== VCC_LO_REG
)
6193 /* We need to use a more complex DWARF expression for this
6194 else if (regno == VCC_HI_REG)
6196 else if (regno
== SCC_REG
)
6198 else if (SGPR_REGNO_P (regno
))
6200 if (regno
- FIRST_SGPR_REG
< 64)
6201 return (regno
- FIRST_SGPR_REG
+ 32);
6203 return (regno
- FIRST_SGPR_REG
+ 1024);
6205 else if (VGPR_REGNO_P (regno
))
6206 return (regno
- FIRST_VGPR_REG
+ 2560);
6208 /* Otherwise, there's nothing sensible to do. */
6209 return regno
+ 100000;
6212 /* Implement TARGET_DWARF_REGISTER_SPAN.
6214 DImode and Vector DImode require additional registers. */
6217 gcn_dwarf_register_span (rtx rtl
)
6219 machine_mode mode
= GET_MODE (rtl
);
6221 if (VECTOR_MODE_P (mode
))
6222 mode
= GET_MODE_INNER (mode
);
6224 if (GET_MODE_SIZE (mode
) != 8)
6227 rtx p
= gen_rtx_PARALLEL (VOIDmode
, rtvec_alloc (2));
6228 unsigned regno
= REGNO (rtl
);
6229 XVECEXP (p
, 0, 0) = gen_rtx_REG (SImode
, regno
);
6230 XVECEXP (p
, 0, 1) = gen_rtx_REG (SImode
, regno
+ 1);
6236 /* {{{ TARGET hook overrides. */
6238 #undef TARGET_ADDR_SPACE_ADDRESS_MODE
6239 #define TARGET_ADDR_SPACE_ADDRESS_MODE gcn_addr_space_address_mode
6240 #undef TARGET_ADDR_SPACE_LEGITIMATE_ADDRESS_P
6241 #define TARGET_ADDR_SPACE_LEGITIMATE_ADDRESS_P \
6242 gcn_addr_space_legitimate_address_p
6243 #undef TARGET_ADDR_SPACE_LEGITIMIZE_ADDRESS
6244 #define TARGET_ADDR_SPACE_LEGITIMIZE_ADDRESS gcn_addr_space_legitimize_address
6245 #undef TARGET_ADDR_SPACE_POINTER_MODE
6246 #define TARGET_ADDR_SPACE_POINTER_MODE gcn_addr_space_pointer_mode
6247 #undef TARGET_ADDR_SPACE_SUBSET_P
6248 #define TARGET_ADDR_SPACE_SUBSET_P gcn_addr_space_subset_p
6249 #undef TARGET_ADDR_SPACE_CONVERT
6250 #define TARGET_ADDR_SPACE_CONVERT gcn_addr_space_convert
6251 #undef TARGET_ARG_PARTIAL_BYTES
6252 #define TARGET_ARG_PARTIAL_BYTES gcn_arg_partial_bytes
6253 #undef TARGET_ASM_ALIGNED_DI_OP
6254 #define TARGET_ASM_ALIGNED_DI_OP "\t.8byte\t"
6255 #undef TARGET_ASM_FILE_START
6256 #define TARGET_ASM_FILE_START output_file_start
6257 #undef TARGET_ASM_FUNCTION_PROLOGUE
6258 #define TARGET_ASM_FUNCTION_PROLOGUE gcn_target_asm_function_prologue
6259 #undef TARGET_ASM_SELECT_SECTION
6260 #define TARGET_ASM_SELECT_SECTION gcn_asm_select_section
6261 #undef TARGET_ASM_TRAMPOLINE_TEMPLATE
6262 #define TARGET_ASM_TRAMPOLINE_TEMPLATE gcn_asm_trampoline_template
6263 #undef TARGET_ATTRIBUTE_TABLE
6264 #define TARGET_ATTRIBUTE_TABLE gcn_attribute_table
6265 #undef TARGET_BUILTIN_DECL
6266 #define TARGET_BUILTIN_DECL gcn_builtin_decl
6267 #undef TARGET_CAN_CHANGE_MODE_CLASS
6268 #define TARGET_CAN_CHANGE_MODE_CLASS gcn_can_change_mode_class
6269 #undef TARGET_CAN_ELIMINATE
6270 #define TARGET_CAN_ELIMINATE gcn_can_eliminate_p
6271 #undef TARGET_CANNOT_COPY_INSN_P
6272 #define TARGET_CANNOT_COPY_INSN_P gcn_cannot_copy_insn_p
6273 #undef TARGET_CLASS_LIKELY_SPILLED_P
6274 #define TARGET_CLASS_LIKELY_SPILLED_P gcn_class_likely_spilled_p
6275 #undef TARGET_CLASS_MAX_NREGS
6276 #define TARGET_CLASS_MAX_NREGS gcn_class_max_nregs
6277 #undef TARGET_CONDITIONAL_REGISTER_USAGE
6278 #define TARGET_CONDITIONAL_REGISTER_USAGE gcn_conditional_register_usage
6279 #undef TARGET_CONSTANT_ALIGNMENT
6280 #define TARGET_CONSTANT_ALIGNMENT gcn_constant_alignment
6281 #undef TARGET_DEBUG_UNWIND_INFO
6282 #define TARGET_DEBUG_UNWIND_INFO gcn_debug_unwind_info
6283 #undef TARGET_DWARF_REGISTER_SPAN
6284 #define TARGET_DWARF_REGISTER_SPAN gcn_dwarf_register_span
6285 #undef TARGET_EMUTLS_VAR_INIT
6286 #define TARGET_EMUTLS_VAR_INIT gcn_emutls_var_init
6287 #undef TARGET_EXPAND_BUILTIN
6288 #define TARGET_EXPAND_BUILTIN gcn_expand_builtin
6289 #undef TARGET_FUNCTION_ARG
6290 #undef TARGET_FUNCTION_ARG_ADVANCE
6291 #define TARGET_FUNCTION_ARG_ADVANCE gcn_function_arg_advance
6292 #define TARGET_FUNCTION_ARG gcn_function_arg
6293 #undef TARGET_FUNCTION_VALUE
6294 #define TARGET_FUNCTION_VALUE gcn_function_value
6295 #undef TARGET_FUNCTION_VALUE_REGNO_P
6296 #define TARGET_FUNCTION_VALUE_REGNO_P gcn_function_value_regno_p
6297 #undef TARGET_GIMPLIFY_VA_ARG_EXPR
6298 #define TARGET_GIMPLIFY_VA_ARG_EXPR gcn_gimplify_va_arg_expr
6299 #undef TARGET_OMP_DEVICE_KIND_ARCH_ISA
6300 #define TARGET_OMP_DEVICE_KIND_ARCH_ISA gcn_omp_device_kind_arch_isa
6301 #undef TARGET_GOACC_ADJUST_PROPAGATION_RECORD
6302 #define TARGET_GOACC_ADJUST_PROPAGATION_RECORD \
6303 gcn_goacc_adjust_propagation_record
6304 #undef TARGET_GOACC_ADJUST_GANGPRIVATE_DECL
6305 #define TARGET_GOACC_ADJUST_GANGPRIVATE_DECL gcn_goacc_adjust_gangprivate_decl
6306 #undef TARGET_GOACC_FORK_JOIN
6307 #define TARGET_GOACC_FORK_JOIN gcn_fork_join
6308 #undef TARGET_GOACC_REDUCTION
6309 #define TARGET_GOACC_REDUCTION gcn_goacc_reduction
6310 #undef TARGET_GOACC_VALIDATE_DIMS
6311 #define TARGET_GOACC_VALIDATE_DIMS gcn_goacc_validate_dims
6312 #undef TARGET_HARD_REGNO_MODE_OK
6313 #define TARGET_HARD_REGNO_MODE_OK gcn_hard_regno_mode_ok
6314 #undef TARGET_HARD_REGNO_NREGS
6315 #define TARGET_HARD_REGNO_NREGS gcn_hard_regno_nregs
6316 #undef TARGET_HAVE_SPECULATION_SAFE_VALUE
6317 #define TARGET_HAVE_SPECULATION_SAFE_VALUE speculation_safe_value_not_needed
6318 #undef TARGET_INIT_BUILTINS
6319 #define TARGET_INIT_BUILTINS gcn_init_builtins
6320 #undef TARGET_IRA_CHANGE_PSEUDO_ALLOCNO_CLASS
6321 #define TARGET_IRA_CHANGE_PSEUDO_ALLOCNO_CLASS \
6322 gcn_ira_change_pseudo_allocno_class
6323 #undef TARGET_LEGITIMATE_CONSTANT_P
6324 #define TARGET_LEGITIMATE_CONSTANT_P gcn_legitimate_constant_p
6326 #define TARGET_LRA_P hook_bool_void_true
6327 #undef TARGET_MACHINE_DEPENDENT_REORG
6328 #define TARGET_MACHINE_DEPENDENT_REORG gcn_md_reorg
6329 #undef TARGET_MEMORY_MOVE_COST
6330 #define TARGET_MEMORY_MOVE_COST gcn_memory_move_cost
6331 #undef TARGET_MODES_TIEABLE_P
6332 #define TARGET_MODES_TIEABLE_P gcn_modes_tieable_p
6333 #undef TARGET_OPTION_OVERRIDE
6334 #define TARGET_OPTION_OVERRIDE gcn_option_override
6335 #undef TARGET_PRETEND_OUTGOING_VARARGS_NAMED
6336 #define TARGET_PRETEND_OUTGOING_VARARGS_NAMED \
6337 gcn_pretend_outgoing_varargs_named
6338 #undef TARGET_PROMOTE_FUNCTION_MODE
6339 #define TARGET_PROMOTE_FUNCTION_MODE gcn_promote_function_mode
6340 #undef TARGET_REGISTER_MOVE_COST
6341 #define TARGET_REGISTER_MOVE_COST gcn_register_move_cost
6342 #undef TARGET_RETURN_IN_MEMORY
6343 #define TARGET_RETURN_IN_MEMORY gcn_return_in_memory
6344 #undef TARGET_RTX_COSTS
6345 #define TARGET_RTX_COSTS gcn_rtx_costs
6346 #undef TARGET_SECONDARY_RELOAD
6347 #define TARGET_SECONDARY_RELOAD gcn_secondary_reload
6348 #undef TARGET_SECTION_TYPE_FLAGS
6349 #define TARGET_SECTION_TYPE_FLAGS gcn_section_type_flags
6350 #undef TARGET_SCALAR_MODE_SUPPORTED_P
6351 #define TARGET_SCALAR_MODE_SUPPORTED_P gcn_scalar_mode_supported_p
6352 #undef TARGET_SMALL_REGISTER_CLASSES_FOR_MODE_P
6353 #define TARGET_SMALL_REGISTER_CLASSES_FOR_MODE_P \
6354 gcn_small_register_classes_for_mode_p
6355 #undef TARGET_SPILL_CLASS
6356 #define TARGET_SPILL_CLASS gcn_spill_class
6357 #undef TARGET_STRICT_ARGUMENT_NAMING
6358 #define TARGET_STRICT_ARGUMENT_NAMING gcn_strict_argument_naming
6359 #undef TARGET_TRAMPOLINE_INIT
6360 #define TARGET_TRAMPOLINE_INIT gcn_trampoline_init
6361 #undef TARGET_TRULY_NOOP_TRUNCATION
6362 #define TARGET_TRULY_NOOP_TRUNCATION gcn_truly_noop_truncation
6363 #undef TARGET_VECTORIZE_BUILTIN_VECTORIZATION_COST
6364 #define TARGET_VECTORIZE_BUILTIN_VECTORIZATION_COST gcn_vectorization_cost
6365 #undef TARGET_VECTORIZE_GET_MASK_MODE
6366 #define TARGET_VECTORIZE_GET_MASK_MODE gcn_vectorize_get_mask_mode
6367 #undef TARGET_VECTORIZE_PREFERRED_SIMD_MODE
6368 #define TARGET_VECTORIZE_PREFERRED_SIMD_MODE gcn_vectorize_preferred_simd_mode
6369 #undef TARGET_VECTORIZE_PREFERRED_VECTOR_ALIGNMENT
6370 #define TARGET_VECTORIZE_PREFERRED_VECTOR_ALIGNMENT \
6371 gcn_preferred_vector_alignment
6372 #undef TARGET_VECTORIZE_RELATED_MODE
6373 #define TARGET_VECTORIZE_RELATED_MODE gcn_related_vector_mode
6374 #undef TARGET_VECTORIZE_SUPPORT_VECTOR_MISALIGNMENT
6375 #define TARGET_VECTORIZE_SUPPORT_VECTOR_MISALIGNMENT \
6376 gcn_vectorize_support_vector_misalignment
6377 #undef TARGET_VECTORIZE_VEC_PERM_CONST
6378 #define TARGET_VECTORIZE_VEC_PERM_CONST gcn_vectorize_vec_perm_const
6379 #undef TARGET_VECTORIZE_VECTOR_ALIGNMENT_REACHABLE
6380 #define TARGET_VECTORIZE_VECTOR_ALIGNMENT_REACHABLE \
6381 gcn_vector_alignment_reachable
6382 #undef TARGET_VECTOR_MODE_SUPPORTED_P
6383 #define TARGET_VECTOR_MODE_SUPPORTED_P gcn_vector_mode_supported_p
6385 struct gcc_target targetm
= TARGET_INITIALIZER
;