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 %qs in %<amdgpu_hsa_kernel%> attribute",
249 if (args
->requested
& (1 << a
))
251 error ("duplicated parameter specifier %qs in %<amdgpu_hsa_kernel%> "
256 args
->requested
|= (1 << a
);
257 args
->order
[args
->nargs
++] = a
;
260 /* Requesting WORK_ITEM_ID_Z_ARG implies requesting WORK_ITEM_ID_X_ARG and
261 WORK_ITEM_ID_Y_ARG. Similarly, requesting WORK_ITEM_ID_Y_ARG implies
262 requesting WORK_ITEM_ID_X_ARG. */
263 if (args
->requested
& (1 << WORK_ITEM_ID_Z_ARG
))
264 args
->requested
|= (1 << WORK_ITEM_ID_Y_ARG
);
265 if (args
->requested
& (1 << WORK_ITEM_ID_Y_ARG
))
266 args
->requested
|= (1 << WORK_ITEM_ID_X_ARG
);
268 int sgpr_regno
= FIRST_SGPR_REG
;
270 for (int a
= 0; a
< GCN_KERNEL_ARG_TYPES
; a
++)
272 if (!(args
->requested
& (1 << a
)))
275 if (gcn_kernel_arg_types
[a
].fixed_regno
>= 0)
276 args
->reg
[a
] = gcn_kernel_arg_types
[a
].fixed_regno
;
281 switch (gcn_kernel_arg_types
[a
].mode
)
295 args
->reg
[a
] = sgpr_regno
;
296 sgpr_regno
+= reg_count
;
297 if (gcn_kernel_arg_types
[a
].fixed_regno
== -1)
298 args
->nsgprs
+= reg_count
;
301 if (sgpr_regno
> FIRST_SGPR_REG
+ 16)
303 error ("too many arguments passed in sgpr registers");
308 /* Referenced by TARGET_ATTRIBUTE_TABLE.
310 Validates target specific attributes. */
313 gcn_handle_amdgpu_hsa_kernel_attribute (tree
*node
, tree name
,
314 tree args
, int, bool *no_add_attrs
)
316 if (!FUNC_OR_METHOD_TYPE_P (*node
))
318 warning (OPT_Wattributes
, "%qE attribute only applies to functions",
320 *no_add_attrs
= true;
324 /* Can combine regparm with all attributes but fastcall, and thiscall. */
325 if (is_attribute_p ("gcnhsa_kernel", name
))
327 struct gcn_kernel_args kernelarg
;
329 if (gcn_parse_amdgpu_hsa_kernel_attribute (&kernelarg
, args
))
330 *no_add_attrs
= true;
338 /* Implement TARGET_ATTRIBUTE_TABLE.
340 Create target-specific __attribute__ types. */
342 static const struct attribute_spec gcn_attribute_table
[] = {
343 /* { name, min_len, max_len, decl_req, type_req, fn_type_req, handler,
344 affects_type_identity } */
345 {"amdgpu_hsa_kernel", 0, GCN_KERNEL_ARG_TYPES
, false, true,
346 true, true, gcn_handle_amdgpu_hsa_kernel_attribute
, NULL
},
348 {NULL
, 0, 0, false, false, false, false, NULL
, NULL
}
352 /* {{{ Registers and modes. */
354 /* Implement TARGET_SCALAR_MODE_SUPPORTED_P. */
357 gcn_scalar_mode_supported_p (scalar_mode mode
)
359 return (mode
== BImode
361 || mode
== HImode
/* || mode == HFmode */
362 || mode
== SImode
|| mode
== SFmode
363 || mode
== DImode
|| mode
== DFmode
367 /* Implement TARGET_CLASS_MAX_NREGS.
369 Return the number of hard registers needed to hold a value of MODE in
370 a register of class RCLASS. */
373 gcn_class_max_nregs (reg_class_t rclass
, machine_mode mode
)
375 /* Scalar registers are 32bit, vector registers are in fact tuples of
377 if (rclass
== VGPR_REGS
)
379 if (vgpr_1reg_mode_p (mode
))
381 if (vgpr_2reg_mode_p (mode
))
383 /* TImode is used by DImode compare_and_swap. */
387 else if (rclass
== VCC_CONDITIONAL_REG
&& mode
== BImode
)
389 return CEIL (GET_MODE_SIZE (mode
), 4);
392 /* Implement TARGET_HARD_REGNO_NREGS.
394 Return the number of hard registers needed to hold a value of MODE in
398 gcn_hard_regno_nregs (unsigned int regno
, machine_mode mode
)
400 return gcn_class_max_nregs (REGNO_REG_CLASS (regno
), mode
);
403 /* Implement TARGET_HARD_REGNO_MODE_OK.
405 Return true if REGNO can hold value in MODE. */
408 gcn_hard_regno_mode_ok (unsigned int regno
, machine_mode mode
)
410 /* Treat a complex mode as if it were a scalar mode of the same overall
411 size for the purposes of allocating hard registers. */
412 if (COMPLEX_MODE_P (mode
))
438 case FLAT_SCRATCH_LO_REG
:
439 case XNACK_MASK_LO_REG
:
442 return (mode
== SImode
|| mode
== DImode
);
445 return (mode
== BImode
|| mode
== SImode
|| mode
== DImode
);
447 case FLAT_SCRATCH_HI_REG
:
448 case XNACK_MASK_HI_REG
:
451 return mode
== SImode
;
455 return mode
== SImode
/*|| mode == V32BImode */ ;
459 return mode
== BImode
;
461 if (regno
== ARG_POINTER_REGNUM
|| regno
== FRAME_POINTER_REGNUM
)
463 if (SGPR_REGNO_P (regno
))
464 /* We restrict double register values to aligned registers. */
465 return (sgpr_1reg_mode_p (mode
)
466 || (!((regno
- FIRST_SGPR_REG
) & 1) && sgpr_2reg_mode_p (mode
))
467 || (((regno
- FIRST_SGPR_REG
) & 3) == 0 && mode
== TImode
));
468 if (VGPR_REGNO_P (regno
))
469 /* Vector instructions do not care about the alignment of register
470 pairs, but where there is no 64-bit instruction, many of the
471 define_split do not work if the input and output registers partially
472 overlap. We tried to fix this with early clobber and match
473 constraints, but it was bug prone, added complexity, and conflicts
474 with the 'U0' constraints on vec_merge.
475 Therefore, we restrict ourselved to aligned registers. */
476 return (vgpr_1reg_mode_p (mode
)
477 || (!((regno
- FIRST_VGPR_REG
) & 1) && vgpr_2reg_mode_p (mode
))
478 /* TImode is used by DImode compare_and_swap. */
480 && !((regno
- FIRST_VGPR_REG
) & 3)));
484 /* Implement REGNO_REG_CLASS via gcn.h.
486 Return smallest class containing REGNO. */
489 gcn_regno_reg_class (int regno
)
494 return SCC_CONDITIONAL_REG
;
497 return VCC_CONDITIONAL_REG
;
499 return VCCZ_CONDITIONAL_REG
;
501 return EXECZ_CONDITIONAL_REG
;
504 return EXEC_MASK_REG
;
506 if (VGPR_REGNO_P (regno
))
508 if (SGPR_REGNO_P (regno
))
510 if (regno
< FIRST_VGPR_REG
)
512 if (regno
== ARG_POINTER_REGNUM
|| regno
== FRAME_POINTER_REGNUM
)
517 /* Implement TARGET_CAN_CHANGE_MODE_CLASS.
519 GCC assumes that lowpart contains first part of value as stored in memory.
520 This is not the case for vector registers. */
523 gcn_can_change_mode_class (machine_mode from
, machine_mode to
,
524 reg_class_t regclass
)
526 if (!vgpr_vector_mode_p (from
) && !vgpr_vector_mode_p (to
))
528 return (gcn_class_max_nregs (regclass
, from
)
529 == gcn_class_max_nregs (regclass
, to
));
532 /* Implement TARGET_SMALL_REGISTER_CLASSES_FOR_MODE_P.
534 When this hook returns true for MODE, the compiler allows
535 registers explicitly used in the rtl to be used as spill registers
536 but prevents the compiler from extending the lifetime of these
540 gcn_small_register_classes_for_mode_p (machine_mode mode
)
542 /* We allocate into exec and vcc regs. Those make small register class. */
543 return mode
== DImode
|| mode
== SImode
;
546 /* Implement TARGET_CLASS_LIKELY_SPILLED_P.
548 Returns true if pseudos that have been assigned to registers of class RCLASS
549 would likely be spilled because registers of RCLASS are needed for spill
553 gcn_class_likely_spilled_p (reg_class_t rclass
)
555 return (rclass
== EXEC_MASK_REG
556 || reg_classes_intersect_p (ALL_CONDITIONAL_REGS
, rclass
));
559 /* Implement TARGET_MODES_TIEABLE_P.
561 Returns true if a value of MODE1 is accessible in MODE2 without
565 gcn_modes_tieable_p (machine_mode mode1
, machine_mode mode2
)
567 return (GET_MODE_BITSIZE (mode1
) <= MAX_FIXED_MODE_SIZE
568 && GET_MODE_BITSIZE (mode2
) <= MAX_FIXED_MODE_SIZE
);
571 /* Implement TARGET_TRULY_NOOP_TRUNCATION.
573 Returns true if it is safe to “convert” a value of INPREC bits to one of
574 OUTPREC bits (where OUTPREC is smaller than INPREC) by merely operating on
575 it as if it had only OUTPREC bits. */
578 gcn_truly_noop_truncation (poly_uint64 outprec
, poly_uint64 inprec
)
580 return ((inprec
<= 32) && (outprec
<= inprec
));
583 /* Return N-th part of value occupying multiple registers. */
586 gcn_operand_part (machine_mode mode
, rtx op
, int n
)
588 if (GET_MODE_SIZE (mode
) >= 256)
590 /*gcc_assert (GET_MODE_SIZE (mode) == 256 || n == 0); */
594 gcc_assert (REGNO (op
) + n
< FIRST_PSEUDO_REGISTER
);
595 return gen_rtx_REG (V64SImode
, REGNO (op
) + n
);
597 if (GET_CODE (op
) == CONST_VECTOR
)
599 int units
= GET_MODE_NUNITS (mode
);
600 rtvec v
= rtvec_alloc (units
);
602 for (int i
= 0; i
< units
; ++i
)
603 RTVEC_ELT (v
, i
) = gcn_operand_part (GET_MODE_INNER (mode
),
604 CONST_VECTOR_ELT (op
, i
), n
);
606 return gen_rtx_CONST_VECTOR (V64SImode
, v
);
608 if (GET_CODE (op
) == UNSPEC
&& XINT (op
, 1) == UNSPEC_VECTOR
)
609 return gcn_gen_undef (V64SImode
);
612 else if (GET_MODE_SIZE (mode
) == 8 && REG_P (op
))
614 gcc_assert (REGNO (op
) + n
< FIRST_PSEUDO_REGISTER
);
615 return gen_rtx_REG (SImode
, REGNO (op
) + n
);
619 if (GET_CODE (op
) == UNSPEC
&& XINT (op
, 1) == UNSPEC_VECTOR
)
620 return gcn_gen_undef (SImode
);
622 /* If it's a constant then let's assume it is of the largest mode
623 available, otherwise simplify_gen_subreg will fail. */
624 if (mode
== VOIDmode
&& CONST_INT_P (op
))
626 return simplify_gen_subreg (SImode
, op
, mode
, n
* 4);
630 /* Return N-th part of value occupying multiple registers. */
633 gcn_operand_doublepart (machine_mode mode
, rtx op
, int n
)
635 return simplify_gen_subreg (DImode
, op
, mode
, n
* 8);
638 /* Return true if OP can be split into subregs or high/low parts.
639 This is always true for scalars, but not normally true for vectors.
640 However, for vectors in hardregs we can use the low and high registers. */
643 gcn_can_split_p (machine_mode
, rtx op
)
645 if (vgpr_vector_mode_p (GET_MODE (op
)))
647 if (GET_CODE (op
) == SUBREG
)
648 op
= SUBREG_REG (op
);
651 return REGNO (op
) <= FIRST_PSEUDO_REGISTER
;
656 /* Implement TARGET_SPILL_CLASS.
658 Return class of registers which could be used for pseudo of MODE
659 and of class RCLASS for spilling instead of memory. Return NO_REGS
660 if it is not possible or non-profitable. */
663 gcn_spill_class (reg_class_t c
, machine_mode
/*mode */ )
665 if (reg_classes_intersect_p (ALL_CONDITIONAL_REGS
, c
)
666 || c
== VCC_CONDITIONAL_REG
)
672 /* Implement TARGET_IRA_CHANGE_PSEUDO_ALLOCNO_CLASS.
674 Change allocno class for given pseudo from allocno and best class
675 calculated by IRA. */
678 gcn_ira_change_pseudo_allocno_class (int regno
, reg_class_t cl
,
681 /* Avoid returning classes that contain both vgpr and sgpr registers. */
682 if (cl
!= ALL_REGS
&& cl
!= SRCDST_REGS
&& cl
!= ALL_GPR_REGS
)
684 if (best_cl
!= ALL_REGS
&& best_cl
!= SRCDST_REGS
685 && best_cl
!= ALL_GPR_REGS
)
688 machine_mode mode
= PSEUDO_REGNO_MODE (regno
);
689 if (vgpr_vector_mode_p (mode
))
695 /* Create a new DImode pseudo reg and emit an instruction to initialize
699 get_exec (int64_t val
)
701 rtx reg
= gen_reg_rtx (DImode
);
702 emit_insn (gen_rtx_SET (reg
, gen_int_mode (val
, DImode
)));
706 /* Return value of scalar exec register. */
714 /* Return pseudo holding scalar exec register. */
717 gcn_scalar_exec_reg ()
722 /* Return value of full exec register. */
730 /* Return pseudo holding full exec register. */
735 return get_exec (-1);
739 /* {{{ Immediate constants. */
741 /* Initialize shared numeric constants. */
744 init_ext_gcn_constants (void)
746 real_from_integer (&dconst4
, DFmode
, 4, SIGNED
);
748 /* FIXME: this constant probably does not match what hardware really loads.
749 Reality check it eventually. */
750 real_from_string (&dconst1over2pi
,
751 "0.1591549430918953357663423455968866839");
752 real_convert (&dconst1over2pi
, SFmode
, &dconst1over2pi
);
754 ext_gcn_constants_init
= 1;
757 /* Return non-zero if X is a constant that can appear as an inline operand.
758 This is 0, 0.5, -0.5, 1, -1, 2, -2, 4,-4, 1/(2*pi)
759 Or a vector of those.
760 The value returned should be the encoding of this constant. */
763 gcn_inline_fp_constant_p (rtx x
, bool allow_vector
)
765 machine_mode mode
= GET_MODE (x
);
767 if ((mode
== V64HFmode
|| mode
== V64SFmode
|| mode
== V64DFmode
)
771 if (GET_CODE (x
) != CONST_VECTOR
)
773 n
= gcn_inline_fp_constant_p (CONST_VECTOR_ELT (x
, 0), false);
776 for (int i
= 1; i
< 64; i
++)
777 if (CONST_VECTOR_ELT (x
, i
) != CONST_VECTOR_ELT (x
, 0))
782 if (mode
!= HFmode
&& mode
!= SFmode
&& mode
!= DFmode
)
785 const REAL_VALUE_TYPE
*r
;
787 if (x
== CONST0_RTX (mode
))
789 if (x
== CONST1_RTX (mode
))
792 r
= CONST_DOUBLE_REAL_VALUE (x
);
794 if (real_identical (r
, &dconstm1
))
797 if (real_identical (r
, &dconsthalf
))
799 if (real_identical (r
, &dconstm1
))
801 if (real_identical (r
, &dconst2
))
803 if (real_identical (r
, &dconst4
))
805 if (real_identical (r
, &dconst1over2pi
))
807 if (!ext_gcn_constants_init
)
808 init_ext_gcn_constants ();
809 real_value_negate (r
);
810 if (real_identical (r
, &dconsthalf
))
812 if (real_identical (r
, &dconst2
))
814 if (real_identical (r
, &dconst4
))
817 /* FIXME: add 4, -4 and 1/(2*PI). */
822 /* Return non-zero if X is a constant that can appear as an immediate operand.
823 This is 0, 0.5, -0.5, 1, -1, 2, -2, 4,-4, 1/(2*pi)
824 Or a vector of those.
825 The value returned should be the encoding of this constant. */
828 gcn_fp_constant_p (rtx x
, bool allow_vector
)
830 machine_mode mode
= GET_MODE (x
);
832 if ((mode
== V64HFmode
|| mode
== V64SFmode
|| mode
== V64DFmode
)
836 if (GET_CODE (x
) != CONST_VECTOR
)
838 n
= gcn_fp_constant_p (CONST_VECTOR_ELT (x
, 0), false);
841 for (int i
= 1; i
< 64; i
++)
842 if (CONST_VECTOR_ELT (x
, i
) != CONST_VECTOR_ELT (x
, 0))
846 if (mode
!= HFmode
&& mode
!= SFmode
&& mode
!= DFmode
)
849 if (gcn_inline_fp_constant_p (x
, false))
851 /* FIXME: It is not clear how 32bit immediates are interpreted here. */
852 return (mode
!= DFmode
);
855 /* Return true if X is a constant representable as an inline immediate
856 constant in a 32-bit instruction encoding. */
859 gcn_inline_constant_p (rtx x
)
861 if (GET_CODE (x
) == CONST_INT
)
862 return INTVAL (x
) >= -16 && INTVAL (x
) <= 64;
863 if (GET_CODE (x
) == CONST_DOUBLE
)
864 return gcn_inline_fp_constant_p (x
, false);
865 if (GET_CODE (x
) == CONST_VECTOR
)
868 if (!vgpr_vector_mode_p (GET_MODE (x
)))
870 n
= gcn_inline_constant_p (CONST_VECTOR_ELT (x
, 0));
873 for (int i
= 1; i
< 64; i
++)
874 if (CONST_VECTOR_ELT (x
, i
) != CONST_VECTOR_ELT (x
, 0))
881 /* Return true if X is a constant representable as an immediate constant
882 in a 32 or 64-bit instruction encoding. */
885 gcn_constant_p (rtx x
)
887 switch (GET_CODE (x
))
893 return gcn_fp_constant_p (x
, false);
898 if (!vgpr_vector_mode_p (GET_MODE (x
)))
900 n
= gcn_constant_p (CONST_VECTOR_ELT (x
, 0));
903 for (int i
= 1; i
< 64; i
++)
904 if (CONST_VECTOR_ELT (x
, i
) != CONST_VECTOR_ELT (x
, 0))
920 /* Return true if X is a constant representable as two inline immediate
921 constants in a 64-bit instruction that is split into two 32-bit
923 When MIXED is set, the low-part is permitted to use the full 32-bits. */
926 gcn_inline_constant64_p (rtx x
, bool mixed
)
928 if (GET_CODE (x
) == CONST_VECTOR
)
930 if (!vgpr_vector_mode_p (GET_MODE (x
)))
932 if (!gcn_inline_constant64_p (CONST_VECTOR_ELT (x
, 0), mixed
))
934 for (int i
= 1; i
< 64; i
++)
935 if (CONST_VECTOR_ELT (x
, i
) != CONST_VECTOR_ELT (x
, 0))
941 if (GET_CODE (x
) != CONST_INT
)
944 rtx val_lo
= gcn_operand_part (DImode
, x
, 0);
945 rtx val_hi
= gcn_operand_part (DImode
, x
, 1);
946 return ((mixed
|| gcn_inline_constant_p (val_lo
))
947 && gcn_inline_constant_p (val_hi
));
950 /* Return true if X is a constant representable as an immediate constant
951 in a 32 or 64-bit instruction encoding where the hardware will
952 extend the immediate to 64-bits. */
955 gcn_constant64_p (rtx x
)
957 if (!gcn_constant_p (x
))
960 if (GET_CODE (x
) != CONST_INT
)
963 /* Negative numbers are only allowed if they can be encoded within src0,
964 because the 32-bit immediates do not get sign-extended.
965 Unsigned numbers must not be encodable as 32-bit -1..-16, because the
966 assembler will use a src0 inline immediate and that will get
968 HOST_WIDE_INT val
= INTVAL (x
);
969 return (((val
& 0xffffffff) == val
/* Positive 32-bit. */
970 && (val
& 0xfffffff0) != 0xfffffff0) /* Not -1..-16. */
971 || gcn_inline_constant_p (x
)); /* Src0. */
974 /* Implement TARGET_LEGITIMATE_CONSTANT_P.
976 Returns true if X is a legitimate constant for a MODE immediate operand. */
979 gcn_legitimate_constant_p (machine_mode
, rtx x
)
981 return gcn_constant_p (x
);
984 /* Return true if X is a CONST_VECTOR of single constant. */
987 single_cst_vector_p (rtx x
)
989 if (GET_CODE (x
) != CONST_VECTOR
)
991 for (int i
= 1; i
< 64; i
++)
992 if (CONST_VECTOR_ELT (x
, i
) != CONST_VECTOR_ELT (x
, 0))
997 /* Create a CONST_VECTOR of duplicated value A. */
1000 gcn_vec_constant (machine_mode mode
, int a
)
1003 return CONST0_RTX (mode);
1005 return CONSTM1_RTX (mode);
1007 return CONST1_RTX (mode);
1009 return CONST2_RTX (mode);*/
1011 int units
= GET_MODE_NUNITS (mode
);
1012 machine_mode innermode
= GET_MODE_INNER (mode
);
1015 if (FLOAT_MODE_P (innermode
))
1018 real_from_integer (&rv
, NULL
, a
, SIGNED
);
1019 tem
= const_double_from_real_value (rv
, innermode
);
1022 tem
= gen_int_mode (a
, innermode
);
1024 rtvec v
= rtvec_alloc (units
);
1025 for (int i
= 0; i
< units
; ++i
)
1026 RTVEC_ELT (v
, i
) = tem
;
1028 return gen_rtx_CONST_VECTOR (mode
, v
);
1031 /* Create a CONST_VECTOR of duplicated value A. */
1034 gcn_vec_constant (machine_mode mode
, rtx a
)
1036 int units
= GET_MODE_NUNITS (mode
);
1037 rtvec v
= rtvec_alloc (units
);
1039 for (int i
= 0; i
< units
; ++i
)
1040 RTVEC_ELT (v
, i
) = a
;
1042 return gen_rtx_CONST_VECTOR (mode
, v
);
1045 /* Create an undefined vector value, used where an insn operand is
1049 gcn_gen_undef (machine_mode mode
)
1051 return gen_rtx_UNSPEC (mode
, gen_rtvec (1, const0_rtx
), UNSPEC_VECTOR
);
1055 /* {{{ Addresses, pointers and moves. */
1057 /* Return true is REG is a valid place to store a pointer,
1058 for instructions that require an SGPR.
1062 gcn_address_register_p (rtx reg
, machine_mode mode
, bool strict
)
1064 if (GET_CODE (reg
) == SUBREG
)
1065 reg
= SUBREG_REG (reg
);
1070 if (GET_MODE (reg
) != mode
)
1073 int regno
= REGNO (reg
);
1075 if (regno
>= FIRST_PSEUDO_REGISTER
)
1083 regno
= reg_renumber
[regno
];
1086 return (SGPR_REGNO_P (regno
) || regno
== M0_REG
1087 || regno
== ARG_POINTER_REGNUM
|| regno
== FRAME_POINTER_REGNUM
);
1090 /* Return true is REG is a valid place to store a pointer,
1091 for instructions that require a VGPR. */
1094 gcn_vec_address_register_p (rtx reg
, machine_mode mode
, bool strict
)
1096 if (GET_CODE (reg
) == SUBREG
)
1097 reg
= SUBREG_REG (reg
);
1102 if (GET_MODE (reg
) != mode
)
1105 int regno
= REGNO (reg
);
1107 if (regno
>= FIRST_PSEUDO_REGISTER
)
1115 regno
= reg_renumber
[regno
];
1118 return VGPR_REGNO_P (regno
);
1121 /* Return true if X would be valid inside a MEM using the Flat address
1125 gcn_flat_address_p (rtx x
, machine_mode mode
)
1127 bool vec_mode
= (GET_MODE_CLASS (mode
) == MODE_VECTOR_INT
1128 || GET_MODE_CLASS (mode
) == MODE_VECTOR_FLOAT
);
1130 if (vec_mode
&& gcn_address_register_p (x
, DImode
, false))
1133 if (!vec_mode
&& gcn_vec_address_register_p (x
, DImode
, false))
1136 if (TARGET_GCN5_PLUS
1137 && GET_CODE (x
) == PLUS
1138 && gcn_vec_address_register_p (XEXP (x
, 0), DImode
, false)
1139 && CONST_INT_P (XEXP (x
, 1)))
1145 /* Return true if X would be valid inside a MEM using the Scalar Flat
1149 gcn_scalar_flat_address_p (rtx x
)
1151 if (gcn_address_register_p (x
, DImode
, false))
1154 if (GET_CODE (x
) == PLUS
1155 && gcn_address_register_p (XEXP (x
, 0), DImode
, false)
1156 && CONST_INT_P (XEXP (x
, 1)))
1162 /* Return true if MEM X would be valid for the Scalar Flat address space. */
1165 gcn_scalar_flat_mem_p (rtx x
)
1170 if (GET_MODE_SIZE (GET_MODE (x
)) < 4)
1173 return gcn_scalar_flat_address_p (XEXP (x
, 0));
1176 /* Return true if X would be valid inside a MEM using the LDS or GDS
1180 gcn_ds_address_p (rtx x
)
1182 if (gcn_vec_address_register_p (x
, SImode
, false))
1185 if (GET_CODE (x
) == PLUS
1186 && gcn_vec_address_register_p (XEXP (x
, 0), SImode
, false)
1187 && CONST_INT_P (XEXP (x
, 1)))
1193 /* Return true if ADDR would be valid inside a MEM using the Global
1197 gcn_global_address_p (rtx addr
)
1199 if (gcn_address_register_p (addr
, DImode
, false)
1200 || gcn_vec_address_register_p (addr
, DImode
, false))
1203 if (GET_CODE (addr
) == PLUS
)
1205 rtx base
= XEXP (addr
, 0);
1206 rtx offset
= XEXP (addr
, 1);
1207 bool immediate_p
= (CONST_INT_P (offset
)
1208 && INTVAL (offset
) >= -(1 << 12)
1209 && INTVAL (offset
) < (1 << 12));
1211 if ((gcn_address_register_p (base
, DImode
, false)
1212 || gcn_vec_address_register_p (base
, DImode
, false))
1214 /* SGPR + CONST or VGPR + CONST */
1217 if (gcn_address_register_p (base
, DImode
, false)
1218 && gcn_vgpr_register_operand (offset
, SImode
))
1222 if (GET_CODE (base
) == PLUS
1223 && gcn_address_register_p (XEXP (base
, 0), DImode
, false)
1224 && gcn_vgpr_register_operand (XEXP (base
, 1), SImode
)
1226 /* (SGPR + VGPR) + CONST */
1233 /* Implement TARGET_ADDR_SPACE_LEGITIMATE_ADDRESS_P.
1235 Recognizes RTL expressions that are valid memory addresses for an
1236 instruction. The MODE argument is the machine mode for the MEM
1237 expression that wants to use this address.
1239 It only recognizes address in canonical form. LEGITIMIZE_ADDRESS should
1240 convert common non-canonical forms to canonical form so that they will
1244 gcn_addr_space_legitimate_address_p (machine_mode mode
, rtx x
, bool strict
,
1247 /* All vector instructions need to work on addresses in registers. */
1248 if (!TARGET_GCN5_PLUS
&& (vgpr_vector_mode_p (mode
) && !REG_P (x
)))
1251 if (AS_SCALAR_FLAT_P (as
))
1253 if (mode
== QImode
|| mode
== HImode
)
1256 switch (GET_CODE (x
))
1259 return gcn_address_register_p (x
, DImode
, strict
);
1260 /* Addresses are in the form BASE+OFFSET
1261 OFFSET is either 20bit unsigned immediate, SGPR or M0.
1262 Writes and atomics do not accept SGPR. */
1265 rtx x0
= XEXP (x
, 0);
1266 rtx x1
= XEXP (x
, 1);
1267 if (!gcn_address_register_p (x0
, DImode
, strict
))
1269 /* FIXME: This is disabled because of the mode mismatch between
1270 SImode (for the address or m0 register) and the DImode PLUS.
1271 We'll need a zero_extend or similar.
1273 if (gcn_m0_register_p (x1, SImode, strict)
1274 || gcn_address_register_p (x1, SImode, strict))
1277 if (GET_CODE (x1
) == CONST_INT
)
1279 if (INTVAL (x1
) >= 0 && INTVAL (x1
) < (1 << 20)
1280 /* The low bits of the offset are ignored, even when
1281 they're meant to realign the pointer. */
1282 && !(INTVAL (x1
) & 0x3))
1292 else if (AS_SCRATCH_P (as
))
1293 return gcn_address_register_p (x
, SImode
, strict
);
1294 else if (AS_FLAT_P (as
) || AS_FLAT_SCRATCH_P (as
))
1296 if (TARGET_GCN3
|| GET_CODE (x
) == REG
)
1297 return ((GET_MODE_CLASS (mode
) == MODE_VECTOR_INT
1298 || GET_MODE_CLASS (mode
) == MODE_VECTOR_FLOAT
)
1299 ? gcn_address_register_p (x
, DImode
, strict
)
1300 : gcn_vec_address_register_p (x
, DImode
, strict
));
1303 gcc_assert (TARGET_GCN5_PLUS
);
1305 if (GET_CODE (x
) == PLUS
)
1307 rtx x1
= XEXP (x
, 1);
1309 if (VECTOR_MODE_P (mode
)
1310 ? !gcn_address_register_p (x
, DImode
, strict
)
1311 : !gcn_vec_address_register_p (x
, DImode
, strict
))
1314 if (GET_CODE (x1
) == CONST_INT
)
1316 if (INTVAL (x1
) >= 0 && INTVAL (x1
) < (1 << 12)
1317 /* The low bits of the offset are ignored, even when
1318 they're meant to realign the pointer. */
1319 && !(INTVAL (x1
) & 0x3))
1326 else if (AS_GLOBAL_P (as
))
1328 gcc_assert (TARGET_GCN5_PLUS
);
1330 if (GET_CODE (x
) == REG
)
1331 return (gcn_address_register_p (x
, DImode
, strict
)
1332 || (!VECTOR_MODE_P (mode
)
1333 && gcn_vec_address_register_p (x
, DImode
, strict
)));
1334 else if (GET_CODE (x
) == PLUS
)
1336 rtx base
= XEXP (x
, 0);
1337 rtx offset
= XEXP (x
, 1);
1339 bool immediate_p
= (GET_CODE (offset
) == CONST_INT
1340 /* Signed 13-bit immediate. */
1341 && INTVAL (offset
) >= -(1 << 12)
1342 && INTVAL (offset
) < (1 << 12)
1343 /* The low bits of the offset are ignored, even
1344 when they're meant to realign the pointer. */
1345 && !(INTVAL (offset
) & 0x3));
1347 if (!VECTOR_MODE_P (mode
))
1349 if ((gcn_address_register_p (base
, DImode
, strict
)
1350 || gcn_vec_address_register_p (base
, DImode
, strict
))
1352 /* SGPR + CONST or VGPR + CONST */
1355 if (gcn_address_register_p (base
, DImode
, strict
)
1356 && gcn_vgpr_register_operand (offset
, SImode
))
1360 if (GET_CODE (base
) == PLUS
1361 && gcn_address_register_p (XEXP (base
, 0), DImode
, strict
)
1362 && gcn_vgpr_register_operand (XEXP (base
, 1), SImode
)
1364 /* (SGPR + VGPR) + CONST */
1369 if (gcn_address_register_p (base
, DImode
, strict
)
1378 else if (AS_ANY_DS_P (as
))
1379 switch (GET_CODE (x
))
1382 return (VECTOR_MODE_P (mode
)
1383 ? gcn_address_register_p (x
, SImode
, strict
)
1384 : gcn_vec_address_register_p (x
, SImode
, strict
));
1385 /* Addresses are in the form BASE+OFFSET
1386 OFFSET is either 20bit unsigned immediate, SGPR or M0.
1387 Writes and atomics do not accept SGPR. */
1390 rtx x0
= XEXP (x
, 0);
1391 rtx x1
= XEXP (x
, 1);
1392 if (!gcn_vec_address_register_p (x0
, DImode
, strict
))
1394 if (GET_CODE (x1
) == REG
)
1396 if (GET_CODE (x1
) != REG
1397 || (REGNO (x1
) <= FIRST_PSEUDO_REGISTER
1398 && !gcn_ssrc_register_operand (x1
, DImode
)))
1401 else if (GET_CODE (x1
) == CONST_VECTOR
1402 && GET_CODE (CONST_VECTOR_ELT (x1
, 0)) == CONST_INT
1403 && single_cst_vector_p (x1
))
1405 x1
= CONST_VECTOR_ELT (x1
, 0);
1406 if (INTVAL (x1
) >= 0 && INTVAL (x1
) < (1 << 20))
1420 /* Implement TARGET_ADDR_SPACE_POINTER_MODE.
1422 Return the appropriate mode for a named address pointer. */
1424 static scalar_int_mode
1425 gcn_addr_space_pointer_mode (addr_space_t addrspace
)
1429 case ADDR_SPACE_SCRATCH
:
1430 case ADDR_SPACE_LDS
:
1431 case ADDR_SPACE_GDS
:
1433 case ADDR_SPACE_DEFAULT
:
1434 case ADDR_SPACE_FLAT
:
1435 case ADDR_SPACE_FLAT_SCRATCH
:
1436 case ADDR_SPACE_SCALAR_FLAT
:
1443 /* Implement TARGET_ADDR_SPACE_ADDRESS_MODE.
1445 Return the appropriate mode for a named address space address. */
1447 static scalar_int_mode
1448 gcn_addr_space_address_mode (addr_space_t addrspace
)
1450 return gcn_addr_space_pointer_mode (addrspace
);
1453 /* Implement TARGET_ADDR_SPACE_SUBSET_P.
1455 Determine if one named address space is a subset of another. */
1458 gcn_addr_space_subset_p (addr_space_t subset
, addr_space_t superset
)
1460 if (subset
== superset
)
1462 /* FIXME is this true? */
1463 if (AS_FLAT_P (superset
) || AS_SCALAR_FLAT_P (superset
))
1468 /* Convert from one address space to another. */
1471 gcn_addr_space_convert (rtx op
, tree from_type
, tree to_type
)
1473 gcc_assert (POINTER_TYPE_P (from_type
));
1474 gcc_assert (POINTER_TYPE_P (to_type
));
1476 addr_space_t as_from
= TYPE_ADDR_SPACE (TREE_TYPE (from_type
));
1477 addr_space_t as_to
= TYPE_ADDR_SPACE (TREE_TYPE (to_type
));
1479 if (AS_LDS_P (as_from
) && AS_FLAT_P (as_to
))
1481 rtx queue
= gen_rtx_REG (DImode
,
1482 cfun
->machine
->args
.reg
[QUEUE_PTR_ARG
]);
1483 rtx group_seg_aperture_hi
= gen_rtx_MEM (SImode
,
1484 gen_rtx_PLUS (DImode
, queue
,
1485 gen_int_mode (64, SImode
)));
1486 rtx tmp
= gen_reg_rtx (DImode
);
1488 emit_move_insn (gen_lowpart (SImode
, tmp
), op
);
1489 emit_move_insn (gen_highpart_mode (SImode
, DImode
, tmp
),
1490 group_seg_aperture_hi
);
1494 else if (as_from
== as_to
)
1501 /* Implement REGNO_MODE_CODE_OK_FOR_BASE_P via gcn.h
1503 Retun true if REGNO is OK for memory adressing. */
1506 gcn_regno_mode_code_ok_for_base_p (int regno
,
1507 machine_mode
, addr_space_t as
, int, int)
1509 if (regno
>= FIRST_PSEUDO_REGISTER
)
1512 regno
= reg_renumber
[regno
];
1517 return (VGPR_REGNO_P (regno
)
1518 || regno
== ARG_POINTER_REGNUM
|| regno
== FRAME_POINTER_REGNUM
);
1519 else if (AS_SCALAR_FLAT_P (as
))
1520 return (SGPR_REGNO_P (regno
)
1521 || regno
== ARG_POINTER_REGNUM
|| regno
== FRAME_POINTER_REGNUM
);
1522 else if (AS_GLOBAL_P (as
))
1524 return (SGPR_REGNO_P (regno
)
1525 || VGPR_REGNO_P (regno
)
1526 || regno
== ARG_POINTER_REGNUM
1527 || regno
== FRAME_POINTER_REGNUM
);
1534 /* Implement MODE_CODE_BASE_REG_CLASS via gcn.h.
1536 Return a suitable register class for memory addressing. */
1539 gcn_mode_code_base_reg_class (machine_mode mode
, addr_space_t as
, int oc
,
1544 case ADDR_SPACE_DEFAULT
:
1545 return gcn_mode_code_base_reg_class (mode
, DEFAULT_ADDR_SPACE
, oc
, ic
);
1546 case ADDR_SPACE_SCALAR_FLAT
:
1547 case ADDR_SPACE_SCRATCH
:
1550 case ADDR_SPACE_FLAT
:
1551 case ADDR_SPACE_FLAT_SCRATCH
:
1552 case ADDR_SPACE_LDS
:
1553 case ADDR_SPACE_GDS
:
1554 return ((GET_MODE_CLASS (mode
) == MODE_VECTOR_INT
1555 || GET_MODE_CLASS (mode
) == MODE_VECTOR_FLOAT
)
1556 ? SGPR_REGS
: VGPR_REGS
);
1557 case ADDR_SPACE_GLOBAL
:
1558 return ((GET_MODE_CLASS (mode
) == MODE_VECTOR_INT
1559 || GET_MODE_CLASS (mode
) == MODE_VECTOR_FLOAT
)
1560 ? SGPR_REGS
: ALL_GPR_REGS
);
1565 /* Implement REGNO_OK_FOR_INDEX_P via gcn.h.
1567 Return true if REGNO is OK for index of memory addressing. */
1570 regno_ok_for_index_p (int regno
)
1572 if (regno
>= FIRST_PSEUDO_REGISTER
)
1575 regno
= reg_renumber
[regno
];
1579 return regno
== M0_REG
|| VGPR_REGNO_P (regno
);
1582 /* Generate move which uses the exec flags. If EXEC is NULL, then it is
1583 assumed that all lanes normally relevant to the mode of the move are
1584 affected. If PREV is NULL, then a sensible default is supplied for
1585 the inactive lanes. */
1588 gen_mov_with_exec (rtx op0
, rtx op1
, rtx exec
= NULL
, rtx prev
= NULL
)
1590 machine_mode mode
= GET_MODE (op0
);
1592 if (vgpr_vector_mode_p (mode
))
1594 if (exec
&& exec
!= CONSTM1_RTX (DImode
))
1602 prev
= gcn_gen_undef (mode
);
1603 exec
= gcn_full_exec_reg ();
1606 rtx set
= gen_rtx_SET (op0
, gen_rtx_VEC_MERGE (mode
, op1
, prev
, exec
));
1608 return gen_rtx_PARALLEL (VOIDmode
,
1610 gen_rtx_CLOBBER (VOIDmode
,
1611 gen_rtx_SCRATCH (V64DImode
))));
1614 return (gen_rtx_PARALLEL
1616 gen_rtvec (2, gen_rtx_SET (op0
, op1
),
1617 gen_rtx_USE (VOIDmode
,
1618 exec
? exec
: gcn_scalar_exec ()))));
1621 /* Generate masked move. */
1624 gen_duplicate_load (rtx op0
, rtx op1
, rtx op2
= NULL
, rtx exec
= NULL
)
1627 return (gen_rtx_SET (op0
,
1628 gen_rtx_VEC_MERGE (GET_MODE (op0
),
1629 gen_rtx_VEC_DUPLICATE (GET_MODE
1633 return (gen_rtx_SET (op0
, gen_rtx_VEC_DUPLICATE (GET_MODE (op0
), op1
)));
1636 /* Expand vector init of OP0 by VEC.
1637 Implements vec_init instruction pattern. */
1640 gcn_expand_vector_init (rtx op0
, rtx vec
)
1642 int64_t initialized_mask
= 0;
1643 int64_t curr_mask
= 1;
1644 machine_mode mode
= GET_MODE (op0
);
1646 rtx val
= XVECEXP (vec
, 0, 0);
1648 for (int i
= 1; i
< 64; i
++)
1649 if (rtx_equal_p (val
, XVECEXP (vec
, 0, i
)))
1650 curr_mask
|= (int64_t) 1 << i
;
1652 if (gcn_constant_p (val
))
1653 emit_move_insn (op0
, gcn_vec_constant (mode
, val
));
1656 val
= force_reg (GET_MODE_INNER (mode
), val
);
1657 emit_insn (gen_duplicate_load (op0
, val
));
1659 initialized_mask
|= curr_mask
;
1660 for (int i
= 1; i
< 64; i
++)
1661 if (!(initialized_mask
& ((int64_t) 1 << i
)))
1663 curr_mask
= (int64_t) 1 << i
;
1664 rtx val
= XVECEXP (vec
, 0, i
);
1666 for (int j
= i
+ 1; j
< 64; j
++)
1667 if (rtx_equal_p (val
, XVECEXP (vec
, 0, j
)))
1668 curr_mask
|= (int64_t) 1 << j
;
1669 if (gcn_constant_p (val
))
1670 emit_insn (gen_mov_with_exec (op0
, gcn_vec_constant (mode
, val
),
1671 get_exec (curr_mask
)));
1674 val
= force_reg (GET_MODE_INNER (mode
), val
);
1675 emit_insn (gen_duplicate_load (op0
, val
, op0
,
1676 get_exec (curr_mask
)));
1678 initialized_mask
|= curr_mask
;
1682 /* Load vector constant where n-th lane contains BASE+n*VAL. */
1685 strided_constant (machine_mode mode
, int base
, int val
)
1687 rtx x
= gen_reg_rtx (mode
);
1688 emit_move_insn (x
, gcn_vec_constant (mode
, base
));
1689 emit_insn (gen_addv64si3_exec (x
, x
, gcn_vec_constant (mode
, val
* 32),
1690 x
, get_exec (0xffffffff00000000)));
1691 emit_insn (gen_addv64si3_exec (x
, x
, gcn_vec_constant (mode
, val
* 16),
1692 x
, get_exec (0xffff0000ffff0000)));
1693 emit_insn (gen_addv64si3_exec (x
, x
, gcn_vec_constant (mode
, val
* 8),
1694 x
, get_exec (0xff00ff00ff00ff00)));
1695 emit_insn (gen_addv64si3_exec (x
, x
, gcn_vec_constant (mode
, val
* 4),
1696 x
, get_exec (0xf0f0f0f0f0f0f0f0)));
1697 emit_insn (gen_addv64si3_exec (x
, x
, gcn_vec_constant (mode
, val
* 2),
1698 x
, get_exec (0xcccccccccccccccc)));
1699 emit_insn (gen_addv64si3_exec (x
, x
, gcn_vec_constant (mode
, val
* 1),
1700 x
, get_exec (0xaaaaaaaaaaaaaaaa)));
1704 /* Implement TARGET_ADDR_SPACE_LEGITIMIZE_ADDRESS. */
1707 gcn_addr_space_legitimize_address (rtx x
, rtx old
, machine_mode mode
,
1712 case ADDR_SPACE_DEFAULT
:
1713 return gcn_addr_space_legitimize_address (x
, old
, mode
,
1714 DEFAULT_ADDR_SPACE
);
1715 case ADDR_SPACE_SCALAR_FLAT
:
1716 case ADDR_SPACE_SCRATCH
:
1717 /* Instructions working on vectors need the address to be in
1719 if (vgpr_vector_mode_p (mode
))
1720 return force_reg (GET_MODE (x
), x
);
1723 case ADDR_SPACE_FLAT
:
1724 case ADDR_SPACE_FLAT_SCRATCH
:
1725 case ADDR_SPACE_GLOBAL
:
1726 return TARGET_GCN3
? force_reg (DImode
, x
) : x
;
1727 case ADDR_SPACE_LDS
:
1728 case ADDR_SPACE_GDS
:
1729 /* FIXME: LDS support offsets, handle them!. */
1730 if (vgpr_vector_mode_p (mode
) && GET_MODE (x
) != V64SImode
)
1732 rtx addrs
= gen_reg_rtx (V64SImode
);
1733 rtx base
= force_reg (SImode
, x
);
1734 rtx offsets
= strided_constant (V64SImode
, 0,
1735 GET_MODE_UNIT_SIZE (mode
));
1737 emit_insn (gen_vec_duplicatev64si (addrs
, base
));
1738 emit_insn (gen_addv64si3 (addrs
, offsets
, addrs
));
1746 /* Convert a (mem:<MODE> (reg:DI)) to (mem:<MODE> (reg:V64DI)) with the
1747 proper vector of stepped addresses.
1749 MEM will be a DImode address of a vector in an SGPR.
1750 TMP will be a V64DImode VGPR pair or (scratch:V64DI). */
1753 gcn_expand_scalar_to_vector_address (machine_mode mode
, rtx exec
, rtx mem
,
1756 gcc_assert (MEM_P (mem
));
1757 rtx mem_base
= XEXP (mem
, 0);
1758 rtx mem_index
= NULL_RTX
;
1760 if (!TARGET_GCN5_PLUS
)
1762 /* gcn_addr_space_legitimize_address should have put the address in a
1763 register. If not, it is too late to do anything about it. */
1764 gcc_assert (REG_P (mem_base
));
1767 if (GET_CODE (mem_base
) == PLUS
)
1769 mem_index
= XEXP (mem_base
, 1);
1770 mem_base
= XEXP (mem_base
, 0);
1773 /* RF and RM base registers for vector modes should be always an SGPR. */
1774 gcc_assert (SGPR_REGNO_P (REGNO (mem_base
))
1775 || REGNO (mem_base
) >= FIRST_PSEUDO_REGISTER
);
1777 machine_mode inner
= GET_MODE_INNER (mode
);
1778 int shift
= exact_log2 (GET_MODE_SIZE (inner
));
1779 rtx ramp
= gen_rtx_REG (V64SImode
, VGPR_REGNO (1));
1780 rtx undef_v64si
= gcn_gen_undef (V64SImode
);
1781 rtx new_base
= NULL_RTX
;
1782 addr_space_t as
= MEM_ADDR_SPACE (mem
);
1784 rtx tmplo
= (REG_P (tmp
)
1785 ? gcn_operand_part (V64DImode
, tmp
, 0)
1786 : gen_reg_rtx (V64SImode
));
1788 /* tmplo[:] = ramp[:] << shift */
1790 emit_insn (gen_ashlv64si3_exec (tmplo
, ramp
,
1791 gen_int_mode (shift
, SImode
),
1792 undef_v64si
, exec
));
1794 emit_insn (gen_ashlv64si3 (tmplo
, ramp
, gen_int_mode (shift
, SImode
)));
1798 rtx vcc
= gen_rtx_REG (DImode
, CC_SAVE_REG
);
1802 rtx mem_base_lo
= gcn_operand_part (DImode
, mem_base
, 0);
1803 rtx mem_base_hi
= gcn_operand_part (DImode
, mem_base
, 1);
1804 rtx tmphi
= gcn_operand_part (V64DImode
, tmp
, 1);
1806 /* tmphi[:] = mem_base_hi */
1808 emit_insn (gen_vec_duplicatev64si_exec (tmphi
, mem_base_hi
,
1809 undef_v64si
, exec
));
1811 emit_insn (gen_vec_duplicatev64si (tmphi
, mem_base_hi
));
1813 /* tmp[:] += zext (mem_base) */
1816 emit_insn (gen_addv64si3_vcc_dup_exec (tmplo
, mem_base_lo
, tmplo
,
1817 vcc
, undef_v64si
, exec
));
1818 emit_insn (gen_addcv64si3_exec (tmphi
, tmphi
, const0_rtx
,
1819 vcc
, vcc
, undef_v64si
, exec
));
1822 emit_insn (gen_addv64di3_vcc_zext_dup (tmp
, mem_base_lo
, tmp
, vcc
));
1826 tmp
= gen_reg_rtx (V64DImode
);
1828 emit_insn (gen_addv64di3_vcc_zext_dup2_exec
1829 (tmp
, tmplo
, mem_base
, vcc
, gcn_gen_undef (V64DImode
),
1832 emit_insn (gen_addv64di3_vcc_zext_dup2 (tmp
, tmplo
, mem_base
, vcc
));
1837 else if (AS_ANY_DS_P (as
))
1840 emit_insn (gen_addv64si3_dup (tmplo
, tmplo
, mem_base
));
1842 emit_insn (gen_addv64si3_dup_exec (tmplo
, tmplo
, mem_base
,
1843 gcn_gen_undef (V64SImode
), exec
));
1848 mem_base
= gen_rtx_VEC_DUPLICATE (V64DImode
, mem_base
);
1849 new_base
= gen_rtx_PLUS (V64DImode
, mem_base
,
1850 gen_rtx_SIGN_EXTEND (V64DImode
, tmplo
));
1853 return gen_rtx_PLUS (GET_MODE (new_base
), new_base
,
1854 gen_rtx_VEC_DUPLICATE (GET_MODE (new_base
),
1855 (mem_index
? mem_index
1859 /* Convert a BASE address, a vector of OFFSETS, and a SCALE, to addresses
1860 suitable for the given address space. This is indented for use in
1861 gather/scatter patterns.
1863 The offsets may be signed or unsigned, according to UNSIGNED_P.
1864 If EXEC is set then _exec patterns will be used, otherwise plain.
1867 ADDR_SPACE_FLAT - return V64DImode vector of absolute addresses.
1868 ADDR_SPACE_GLOBAL - return V64SImode vector of offsets. */
1871 gcn_expand_scaled_offsets (addr_space_t as
, rtx base
, rtx offsets
, rtx scale
,
1872 bool unsigned_p
, rtx exec
)
1874 rtx tmpsi
= gen_reg_rtx (V64SImode
);
1875 rtx tmpdi
= gen_reg_rtx (V64DImode
);
1876 rtx undefsi
= exec
? gcn_gen_undef (V64SImode
) : NULL
;
1877 rtx undefdi
= exec
? gcn_gen_undef (V64DImode
) : NULL
;
1879 if (CONST_INT_P (scale
)
1880 && INTVAL (scale
) > 0
1881 && exact_log2 (INTVAL (scale
)) >= 0)
1882 emit_insn (gen_ashlv64si3 (tmpsi
, offsets
,
1883 GEN_INT (exact_log2 (INTVAL (scale
)))));
1886 ? emit_insn (gen_mulv64si3_dup_exec (tmpsi
, offsets
, scale
, undefsi
,
1888 : emit_insn (gen_mulv64si3_dup (tmpsi
, offsets
, scale
)));
1890 /* "Global" instructions do not support negative register offsets. */
1891 if (as
== ADDR_SPACE_FLAT
|| !unsigned_p
)
1895 ? emit_insn (gen_addv64di3_zext_dup2_exec (tmpdi
, tmpsi
, base
,
1897 : emit_insn (gen_addv64di3_zext_dup2 (tmpdi
, tmpsi
, base
)));
1900 ? emit_insn (gen_addv64di3_sext_dup2_exec (tmpdi
, tmpsi
, base
,
1902 : emit_insn (gen_addv64di3_sext_dup2 (tmpdi
, tmpsi
, base
)));
1905 else if (as
== ADDR_SPACE_GLOBAL
)
1911 /* Return true if move from OP0 to OP1 is known to be executed in vector
1915 gcn_vgpr_move_p (rtx op0
, rtx op1
)
1917 if (MEM_P (op0
) && AS_SCALAR_FLAT_P (MEM_ADDR_SPACE (op0
)))
1919 if (MEM_P (op1
) && AS_SCALAR_FLAT_P (MEM_ADDR_SPACE (op1
)))
1921 return ((REG_P (op0
) && VGPR_REGNO_P (REGNO (op0
)))
1922 || (REG_P (op1
) && VGPR_REGNO_P (REGNO (op1
)))
1923 || vgpr_vector_mode_p (GET_MODE (op0
)));
1926 /* Return true if move from OP0 to OP1 is known to be executed in scalar
1927 unit. Used in the machine description. */
1930 gcn_sgpr_move_p (rtx op0
, rtx op1
)
1932 if (MEM_P (op0
) && AS_SCALAR_FLAT_P (MEM_ADDR_SPACE (op0
)))
1934 if (MEM_P (op1
) && AS_SCALAR_FLAT_P (MEM_ADDR_SPACE (op1
)))
1936 if (!REG_P (op0
) || REGNO (op0
) >= FIRST_PSEUDO_REGISTER
1937 || VGPR_REGNO_P (REGNO (op0
)))
1940 && REGNO (op1
) < FIRST_PSEUDO_REGISTER
1941 && !VGPR_REGNO_P (REGNO (op1
)))
1943 return immediate_operand (op1
, VOIDmode
) || memory_operand (op1
, VOIDmode
);
1946 /* Implement TARGET_SECONDARY_RELOAD.
1948 The address space determines which registers can be used for loads and
1952 gcn_secondary_reload (bool in_p
, rtx x
, reg_class_t rclass
,
1953 machine_mode reload_mode
, secondary_reload_info
*sri
)
1955 reg_class_t result
= NO_REGS
;
1956 bool spilled_pseudo
=
1957 (REG_P (x
) || GET_CODE (x
) == SUBREG
) && true_regnum (x
) == -1;
1959 if (dump_file
&& (dump_flags
& TDF_DETAILS
))
1961 fprintf (dump_file
, "gcn_secondary_reload: ");
1962 dump_value_slim (dump_file
, x
, 1);
1963 fprintf (dump_file
, " %s %s:%s", (in_p
? "->" : "<-"),
1964 reg_class_names
[rclass
], GET_MODE_NAME (reload_mode
));
1965 if (REG_P (x
) || GET_CODE (x
) == SUBREG
)
1966 fprintf (dump_file
, " (true regnum: %d \"%s\")", true_regnum (x
),
1967 (true_regnum (x
) >= 0
1968 && true_regnum (x
) < FIRST_PSEUDO_REGISTER
1969 ? reg_names
[true_regnum (x
)]
1970 : (spilled_pseudo
? "stack spill" : "??")));
1971 fprintf (dump_file
, "\n");
1974 /* Some callers don't use or initialize icode. */
1975 sri
->icode
= CODE_FOR_nothing
;
1977 if (MEM_P (x
) || spilled_pseudo
)
1979 addr_space_t as
= DEFAULT_ADDR_SPACE
;
1981 /* If we have a spilled pseudo, we can't find the address space
1982 directly, but we know it's in ADDR_SPACE_FLAT space for GCN3 or
1983 ADDR_SPACE_GLOBAL for GCN5. */
1985 as
= MEM_ADDR_SPACE (x
);
1987 if (as
== ADDR_SPACE_DEFAULT
)
1988 as
= DEFAULT_ADDR_SPACE
;
1992 case ADDR_SPACE_SCALAR_FLAT
:
1994 ((!MEM_P (x
) || rclass
== SGPR_REGS
) ? NO_REGS
: SGPR_REGS
);
1996 case ADDR_SPACE_FLAT
:
1997 case ADDR_SPACE_FLAT_SCRATCH
:
1998 case ADDR_SPACE_GLOBAL
:
1999 if (GET_MODE_CLASS (reload_mode
) == MODE_VECTOR_INT
2000 || GET_MODE_CLASS (reload_mode
) == MODE_VECTOR_FLOAT
)
2003 switch (reload_mode
)
2006 sri
->icode
= CODE_FOR_reload_inv64si
;
2009 sri
->icode
= CODE_FOR_reload_inv64sf
;
2012 sri
->icode
= CODE_FOR_reload_inv64hi
;
2015 sri
->icode
= CODE_FOR_reload_inv64hf
;
2018 sri
->icode
= CODE_FOR_reload_inv64qi
;
2021 sri
->icode
= CODE_FOR_reload_inv64di
;
2024 sri
->icode
= CODE_FOR_reload_inv64df
;
2030 switch (reload_mode
)
2033 sri
->icode
= CODE_FOR_reload_outv64si
;
2036 sri
->icode
= CODE_FOR_reload_outv64sf
;
2039 sri
->icode
= CODE_FOR_reload_outv64hi
;
2042 sri
->icode
= CODE_FOR_reload_outv64hf
;
2045 sri
->icode
= CODE_FOR_reload_outv64qi
;
2048 sri
->icode
= CODE_FOR_reload_outv64di
;
2051 sri
->icode
= CODE_FOR_reload_outv64df
;
2059 case ADDR_SPACE_LDS
:
2060 case ADDR_SPACE_GDS
:
2061 case ADDR_SPACE_SCRATCH
:
2062 result
= (rclass
== VGPR_REGS
? NO_REGS
: VGPR_REGS
);
2067 if (dump_file
&& (dump_flags
& TDF_DETAILS
))
2068 fprintf (dump_file
, " <= %s (icode: %s)\n", reg_class_names
[result
],
2069 get_insn_name (sri
->icode
));
2074 /* Update register usage after having seen the compiler flags and kernel
2075 attributes. We typically want to fix registers that contain values
2076 set by the HSA runtime. */
2079 gcn_conditional_register_usage (void)
2081 if (!cfun
|| !cfun
->machine
)
2084 if (cfun
->machine
->normal_function
)
2086 /* Restrict the set of SGPRs and VGPRs used by non-kernel functions. */
2087 for (int i
= SGPR_REGNO (MAX_NORMAL_SGPR_COUNT
);
2088 i
<= LAST_SGPR_REG
; i
++)
2089 fixed_regs
[i
] = 1, call_used_regs
[i
] = 1;
2091 for (int i
= VGPR_REGNO (MAX_NORMAL_VGPR_COUNT
);
2092 i
<= LAST_VGPR_REG
; i
++)
2093 fixed_regs
[i
] = 1, call_used_regs
[i
] = 1;
2098 /* If the set of requested args is the default set, nothing more needs to
2100 if (cfun
->machine
->args
.requested
== default_requested_args
)
2103 /* Requesting a set of args different from the default violates the ABI. */
2104 if (!leaf_function_p ())
2105 warning (0, "A non-default set of initial values has been requested, "
2106 "which violates the ABI");
2108 for (int i
= SGPR_REGNO (0); i
< SGPR_REGNO (14); i
++)
2111 /* Fix the runtime argument register containing values that may be
2112 needed later. DISPATCH_PTR_ARG and FLAT_SCRATCH_* should not be
2113 needed after the prologue so there's no need to fix them. */
2114 if (cfun
->machine
->args
.reg
[PRIVATE_SEGMENT_WAVE_OFFSET_ARG
] >= 0)
2115 fixed_regs
[cfun
->machine
->args
.reg
[PRIVATE_SEGMENT_WAVE_OFFSET_ARG
]] = 1;
2116 if (cfun
->machine
->args
.reg
[PRIVATE_SEGMENT_BUFFER_ARG
] >= 0)
2118 /* The upper 32-bits of the 64-bit descriptor are not used, so allow
2119 the containing registers to be used for other purposes. */
2120 fixed_regs
[cfun
->machine
->args
.reg
[PRIVATE_SEGMENT_BUFFER_ARG
]] = 1;
2121 fixed_regs
[cfun
->machine
->args
.reg
[PRIVATE_SEGMENT_BUFFER_ARG
] + 1] = 1;
2123 if (cfun
->machine
->args
.reg
[KERNARG_SEGMENT_PTR_ARG
] >= 0)
2125 fixed_regs
[cfun
->machine
->args
.reg
[KERNARG_SEGMENT_PTR_ARG
]] = 1;
2126 fixed_regs
[cfun
->machine
->args
.reg
[KERNARG_SEGMENT_PTR_ARG
] + 1] = 1;
2128 if (cfun
->machine
->args
.reg
[DISPATCH_PTR_ARG
] >= 0)
2130 fixed_regs
[cfun
->machine
->args
.reg
[DISPATCH_PTR_ARG
]] = 1;
2131 fixed_regs
[cfun
->machine
->args
.reg
[DISPATCH_PTR_ARG
] + 1] = 1;
2133 if (cfun
->machine
->args
.reg
[WORKGROUP_ID_X_ARG
] >= 0)
2134 fixed_regs
[cfun
->machine
->args
.reg
[WORKGROUP_ID_X_ARG
]] = 1;
2135 if (cfun
->machine
->args
.reg
[WORK_ITEM_ID_X_ARG
] >= 0)
2136 fixed_regs
[cfun
->machine
->args
.reg
[WORK_ITEM_ID_X_ARG
]] = 1;
2137 if (cfun
->machine
->args
.reg
[WORK_ITEM_ID_Y_ARG
] >= 0)
2138 fixed_regs
[cfun
->machine
->args
.reg
[WORK_ITEM_ID_Y_ARG
]] = 1;
2139 if (cfun
->machine
->args
.reg
[WORK_ITEM_ID_Z_ARG
] >= 0)
2140 fixed_regs
[cfun
->machine
->args
.reg
[WORK_ITEM_ID_Z_ARG
]] = 1;
2143 /* Determine if a load or store is valid, according to the register classes
2144 and address space. Used primarily by the machine description to decide
2145 when to split a move into two steps. */
2148 gcn_valid_move_p (machine_mode mode
, rtx dest
, rtx src
)
2150 if (!MEM_P (dest
) && !MEM_P (src
))
2154 && AS_FLAT_P (MEM_ADDR_SPACE (dest
))
2155 && (gcn_flat_address_p (XEXP (dest
, 0), mode
)
2156 || GET_CODE (XEXP (dest
, 0)) == SYMBOL_REF
2157 || GET_CODE (XEXP (dest
, 0)) == LABEL_REF
)
2158 && gcn_vgpr_register_operand (src
, mode
))
2160 else if (MEM_P (src
)
2161 && AS_FLAT_P (MEM_ADDR_SPACE (src
))
2162 && (gcn_flat_address_p (XEXP (src
, 0), mode
)
2163 || GET_CODE (XEXP (src
, 0)) == SYMBOL_REF
2164 || GET_CODE (XEXP (src
, 0)) == LABEL_REF
)
2165 && gcn_vgpr_register_operand (dest
, mode
))
2169 && AS_GLOBAL_P (MEM_ADDR_SPACE (dest
))
2170 && (gcn_global_address_p (XEXP (dest
, 0))
2171 || GET_CODE (XEXP (dest
, 0)) == SYMBOL_REF
2172 || GET_CODE (XEXP (dest
, 0)) == LABEL_REF
)
2173 && gcn_vgpr_register_operand (src
, mode
))
2175 else if (MEM_P (src
)
2176 && AS_GLOBAL_P (MEM_ADDR_SPACE (src
))
2177 && (gcn_global_address_p (XEXP (src
, 0))
2178 || GET_CODE (XEXP (src
, 0)) == SYMBOL_REF
2179 || GET_CODE (XEXP (src
, 0)) == LABEL_REF
)
2180 && gcn_vgpr_register_operand (dest
, mode
))
2184 && MEM_ADDR_SPACE (dest
) == ADDR_SPACE_SCALAR_FLAT
2185 && (gcn_scalar_flat_address_p (XEXP (dest
, 0))
2186 || GET_CODE (XEXP (dest
, 0)) == SYMBOL_REF
2187 || GET_CODE (XEXP (dest
, 0)) == LABEL_REF
)
2188 && gcn_ssrc_register_operand (src
, mode
))
2190 else if (MEM_P (src
)
2191 && MEM_ADDR_SPACE (src
) == ADDR_SPACE_SCALAR_FLAT
2192 && (gcn_scalar_flat_address_p (XEXP (src
, 0))
2193 || GET_CODE (XEXP (src
, 0)) == SYMBOL_REF
2194 || GET_CODE (XEXP (src
, 0)) == LABEL_REF
)
2195 && gcn_sdst_register_operand (dest
, mode
))
2199 && AS_ANY_DS_P (MEM_ADDR_SPACE (dest
))
2200 && gcn_ds_address_p (XEXP (dest
, 0))
2201 && gcn_vgpr_register_operand (src
, mode
))
2203 else if (MEM_P (src
)
2204 && AS_ANY_DS_P (MEM_ADDR_SPACE (src
))
2205 && gcn_ds_address_p (XEXP (src
, 0))
2206 && gcn_vgpr_register_operand (dest
, mode
))
2213 /* {{{ Functions and ABI. */
2215 /* Implement TARGET_FUNCTION_VALUE.
2217 Define how to find the value returned by a function.
2218 The register location is always the same, but the mode depends on
2222 gcn_function_value (const_tree valtype
, const_tree
, bool)
2224 machine_mode mode
= TYPE_MODE (valtype
);
2226 if (INTEGRAL_TYPE_P (valtype
)
2227 && GET_MODE_CLASS (mode
) == MODE_INT
2228 && GET_MODE_SIZE (mode
) < 4)
2231 return gen_rtx_REG (mode
, SGPR_REGNO (RETURN_VALUE_REG
));
2234 /* Implement TARGET_FUNCTION_VALUE_REGNO_P.
2236 Return true if N is a possible register number for the function return
2240 gcn_function_value_regno_p (const unsigned int n
)
2242 return n
== RETURN_VALUE_REG
;
2245 /* Calculate the number of registers required to hold function argument
2249 num_arg_regs (const function_arg_info
&arg
)
2251 if (targetm
.calls
.must_pass_in_stack (arg
))
2254 int size
= arg
.promoted_size_in_bytes ();
2255 return (size
+ UNITS_PER_WORD
- 1) / UNITS_PER_WORD
;
2258 /* Implement TARGET_STRICT_ARGUMENT_NAMING.
2260 Return true if the location where a function argument is passed
2261 depends on whether or not it is a named argument
2263 For gcn, we know how to handle functions declared as stdarg: by
2264 passing an extra pointer to the unnamed arguments. However, the
2265 Fortran frontend can produce a different situation, where a
2266 function pointer is declared with no arguments, but the actual
2267 function and calls to it take more arguments. In that case, we
2268 want to ensure the call matches the definition of the function. */
2271 gcn_strict_argument_naming (cumulative_args_t cum_v
)
2273 CUMULATIVE_ARGS
*cum
= get_cumulative_args (cum_v
);
2275 return cum
->fntype
== NULL_TREE
|| stdarg_p (cum
->fntype
);
2278 /* Implement TARGET_PRETEND_OUTGOING_VARARGS_NAMED.
2280 See comment on gcn_strict_argument_naming. */
2283 gcn_pretend_outgoing_varargs_named (cumulative_args_t cum_v
)
2285 return !gcn_strict_argument_naming (cum_v
);
2288 /* Implement TARGET_FUNCTION_ARG.
2290 Return an RTX indicating whether a function argument is passed in a register
2291 and if so, which register. */
2294 gcn_function_arg (cumulative_args_t cum_v
, const function_arg_info
&arg
)
2296 CUMULATIVE_ARGS
*cum
= get_cumulative_args (cum_v
);
2297 if (cum
->normal_function
)
2299 if (!arg
.named
|| arg
.end_marker_p ())
2302 if (targetm
.calls
.must_pass_in_stack (arg
))
2305 /* Vector parameters are not supported yet. */
2306 if (VECTOR_MODE_P (arg
.mode
))
2309 int reg_num
= FIRST_PARM_REG
+ cum
->num
;
2310 int num_regs
= num_arg_regs (arg
);
2312 while (reg_num
% num_regs
!= 0)
2314 if (reg_num
+ num_regs
<= FIRST_PARM_REG
+ NUM_PARM_REGS
)
2315 return gen_rtx_REG (arg
.mode
, reg_num
);
2319 if (cum
->num
>= cum
->args
.nargs
)
2321 cum
->offset
= (cum
->offset
+ TYPE_ALIGN (arg
.type
) / 8 - 1)
2322 & -(TYPE_ALIGN (arg
.type
) / 8);
2323 cfun
->machine
->kernarg_segment_alignment
2324 = MAX ((unsigned) cfun
->machine
->kernarg_segment_alignment
,
2325 TYPE_ALIGN (arg
.type
) / 8);
2326 rtx addr
= gen_rtx_REG (DImode
,
2327 cum
->args
.reg
[KERNARG_SEGMENT_PTR_ARG
]);
2329 addr
= gen_rtx_PLUS (DImode
, addr
,
2330 gen_int_mode (cum
->offset
, DImode
));
2331 rtx mem
= gen_rtx_MEM (arg
.mode
, addr
);
2332 set_mem_attributes (mem
, arg
.type
, 1);
2333 set_mem_addr_space (mem
, ADDR_SPACE_SCALAR_FLAT
);
2334 MEM_READONLY_P (mem
) = 1;
2338 int a
= cum
->args
.order
[cum
->num
];
2339 if (arg
.mode
!= gcn_kernel_arg_types
[a
].mode
)
2341 error ("wrong type of argument %s", gcn_kernel_arg_types
[a
].name
);
2344 return gen_rtx_REG ((machine_mode
) gcn_kernel_arg_types
[a
].mode
,
2350 /* Implement TARGET_FUNCTION_ARG_ADVANCE.
2352 Updates the summarizer variable pointed to by CUM_V to advance past an
2353 argument in the argument list. */
2356 gcn_function_arg_advance (cumulative_args_t cum_v
,
2357 const function_arg_info
&arg
)
2359 CUMULATIVE_ARGS
*cum
= get_cumulative_args (cum_v
);
2361 if (cum
->normal_function
)
2366 int num_regs
= num_arg_regs (arg
);
2368 while ((FIRST_PARM_REG
+ cum
->num
) % num_regs
!= 0)
2370 cum
->num
+= num_regs
;
2374 if (cum
->num
< cum
->args
.nargs
)
2378 cum
->offset
+= tree_to_uhwi (TYPE_SIZE_UNIT (arg
.type
));
2379 cfun
->machine
->kernarg_segment_byte_size
= cum
->offset
;
2384 /* Implement TARGET_ARG_PARTIAL_BYTES.
2386 Returns the number of bytes at the beginning of an argument that must be put
2387 in registers. The value must be zero for arguments that are passed entirely
2388 in registers or that are entirely pushed on the stack. */
2391 gcn_arg_partial_bytes (cumulative_args_t cum_v
, const function_arg_info
&arg
)
2393 CUMULATIVE_ARGS
*cum
= get_cumulative_args (cum_v
);
2398 if (targetm
.calls
.must_pass_in_stack (arg
))
2401 if (cum
->num
>= NUM_PARM_REGS
)
2404 /* If the argument fits entirely in registers, return 0. */
2405 if (cum
->num
+ num_arg_regs (arg
) <= NUM_PARM_REGS
)
2408 return (NUM_PARM_REGS
- cum
->num
) * UNITS_PER_WORD
;
2411 /* A normal function which takes a pointer argument (to a scalar) may be
2412 passed a pointer to LDS space (via a high-bits-set aperture), and that only
2413 works with FLAT addressing, not GLOBAL. Force FLAT addressing if the
2414 function has an incoming pointer-to-scalar parameter. */
2417 gcn_detect_incoming_pointer_arg (tree fndecl
)
2419 gcc_assert (cfun
&& cfun
->machine
);
2421 for (tree arg
= TYPE_ARG_TYPES (TREE_TYPE (fndecl
));
2423 arg
= TREE_CHAIN (arg
))
2424 if (POINTER_TYPE_P (TREE_VALUE (arg
))
2425 && !AGGREGATE_TYPE_P (TREE_TYPE (TREE_VALUE (arg
))))
2426 cfun
->machine
->use_flat_addressing
= true;
2429 /* Implement INIT_CUMULATIVE_ARGS, via gcn.h.
2431 Initialize a variable CUM of type CUMULATIVE_ARGS for a call to a function
2432 whose data type is FNTYPE. For a library call, FNTYPE is 0. */
2435 gcn_init_cumulative_args (CUMULATIVE_ARGS
*cum
/* Argument info to init */ ,
2436 tree fntype
/* tree ptr for function decl */ ,
2437 rtx libname
/* SYMBOL_REF of library name or 0 */ ,
2438 tree fndecl
, int caller
)
2440 memset (cum
, 0, sizeof (*cum
));
2441 cum
->fntype
= fntype
;
2444 gcc_assert (cfun
&& cfun
->machine
);
2445 cum
->normal_function
= true;
2448 cfun
->machine
->normal_function
= true;
2449 gcn_detect_incoming_pointer_arg (fndecl
);
2455 attr
= lookup_attribute ("amdgpu_hsa_kernel", DECL_ATTRIBUTES (fndecl
));
2456 if (fndecl
&& !attr
)
2457 attr
= lookup_attribute ("amdgpu_hsa_kernel",
2458 TYPE_ATTRIBUTES (TREE_TYPE (fndecl
)));
2459 if (!attr
&& fntype
)
2460 attr
= lookup_attribute ("amdgpu_hsa_kernel", TYPE_ATTRIBUTES (fntype
));
2461 /* Handle main () as kernel, so we can run testsuite.
2462 Handle OpenACC kernels similarly to main. */
2463 if (!attr
&& !caller
&& fndecl
2464 && (MAIN_NAME_P (DECL_NAME (fndecl
))
2465 || lookup_attribute ("omp target entrypoint",
2466 DECL_ATTRIBUTES (fndecl
)) != NULL_TREE
))
2467 gcn_parse_amdgpu_hsa_kernel_attribute (&cum
->args
, NULL_TREE
);
2470 if (!attr
|| caller
)
2472 gcc_assert (cfun
&& cfun
->machine
);
2473 cum
->normal_function
= true;
2475 cfun
->machine
->normal_function
= true;
2477 gcn_parse_amdgpu_hsa_kernel_attribute
2478 (&cum
->args
, attr
? TREE_VALUE (attr
) : NULL_TREE
);
2480 cfun
->machine
->args
= cum
->args
;
2481 if (!caller
&& cfun
->machine
->normal_function
)
2482 gcn_detect_incoming_pointer_arg (fndecl
);
2488 gcn_return_in_memory (const_tree type
, const_tree
ARG_UNUSED (fntype
))
2490 machine_mode mode
= TYPE_MODE (type
);
2491 HOST_WIDE_INT size
= int_size_in_bytes (type
);
2493 if (AGGREGATE_TYPE_P (type
))
2496 /* Vector return values are not supported yet. */
2497 if (VECTOR_TYPE_P (type
))
2500 if (mode
== BLKmode
)
2503 if (size
> 2 * UNITS_PER_WORD
)
2509 /* Implement TARGET_PROMOTE_FUNCTION_MODE.
2511 Return the mode to use for outgoing function arguments. */
2514 gcn_promote_function_mode (const_tree
ARG_UNUSED (type
), machine_mode mode
,
2515 int *ARG_UNUSED (punsignedp
),
2516 const_tree
ARG_UNUSED (funtype
),
2517 int ARG_UNUSED (for_return
))
2519 if (GET_MODE_CLASS (mode
) == MODE_INT
&& GET_MODE_SIZE (mode
) < 4)
2525 /* Implement TARGET_GIMPLIFY_VA_ARG_EXPR.
2527 Derived from hppa_gimplify_va_arg_expr. The generic routine doesn't handle
2528 ARGS_GROW_DOWNWARDS. */
2531 gcn_gimplify_va_arg_expr (tree valist
, tree type
,
2532 gimple_seq
*ARG_UNUSED (pre_p
),
2533 gimple_seq
*ARG_UNUSED (post_p
))
2535 tree ptr
= build_pointer_type (type
);
2540 indirect
= pass_va_arg_by_reference (type
);
2544 ptr
= build_pointer_type (type
);
2546 valist_type
= TREE_TYPE (valist
);
2548 /* Args grow down. Not handled by generic routines. */
2550 u
= fold_convert (sizetype
, size_in_bytes (type
));
2551 u
= fold_build1 (NEGATE_EXPR
, sizetype
, u
);
2552 t
= fold_build_pointer_plus (valist
, u
);
2554 /* Align to 8 byte boundary. */
2556 u
= build_int_cst (TREE_TYPE (t
), -8);
2557 t
= build2 (BIT_AND_EXPR
, TREE_TYPE (t
), t
, u
);
2558 t
= fold_convert (valist_type
, t
);
2560 t
= build2 (MODIFY_EXPR
, valist_type
, valist
, t
);
2562 t
= fold_convert (ptr
, t
);
2563 t
= build_va_arg_indirect_ref (t
);
2566 t
= build_va_arg_indirect_ref (t
);
2571 /* Return 1 if TRAIT NAME is present in the OpenMP context's
2572 device trait set, return 0 if not present in any OpenMP context in the
2573 whole translation unit, or -1 if not present in the current OpenMP context
2574 but might be present in another OpenMP context in the same TU. */
2577 gcn_omp_device_kind_arch_isa (enum omp_device_kind_arch_isa trait
,
2582 case omp_device_kind
:
2583 return strcmp (name
, "gpu") == 0;
2584 case omp_device_arch
:
2585 return strcmp (name
, "gcn") == 0;
2586 case omp_device_isa
:
2587 if (strcmp (name
, "fiji") == 0)
2588 return gcn_arch
== PROCESSOR_FIJI
;
2589 if (strcmp (name
, "gfx900") == 0)
2590 return gcn_arch
== PROCESSOR_VEGA10
;
2591 if (strcmp (name
, "gfx906") == 0)
2592 return gcn_arch
== PROCESSOR_VEGA20
;
2593 if (strcmp (name
, "gfx908") == 0)
2594 return gcn_arch
== PROCESSOR_GFX908
;
2601 /* Calculate stack offsets needed to create prologues and epilogues. */
2603 static struct machine_function
*
2604 gcn_compute_frame_offsets (void)
2606 machine_function
*offsets
= cfun
->machine
;
2608 if (reload_completed
)
2611 offsets
->need_frame_pointer
= frame_pointer_needed
;
2613 offsets
->outgoing_args_size
= crtl
->outgoing_args_size
;
2614 offsets
->pretend_size
= crtl
->args
.pretend_args_size
;
2616 offsets
->local_vars
= get_frame_size ();
2618 offsets
->lr_needs_saving
= (!leaf_function_p ()
2619 || df_regs_ever_live_p (LR_REGNUM
)
2620 || df_regs_ever_live_p (LR_REGNUM
+ 1));
2622 offsets
->callee_saves
= offsets
->lr_needs_saving
? 8 : 0;
2624 for (int regno
= 0; regno
< FIRST_PSEUDO_REGISTER
; regno
++)
2625 if ((df_regs_ever_live_p (regno
) && !call_used_or_fixed_reg_p (regno
))
2626 || ((regno
& ~1) == HARD_FRAME_POINTER_REGNUM
2627 && frame_pointer_needed
))
2628 offsets
->callee_saves
+= (VGPR_REGNO_P (regno
) ? 256 : 4);
2630 /* Round up to 64-bit boundary to maintain stack alignment. */
2631 offsets
->callee_saves
= (offsets
->callee_saves
+ 7) & ~7;
2636 /* Insert code into the prologue or epilogue to store or load any
2637 callee-save register to/from the stack.
2639 Helper function for gcn_expand_prologue and gcn_expand_epilogue. */
2642 move_callee_saved_registers (rtx sp
, machine_function
*offsets
,
2645 int regno
, offset
, saved_scalars
;
2646 rtx exec
= gen_rtx_REG (DImode
, EXEC_REG
);
2647 rtx vcc
= gen_rtx_REG (DImode
, VCC_LO_REG
);
2648 rtx offreg
= gen_rtx_REG (SImode
, SGPR_REGNO (22));
2649 rtx as
= gen_rtx_CONST_INT (VOIDmode
, STACK_ADDR_SPACE
);
2650 HOST_WIDE_INT exec_set
= 0;
2655 /* Move scalars into two vector registers. */
2656 for (regno
= 0, saved_scalars
= 0; regno
< FIRST_VGPR_REG
; regno
++)
2657 if ((df_regs_ever_live_p (regno
) && !call_used_or_fixed_reg_p (regno
))
2658 || ((regno
& ~1) == LINK_REGNUM
&& offsets
->lr_needs_saving
)
2659 || ((regno
& ~1) == HARD_FRAME_POINTER_REGNUM
2660 && offsets
->need_frame_pointer
))
2662 rtx reg
= gen_rtx_REG (SImode
, regno
);
2663 rtx vreg
= gen_rtx_REG (V64SImode
,
2664 VGPR_REGNO (6 + (saved_scalars
/ 64)));
2665 int lane
= saved_scalars
% 64;
2668 emit_insn (gen_vec_setv64si (vreg
, reg
, GEN_INT (lane
)));
2670 emit_insn (gen_vec_extractv64sisi (reg
, vreg
, GEN_INT (lane
)));
2675 rtx move_scalars
= get_insns ();
2679 /* Ensure that all vector lanes are moved. */
2681 emit_move_insn (exec
, GEN_INT (exec_set
));
2683 /* Set up a vector stack pointer. */
2684 rtx _0_1_2_3
= gen_rtx_REG (V64SImode
, VGPR_REGNO (1));
2685 rtx _0_4_8_12
= gen_rtx_REG (V64SImode
, VGPR_REGNO (3));
2686 emit_insn (gen_ashlv64si3_exec (_0_4_8_12
, _0_1_2_3
, GEN_INT (2),
2687 gcn_gen_undef (V64SImode
), exec
));
2688 rtx vsp
= gen_rtx_REG (V64DImode
, VGPR_REGNO (4));
2689 emit_insn (gen_vec_duplicatev64di_exec (vsp
, sp
, gcn_gen_undef (V64DImode
),
2691 emit_insn (gen_addv64si3_vcc_exec (gcn_operand_part (V64SImode
, vsp
, 0),
2692 gcn_operand_part (V64SImode
, vsp
, 0),
2693 _0_4_8_12
, vcc
, gcn_gen_undef (V64SImode
),
2695 emit_insn (gen_addcv64si3_exec (gcn_operand_part (V64SImode
, vsp
, 1),
2696 gcn_operand_part (V64SImode
, vsp
, 1),
2697 const0_rtx
, vcc
, vcc
,
2698 gcn_gen_undef (V64SImode
), exec
));
2701 for (regno
= FIRST_VGPR_REG
, offset
= offsets
->pretend_size
;
2702 regno
< FIRST_PSEUDO_REGISTER
; regno
++)
2703 if ((df_regs_ever_live_p (regno
) && !call_used_or_fixed_reg_p (regno
))
2704 || (regno
== VGPR_REGNO (6) && saved_scalars
> 0)
2705 || (regno
== VGPR_REGNO (7) && saved_scalars
> 63))
2707 rtx reg
= gen_rtx_REG (V64SImode
, regno
);
2710 if (regno
== VGPR_REGNO (6) && saved_scalars
< 64)
2711 size
= saved_scalars
* 4;
2712 else if (regno
== VGPR_REGNO (7) && saved_scalars
< 128)
2713 size
= (saved_scalars
- 64) * 4;
2715 if (size
!= 256 || exec_set
!= -1)
2717 exec_set
= ((unsigned HOST_WIDE_INT
) 1 << (size
/ 4)) - 1;
2718 emit_move_insn (exec
, gen_int_mode (exec_set
, DImode
));
2722 emit_insn (gen_scatterv64si_insn_1offset_exec (vsp
, const0_rtx
, reg
,
2723 as
, const0_rtx
, exec
));
2725 emit_insn (gen_gatherv64si_insn_1offset_exec
2726 (reg
, vsp
, const0_rtx
, as
, const0_rtx
,
2727 gcn_gen_undef (V64SImode
), exec
));
2729 /* Move our VSP to the next stack entry. */
2730 if (offreg_set
!= size
)
2733 emit_move_insn (offreg
, GEN_INT (size
));
2738 emit_move_insn (exec
, GEN_INT (exec_set
));
2740 emit_insn (gen_addv64si3_vcc_dup_exec
2741 (gcn_operand_part (V64SImode
, vsp
, 0),
2742 offreg
, gcn_operand_part (V64SImode
, vsp
, 0),
2743 vcc
, gcn_gen_undef (V64SImode
), exec
));
2744 emit_insn (gen_addcv64si3_exec
2745 (gcn_operand_part (V64SImode
, vsp
, 1),
2746 gcn_operand_part (V64SImode
, vsp
, 1),
2747 const0_rtx
, vcc
, vcc
, gcn_gen_undef (V64SImode
), exec
));
2752 rtx move_vectors
= get_insns ();
2757 emit_insn (move_scalars
);
2758 emit_insn (move_vectors
);
2762 emit_insn (move_vectors
);
2763 emit_insn (move_scalars
);
2767 /* Generate prologue. Called from gen_prologue during pro_and_epilogue pass.
2769 For a non-kernel function, the stack layout looks like this (interim),
2773 |__________________| <-- current SP
2775 |__________________|
2777 |__________________|
2779 |__________________| <-- FP/hard FP
2780 | callee-save regs |
2781 |__________________| <-- soft arg pointer
2783 |__________________| <-- incoming SP
2785 lo |..................|
2787 This implies arguments (beyond the first N in registers) must grow
2788 downwards (as, apparently, PA has them do).
2790 For a kernel function we have the simpler:
2793 |__________________| <-- current SP
2795 |__________________|
2797 |__________________|
2799 lo |__________________| <-- FP/hard FP
2804 gcn_expand_prologue ()
2806 machine_function
*offsets
= gcn_compute_frame_offsets ();
2808 if (!cfun
|| !cfun
->machine
|| cfun
->machine
->normal_function
)
2810 rtx sp
= gen_rtx_REG (Pmode
, STACK_POINTER_REGNUM
);
2811 rtx sp_hi
= gcn_operand_part (Pmode
, sp
, 1);
2812 rtx sp_lo
= gcn_operand_part (Pmode
, sp
, 0);
2813 rtx fp
= gen_rtx_REG (Pmode
, HARD_FRAME_POINTER_REGNUM
);
2814 rtx fp_hi
= gcn_operand_part (Pmode
, fp
, 1);
2815 rtx fp_lo
= gcn_operand_part (Pmode
, fp
, 0);
2819 if (offsets
->pretend_size
> 0)
2821 /* FIXME: Do the actual saving of register pretend args to the stack.
2822 Register order needs consideration. */
2825 /* Save callee-save regs. */
2826 move_callee_saved_registers (sp
, offsets
, true);
2828 HOST_WIDE_INT sp_adjust
= offsets
->pretend_size
2829 + offsets
->callee_saves
2830 + offsets
->local_vars
+ offsets
->outgoing_args_size
;
2833 /* Adding RTX_FRAME_RELATED_P effectively disables spliting, so
2834 we use split add explictly, and specify the DImode add in
2836 rtx scc
= gen_rtx_REG (BImode
, SCC_REG
);
2837 rtx adjustment
= gen_int_mode (sp_adjust
, SImode
);
2838 rtx insn
= emit_insn (gen_addsi3_scalar_carry (sp_lo
, sp_lo
,
2840 RTX_FRAME_RELATED_P (insn
) = 1;
2841 add_reg_note (insn
, REG_FRAME_RELATED_EXPR
,
2843 gen_rtx_PLUS (DImode
, sp
, adjustment
)));
2844 emit_insn (gen_addcsi3_scalar_zero (sp_hi
, sp_hi
, scc
));
2847 if (offsets
->need_frame_pointer
)
2849 /* Adding RTX_FRAME_RELATED_P effectively disables spliting, so
2850 we use split add explictly, and specify the DImode add in
2852 rtx scc
= gen_rtx_REG (BImode
, SCC_REG
);
2853 int fp_adjust
= -(offsets
->local_vars
+ offsets
->outgoing_args_size
);
2854 rtx adjustment
= gen_int_mode (fp_adjust
, SImode
);
2855 rtx insn
= emit_insn (gen_addsi3_scalar_carry(fp_lo
, sp_lo
,
2857 RTX_FRAME_RELATED_P (insn
) = 1;
2858 add_reg_note (insn
, REG_FRAME_RELATED_EXPR
,
2860 gen_rtx_PLUS (DImode
, sp
, adjustment
)));
2861 emit_insn (gen_addcsi3_scalar (fp_hi
, sp_hi
,
2862 (fp_adjust
< 0 ? GEN_INT (-1)
2867 rtx_insn
*seq
= get_insns ();
2870 /* FIXME: Prologue insns should have this flag set for debug output, etc.
2871 but it causes issues for now.
2872 for (insn = seq; insn; insn = NEXT_INSN (insn))
2874 RTX_FRAME_RELATED_P (insn) = 1;*/
2880 rtx wave_offset
= gen_rtx_REG (SImode
,
2881 cfun
->machine
->args
.
2882 reg
[PRIVATE_SEGMENT_WAVE_OFFSET_ARG
]);
2884 if (cfun
->machine
->args
.requested
& (1 << FLAT_SCRATCH_INIT_ARG
))
2887 gen_rtx_REG (SImode
,
2888 cfun
->machine
->args
.reg
[FLAT_SCRATCH_INIT_ARG
]);
2890 gen_rtx_REG (SImode
,
2891 cfun
->machine
->args
.reg
[FLAT_SCRATCH_INIT_ARG
] + 1);
2892 rtx fs_reg_lo
= gen_rtx_REG (SImode
, FLAT_SCRATCH_REG
);
2893 rtx fs_reg_hi
= gen_rtx_REG (SImode
, FLAT_SCRATCH_REG
+ 1);
2895 /*rtx queue = gen_rtx_REG(DImode,
2896 cfun->machine->args.reg[QUEUE_PTR_ARG]);
2897 rtx aperture = gen_rtx_MEM (SImode,
2898 gen_rtx_PLUS (DImode, queue,
2899 gen_int_mode (68, SImode)));
2900 set_mem_addr_space (aperture, ADDR_SPACE_SCALAR_FLAT);*/
2902 /* Set up flat_scratch. */
2903 emit_insn (gen_addsi3_scc (fs_reg_hi
, fs_init_lo
, wave_offset
));
2904 emit_insn (gen_lshrsi3_scc (fs_reg_hi
, fs_reg_hi
,
2905 gen_int_mode (8, SImode
)));
2906 emit_move_insn (fs_reg_lo
, fs_init_hi
);
2909 /* Set up frame pointer and stack pointer. */
2910 rtx sp
= gen_rtx_REG (DImode
, STACK_POINTER_REGNUM
);
2911 rtx sp_hi
= simplify_gen_subreg (SImode
, sp
, DImode
, 4);
2912 rtx sp_lo
= simplify_gen_subreg (SImode
, sp
, DImode
, 0);
2913 rtx fp
= gen_rtx_REG (DImode
, HARD_FRAME_POINTER_REGNUM
);
2914 rtx fp_hi
= simplify_gen_subreg (SImode
, fp
, DImode
, 4);
2915 rtx fp_lo
= simplify_gen_subreg (SImode
, fp
, DImode
, 0);
2917 HOST_WIDE_INT sp_adjust
= (offsets
->local_vars
2918 + offsets
->outgoing_args_size
);
2920 /* Initialise FP and SP from the buffer descriptor in s[0:3]. */
2921 emit_move_insn (fp_lo
, gen_rtx_REG (SImode
, 0));
2922 emit_insn (gen_andsi3_scc (fp_hi
, gen_rtx_REG (SImode
, 1),
2923 gen_int_mode (0xffff, SImode
)));
2924 rtx scc
= gen_rtx_REG (BImode
, SCC_REG
);
2925 emit_insn (gen_addsi3_scalar_carry (fp_lo
, fp_lo
, wave_offset
, scc
));
2926 emit_insn (gen_addcsi3_scalar_zero (fp_hi
, fp_hi
, scc
));
2928 /* Adding RTX_FRAME_RELATED_P effectively disables spliting, so we use
2929 split add explictly, and specify the DImode add in the note.
2930 The DWARF info expects that the callee-save data is in the frame,
2931 even though it isn't (because this is the entry point), so we
2932 make a notional adjustment to the DWARF frame offset here. */
2933 rtx dbg_adjustment
= gen_int_mode (sp_adjust
+ offsets
->callee_saves
,
2938 rtx scc
= gen_rtx_REG (BImode
, SCC_REG
);
2939 rtx adjustment
= gen_int_mode (sp_adjust
, DImode
);
2940 insn
= emit_insn (gen_addsi3_scalar_carry(sp_lo
, fp_lo
, adjustment
,
2942 emit_insn (gen_addcsi3_scalar_zero (sp_hi
, fp_hi
, scc
));
2945 insn
= emit_move_insn (sp
, fp
);
2946 RTX_FRAME_RELATED_P (insn
) = 1;
2947 add_reg_note (insn
, REG_FRAME_RELATED_EXPR
,
2948 gen_rtx_SET (sp
, gen_rtx_PLUS (DImode
, sp
,
2951 /* Make sure the flat scratch reg doesn't get optimised away. */
2952 emit_insn (gen_prologue_use (gen_rtx_REG (DImode
, FLAT_SCRATCH_REG
)));
2955 /* Ensure that the scheduler doesn't do anything unexpected. */
2956 emit_insn (gen_blockage ());
2958 /* m0 is initialized for the usual LDS DS and FLAT memory case.
2959 The low-part is the address of the topmost addressable byte, which is
2960 size-1. The high-part is an offset and should be zero. */
2961 emit_move_insn (gen_rtx_REG (SImode
, M0_REG
),
2962 gen_int_mode (LDS_SIZE
-1, SImode
));
2964 emit_insn (gen_prologue_use (gen_rtx_REG (SImode
, M0_REG
)));
2966 if (cfun
&& cfun
->machine
&& !cfun
->machine
->normal_function
&& flag_openmp
)
2968 /* OpenMP kernels have an implicit call to gomp_gcn_enter_kernel. */
2969 rtx fn_reg
= gen_rtx_REG (Pmode
, FIRST_PARM_REG
);
2970 emit_move_insn (fn_reg
, gen_rtx_SYMBOL_REF (Pmode
,
2971 "gomp_gcn_enter_kernel"));
2972 emit_call_insn (gen_gcn_indirect_call (fn_reg
, const0_rtx
));
2976 /* Generate epilogue. Called from gen_epilogue during pro_and_epilogue pass.
2978 See gcn_expand_prologue for stack details. */
2981 gcn_expand_epilogue (void)
2983 /* Ensure that the scheduler doesn't do anything unexpected. */
2984 emit_insn (gen_blockage ());
2986 if (!cfun
|| !cfun
->machine
|| cfun
->machine
->normal_function
)
2988 machine_function
*offsets
= gcn_compute_frame_offsets ();
2989 rtx sp
= gen_rtx_REG (Pmode
, STACK_POINTER_REGNUM
);
2990 rtx fp
= gen_rtx_REG (Pmode
, HARD_FRAME_POINTER_REGNUM
);
2992 HOST_WIDE_INT sp_adjust
= offsets
->callee_saves
+ offsets
->pretend_size
;
2994 if (offsets
->need_frame_pointer
)
2996 /* Restore old SP from the frame pointer. */
2998 emit_insn (gen_subdi3 (sp
, fp
, gen_int_mode (sp_adjust
, DImode
)));
3000 emit_move_insn (sp
, fp
);
3004 /* Restore old SP from current SP. */
3005 sp_adjust
+= offsets
->outgoing_args_size
+ offsets
->local_vars
;
3008 emit_insn (gen_subdi3 (sp
, sp
, gen_int_mode (sp_adjust
, DImode
)));
3011 move_callee_saved_registers (sp
, offsets
, false);
3013 /* There's no explicit use of the link register on the return insn. Emit
3014 one here instead. */
3015 if (offsets
->lr_needs_saving
)
3016 emit_use (gen_rtx_REG (DImode
, LINK_REGNUM
));
3018 /* Similar for frame pointer. */
3019 if (offsets
->need_frame_pointer
)
3020 emit_use (gen_rtx_REG (DImode
, HARD_FRAME_POINTER_REGNUM
));
3022 else if (flag_openmp
)
3024 /* OpenMP kernels have an implicit call to gomp_gcn_exit_kernel. */
3025 rtx fn_reg
= gen_rtx_REG (Pmode
, FIRST_PARM_REG
);
3026 emit_move_insn (fn_reg
,
3027 gen_rtx_SYMBOL_REF (Pmode
, "gomp_gcn_exit_kernel"));
3028 emit_call_insn (gen_gcn_indirect_call (fn_reg
, const0_rtx
));
3030 else if (TREE_CODE (TREE_TYPE (DECL_RESULT (cfun
->decl
))) != VOID_TYPE
)
3032 /* Assume that an exit value compatible with gcn-run is expected.
3033 That is, the third input parameter is an int*.
3035 We can't allocate any new registers, but the kernarg_reg is
3036 dead after this, so we'll use that. */
3037 rtx kernarg_reg
= gen_rtx_REG (DImode
, cfun
->machine
->args
.reg
3038 [KERNARG_SEGMENT_PTR_ARG
]);
3039 rtx retptr_mem
= gen_rtx_MEM (DImode
,
3040 gen_rtx_PLUS (DImode
, kernarg_reg
,
3042 set_mem_addr_space (retptr_mem
, ADDR_SPACE_SCALAR_FLAT
);
3043 emit_move_insn (kernarg_reg
, retptr_mem
);
3045 rtx retval_mem
= gen_rtx_MEM (SImode
, kernarg_reg
);
3046 set_mem_addr_space (retval_mem
, ADDR_SPACE_SCALAR_FLAT
);
3047 emit_move_insn (retval_mem
,
3048 gen_rtx_REG (SImode
, SGPR_REGNO (RETURN_VALUE_REG
)));
3051 emit_jump_insn (gen_gcn_return ());
3054 /* Implement TARGET_CAN_ELIMINATE.
3056 Return true if the compiler is allowed to try to replace register number
3057 FROM_REG with register number TO_REG.
3059 FIXME: is the default "true" not enough? Should this be a negative set? */
3062 gcn_can_eliminate_p (int /*from_reg */ , int to_reg
)
3064 return (to_reg
== HARD_FRAME_POINTER_REGNUM
3065 || to_reg
== STACK_POINTER_REGNUM
);
3068 /* Implement INITIAL_ELIMINATION_OFFSET.
3070 Returns the initial difference between the specified pair of registers, in
3071 terms of stack position. */
3074 gcn_initial_elimination_offset (int from
, int to
)
3076 machine_function
*offsets
= gcn_compute_frame_offsets ();
3080 case ARG_POINTER_REGNUM
:
3081 if (to
== STACK_POINTER_REGNUM
)
3082 return -(offsets
->callee_saves
+ offsets
->local_vars
3083 + offsets
->outgoing_args_size
);
3084 else if (to
== FRAME_POINTER_REGNUM
|| to
== HARD_FRAME_POINTER_REGNUM
)
3085 return -offsets
->callee_saves
;
3090 case FRAME_POINTER_REGNUM
:
3091 if (to
== STACK_POINTER_REGNUM
)
3092 return -(offsets
->local_vars
+ offsets
->outgoing_args_size
);
3093 else if (to
== HARD_FRAME_POINTER_REGNUM
)
3104 /* Implement HARD_REGNO_RENAME_OK.
3106 Return true if it is permissible to rename a hard register from
3107 FROM_REG to TO_REG. */
3110 gcn_hard_regno_rename_ok (unsigned int from_reg
, unsigned int to_reg
)
3112 if (from_reg
== SCC_REG
3113 || from_reg
== VCC_LO_REG
|| from_reg
== VCC_HI_REG
3114 || from_reg
== EXEC_LO_REG
|| from_reg
== EXEC_HI_REG
3115 || to_reg
== SCC_REG
3116 || to_reg
== VCC_LO_REG
|| to_reg
== VCC_HI_REG
3117 || to_reg
== EXEC_LO_REG
|| to_reg
== EXEC_HI_REG
)
3120 /* Allow the link register to be used if it was saved. */
3121 if ((to_reg
& ~1) == LINK_REGNUM
)
3122 return !cfun
|| cfun
->machine
->lr_needs_saving
;
3124 /* Allow the registers used for the static chain to be used if the chain is
3125 not in active use. */
3126 if ((to_reg
& ~1) == STATIC_CHAIN_REGNUM
)
3128 || !(cfun
->static_chain_decl
3129 && df_regs_ever_live_p (STATIC_CHAIN_REGNUM
)
3130 && df_regs_ever_live_p (STATIC_CHAIN_REGNUM
+ 1));
3135 /* Implement HARD_REGNO_CALLER_SAVE_MODE.
3137 Which mode is required for saving NREGS of a pseudo-register in
3138 call-clobbered hard register REGNO. */
3141 gcn_hard_regno_caller_save_mode (unsigned int regno
, unsigned int nregs
,
3142 machine_mode regmode
)
3144 machine_mode result
= choose_hard_reg_mode (regno
, nregs
, NULL
);
3146 if (VECTOR_MODE_P (result
) && !VECTOR_MODE_P (regmode
))
3147 result
= (nregs
== 1 ? SImode
: DImode
);
3152 /* Implement TARGET_ASM_TRAMPOLINE_TEMPLATE.
3154 Output assembler code for a block containing the constant parts
3155 of a trampoline, leaving space for the variable parts. */
3158 gcn_asm_trampoline_template (FILE *f
)
3160 /* The source operand of the move instructions must be a 32-bit
3161 constant following the opcode. */
3162 asm_fprintf (f
, "\ts_mov_b32\ts%i, 0xffff\n", STATIC_CHAIN_REGNUM
);
3163 asm_fprintf (f
, "\ts_mov_b32\ts%i, 0xffff\n", STATIC_CHAIN_REGNUM
+ 1);
3164 asm_fprintf (f
, "\ts_mov_b32\ts%i, 0xffff\n", CC_SAVE_REG
);
3165 asm_fprintf (f
, "\ts_mov_b32\ts%i, 0xffff\n", CC_SAVE_REG
+ 1);
3166 asm_fprintf (f
, "\ts_setpc_b64\ts[%i:%i]\n", CC_SAVE_REG
, CC_SAVE_REG
+ 1);
3169 /* Implement TARGET_TRAMPOLINE_INIT.
3171 Emit RTL insns to initialize the variable parts of a trampoline.
3172 FNDECL is the decl of the target address, M_TRAMP is a MEM for
3173 the trampoline, and CHAIN_VALUE is an RTX for the static chain
3174 to be passed to the target function. */
3177 gcn_trampoline_init (rtx m_tramp
, tree fndecl
, rtx chain_value
)
3179 if (TARGET_GCN5_PLUS
)
3180 sorry ("nested function trampolines not supported on GCN5 due to"
3181 " non-executable stacks");
3183 emit_block_move (m_tramp
, assemble_trampoline_template (),
3184 GEN_INT (TRAMPOLINE_SIZE
), BLOCK_OP_NORMAL
);
3186 rtx fnaddr
= XEXP (DECL_RTL (fndecl
), 0);
3187 rtx chain_value_reg
= copy_to_reg (chain_value
);
3188 rtx fnaddr_reg
= copy_to_reg (fnaddr
);
3190 for (int i
= 0; i
< 4; i
++)
3192 rtx mem
= adjust_address (m_tramp
, SImode
, i
* 8 + 4);
3193 rtx reg
= i
< 2 ? chain_value_reg
: fnaddr_reg
;
3194 emit_move_insn (mem
, gen_rtx_SUBREG (SImode
, reg
, (i
% 2) * 4));
3197 rtx tramp_addr
= XEXP (m_tramp
, 0);
3198 emit_insn (gen_clear_icache (tramp_addr
,
3199 plus_constant (ptr_mode
, tramp_addr
,
3204 /* {{{ Miscellaneous. */
3206 /* Implement TARGET_CANNOT_COPY_INSN_P.
3208 Return true if INSN must not be duplicated. */
3211 gcn_cannot_copy_insn_p (rtx_insn
*insn
)
3213 if (recog_memoized (insn
) == CODE_FOR_gcn_wavefront_barrier
)
3219 /* Implement TARGET_DEBUG_UNWIND_INFO.
3221 Defines the mechanism that will be used for describing frame unwind
3222 information to the debugger. */
3224 static enum unwind_info_type
3225 gcn_debug_unwind_info ()
3227 /* No support for debug info, yet. */
3231 /* Determine if there is a suitable hardware conversion instruction.
3232 Used primarily by the machine description. */
3235 gcn_valid_cvt_p (machine_mode from
, machine_mode to
, enum gcn_cvt_t op
)
3237 if (VECTOR_MODE_P (from
) != VECTOR_MODE_P (to
))
3240 if (VECTOR_MODE_P (from
))
3242 from
= GET_MODE_INNER (from
);
3243 to
= GET_MODE_INNER (to
);
3249 case fixuns_trunc_cvt
:
3250 if (GET_MODE_CLASS (from
) != MODE_FLOAT
3251 || GET_MODE_CLASS (to
) != MODE_INT
)
3256 if (GET_MODE_CLASS (from
) != MODE_INT
3257 || GET_MODE_CLASS (to
) != MODE_FLOAT
)
3261 if (GET_MODE_CLASS (from
) != MODE_FLOAT
3262 || GET_MODE_CLASS (to
) != MODE_FLOAT
3263 || GET_MODE_SIZE (from
) >= GET_MODE_SIZE (to
))
3267 if (GET_MODE_CLASS (from
) != MODE_FLOAT
3268 || GET_MODE_CLASS (to
) != MODE_FLOAT
3269 || GET_MODE_SIZE (from
) <= GET_MODE_SIZE (to
))
3274 return ((to
== HImode
&& from
== HFmode
)
3275 || (to
== SImode
&& (from
== SFmode
|| from
== DFmode
))
3276 || (to
== HFmode
&& (from
== HImode
|| from
== SFmode
))
3277 || (to
== SFmode
&& (from
== SImode
|| from
== HFmode
3279 || (to
== DFmode
&& (from
== SImode
|| from
== SFmode
)));
3282 /* Implement TARGET_EMUTLS_VAR_INIT.
3284 Disable emutls (gthr-gcn.h does not support it, yet). */
3287 gcn_emutls_var_init (tree
, tree decl
, tree
)
3289 sorry_at (DECL_SOURCE_LOCATION (decl
), "TLS is not implemented for GCN.");
3296 /* Implement TARGET_RTX_COSTS.
3298 Compute a (partial) cost for rtx X. Return true if the complete
3299 cost has been computed, and false if subexpressions should be
3300 scanned. In either case, *TOTAL contains the cost result. */
3303 gcn_rtx_costs (rtx x
, machine_mode
, int, int, int *total
, bool)
3305 enum rtx_code code
= GET_CODE (x
);
3312 if (gcn_inline_constant_p (x
))
3314 else if (code
== CONST_INT
3315 && ((unsigned HOST_WIDE_INT
) INTVAL (x
) + 0x8000) < 0x10000)
3317 else if (gcn_constant_p (x
))
3320 *total
= vgpr_vector_mode_p (GET_MODE (x
)) ? 64 : 4;
3333 /* Implement TARGET_MEMORY_MOVE_COST.
3335 Return the cost of moving data of mode M between a
3336 register and memory. A value of 2 is the default; this cost is
3337 relative to those in `REGISTER_MOVE_COST'.
3339 This function is used extensively by register_move_cost that is used to
3340 build tables at startup. Make it inline in this case.
3341 When IN is 2, return maximum of in and out move cost.
3343 If moving between registers and memory is more expensive than
3344 between two registers, you should define this macro to express the
3347 Model also increased moving costs of QImode registers in non
3350 #define LOAD_COST 32
3351 #define STORE_COST 32
3353 gcn_memory_move_cost (machine_mode mode
, reg_class_t regclass
, bool in
)
3355 int nregs
= CEIL (GET_MODE_SIZE (mode
), 4);
3358 case SCC_CONDITIONAL_REG
:
3359 case VCCZ_CONDITIONAL_REG
:
3360 case VCC_CONDITIONAL_REG
:
3361 case EXECZ_CONDITIONAL_REG
:
3362 case ALL_CONDITIONAL_REGS
:
3364 case SGPR_EXEC_REGS
:
3366 case SGPR_VOP_SRC_REGS
:
3367 case SGPR_MEM_SRC_REGS
:
3373 return (STORE_COST
+ 2) * nregs
;
3374 return LOAD_COST
* nregs
;
3377 return (LOAD_COST
+ 2) * nregs
;
3378 return STORE_COST
* nregs
;
3383 return (LOAD_COST
+ 2) * nregs
;
3384 return (STORE_COST
+ 2) * nregs
;
3390 /* Implement TARGET_REGISTER_MOVE_COST.
3392 Return the cost of moving data from a register in class CLASS1 to
3393 one in class CLASS2. Base value is 2. */
3396 gcn_register_move_cost (machine_mode
, reg_class_t dst
, reg_class_t src
)
3398 /* Increase cost of moving from and to vector registers. While this is
3399 fast in hardware (I think), it has hidden cost of setting up the exec
3401 if ((src
< VGPR_REGS
) != (dst
< VGPR_REGS
))
3409 /* Type codes used by GCN built-in definitions. */
3411 enum gcn_builtin_type_index
3413 GCN_BTI_END_OF_PARAMS
,
3432 GCN_BTI_LDS_VOIDPTR
,
3437 static GTY(()) tree gcn_builtin_types
[GCN_BTI_MAX
];
3439 #define exec_type_node (gcn_builtin_types[GCN_BTI_EXEC])
3440 #define sf_type_node (gcn_builtin_types[GCN_BTI_SF])
3441 #define v64si_type_node (gcn_builtin_types[GCN_BTI_V64SI])
3442 #define v64sf_type_node (gcn_builtin_types[GCN_BTI_V64SF])
3443 #define v64ptr_type_node (gcn_builtin_types[GCN_BTI_V64PTR])
3444 #define siptr_type_node (gcn_builtin_types[GCN_BTI_SIPTR])
3445 #define sfptr_type_node (gcn_builtin_types[GCN_BTI_SFPTR])
3446 #define voidptr_type_node (gcn_builtin_types[GCN_BTI_VOIDPTR])
3447 #define size_t_type_node (gcn_builtin_types[GCN_BTI_SIZE_T])
3449 static rtx
gcn_expand_builtin_1 (tree
, rtx
, rtx
, machine_mode
, int,
3450 struct gcn_builtin_description
*);
3451 static rtx
gcn_expand_builtin_binop (tree
, rtx
, rtx
, machine_mode
, int,
3452 struct gcn_builtin_description
*);
3454 struct gcn_builtin_description
;
3455 typedef rtx (*gcn_builtin_expander
) (tree
, rtx
, rtx
, machine_mode
, int,
3456 struct gcn_builtin_description
*);
3458 enum gcn_builtin_type
3460 B_UNIMPLEMENTED
, /* Sorry out */
3461 B_INSN
, /* Emit a pattern */
3462 B_OVERLOAD
/* Placeholder for an overloaded function */
3465 struct gcn_builtin_description
3470 enum gcn_builtin_type type
;
3471 /* The first element of parm is always the return type. The rest
3472 are a zero terminated list of parameters. */
3474 gcn_builtin_expander expander
;
3477 /* Read in the GCN builtins from gcn-builtins.def. */
3479 extern GTY(()) struct gcn_builtin_description gcn_builtins
[GCN_BUILTIN_MAX
];
3481 struct gcn_builtin_description gcn_builtins
[] = {
3482 #define DEF_BUILTIN(fcode, icode, name, type, params, expander) \
3483 {GCN_BUILTIN_ ## fcode, icode, name, type, params, expander},
3485 #define DEF_BUILTIN_BINOP_INT_FP(fcode, ic, name) \
3486 {GCN_BUILTIN_ ## fcode ## _V64SI, \
3487 CODE_FOR_ ## ic ##v64si3_exec, name "_v64int", B_INSN, \
3488 {GCN_BTI_V64SI, GCN_BTI_EXEC, GCN_BTI_V64SI, GCN_BTI_V64SI, \
3489 GCN_BTI_V64SI, GCN_BTI_END_OF_PARAMS}, gcn_expand_builtin_binop}, \
3490 {GCN_BUILTIN_ ## fcode ## _V64SI_unspec, \
3491 CODE_FOR_ ## ic ##v64si3_exec, name "_v64int_unspec", B_INSN, \
3492 {GCN_BTI_V64SI, GCN_BTI_EXEC, GCN_BTI_V64SI, GCN_BTI_V64SI, \
3493 GCN_BTI_END_OF_PARAMS}, gcn_expand_builtin_binop},
3495 #include "gcn-builtins.def"
3496 #undef DEF_BUILTIN_BINOP_INT_FP
3500 static GTY(()) tree gcn_builtin_decls
[GCN_BUILTIN_MAX
];
3502 /* Implement TARGET_BUILTIN_DECL.
3504 Return the GCN builtin for CODE. */
3507 gcn_builtin_decl (unsigned code
, bool ARG_UNUSED (initialize_p
))
3509 if (code
>= GCN_BUILTIN_MAX
)
3510 return error_mark_node
;
3512 return gcn_builtin_decls
[code
];
3515 /* Helper function for gcn_init_builtins. */
3518 gcn_init_builtin_types (void)
3520 gcn_builtin_types
[GCN_BTI_VOID
] = void_type_node
;
3521 gcn_builtin_types
[GCN_BTI_BOOL
] = boolean_type_node
;
3522 gcn_builtin_types
[GCN_BTI_INT
] = intSI_type_node
;
3523 gcn_builtin_types
[GCN_BTI_UINT
] = unsigned_type_for (intSI_type_node
);
3524 gcn_builtin_types
[GCN_BTI_SIZE_T
] = size_type_node
;
3525 gcn_builtin_types
[GCN_BTI_LLINT
] = intDI_type_node
;
3526 gcn_builtin_types
[GCN_BTI_LLUINT
] = unsigned_type_for (intDI_type_node
);
3528 exec_type_node
= unsigned_intDI_type_node
;
3529 sf_type_node
= float32_type_node
;
3530 v64si_type_node
= build_vector_type (intSI_type_node
, 64);
3531 v64sf_type_node
= build_vector_type (float_type_node
, 64);
3532 v64ptr_type_node
= build_vector_type (unsigned_intDI_type_node
3533 /*build_pointer_type
3534 (integer_type_node) */
3536 tree tmp
= build_distinct_type_copy (intSI_type_node
);
3537 TYPE_ADDR_SPACE (tmp
) = ADDR_SPACE_FLAT
;
3538 siptr_type_node
= build_pointer_type (tmp
);
3540 tmp
= build_distinct_type_copy (float_type_node
);
3541 TYPE_ADDR_SPACE (tmp
) = ADDR_SPACE_FLAT
;
3542 sfptr_type_node
= build_pointer_type (tmp
);
3544 tmp
= build_distinct_type_copy (void_type_node
);
3545 TYPE_ADDR_SPACE (tmp
) = ADDR_SPACE_FLAT
;
3546 voidptr_type_node
= build_pointer_type (tmp
);
3548 tmp
= build_distinct_type_copy (void_type_node
);
3549 TYPE_ADDR_SPACE (tmp
) = ADDR_SPACE_LDS
;
3550 gcn_builtin_types
[GCN_BTI_LDS_VOIDPTR
] = build_pointer_type (tmp
);
3553 /* Implement TARGET_INIT_BUILTINS.
3555 Set up all builtin functions for this target. */
3558 gcn_init_builtins (void)
3560 gcn_init_builtin_types ();
3562 struct gcn_builtin_description
*d
;
3564 for (i
= 0, d
= gcn_builtins
; i
< GCN_BUILTIN_MAX
; i
++, d
++)
3567 char name
[64]; /* build_function will make a copy. */
3570 /* FIXME: Is this necessary/useful? */
3574 /* Find last parm. */
3575 for (parm
= 1; d
->parm
[parm
] != GCN_BTI_END_OF_PARAMS
; parm
++)
3580 p
= tree_cons (NULL_TREE
, gcn_builtin_types
[d
->parm
[--parm
]], p
);
3582 p
= build_function_type (gcn_builtin_types
[d
->parm
[0]], p
);
3584 sprintf (name
, "__builtin_gcn_%s", d
->name
);
3585 gcn_builtin_decls
[i
]
3586 = add_builtin_function (name
, p
, i
, BUILT_IN_MD
, NULL
, NULL_TREE
);
3588 /* These builtins don't throw. */
3589 TREE_NOTHROW (gcn_builtin_decls
[i
]) = 1;
3592 /* FIXME: remove the ifdef once OpenACC support is merged upstream. */
3593 #ifdef BUILT_IN_GOACC_SINGLE_START
3594 /* These builtins need to take/return an LDS pointer: override the generic
3597 set_builtin_decl (BUILT_IN_GOACC_SINGLE_START
,
3598 gcn_builtin_decls
[GCN_BUILTIN_ACC_SINGLE_START
], false);
3600 set_builtin_decl (BUILT_IN_GOACC_SINGLE_COPY_START
,
3601 gcn_builtin_decls
[GCN_BUILTIN_ACC_SINGLE_COPY_START
],
3604 set_builtin_decl (BUILT_IN_GOACC_SINGLE_COPY_END
,
3605 gcn_builtin_decls
[GCN_BUILTIN_ACC_SINGLE_COPY_END
],
3608 set_builtin_decl (BUILT_IN_GOACC_BARRIER
,
3609 gcn_builtin_decls
[GCN_BUILTIN_ACC_BARRIER
], false);
3613 /* Expand the CMP_SWAP GCN builtins. We have our own versions that do
3614 not require taking the address of any object, other than the memory
3615 cell being operated on.
3617 Helper function for gcn_expand_builtin_1. */
3620 gcn_expand_cmp_swap (tree exp
, rtx target
)
3622 machine_mode mode
= TYPE_MODE (TREE_TYPE (exp
));
3624 = TYPE_ADDR_SPACE (TREE_TYPE (TREE_TYPE (CALL_EXPR_ARG (exp
, 0))));
3625 machine_mode as_mode
= gcn_addr_space_address_mode (as
);
3628 target
= gen_reg_rtx (mode
);
3630 rtx addr
= expand_expr (CALL_EXPR_ARG (exp
, 0),
3631 NULL_RTX
, as_mode
, EXPAND_NORMAL
);
3632 rtx cmp
= expand_expr (CALL_EXPR_ARG (exp
, 1),
3633 NULL_RTX
, mode
, EXPAND_NORMAL
);
3634 rtx src
= expand_expr (CALL_EXPR_ARG (exp
, 2),
3635 NULL_RTX
, mode
, EXPAND_NORMAL
);
3638 rtx mem
= gen_rtx_MEM (mode
, force_reg (as_mode
, addr
));
3639 set_mem_addr_space (mem
, as
);
3642 cmp
= copy_to_mode_reg (mode
, cmp
);
3644 src
= copy_to_mode_reg (mode
, src
);
3647 pat
= gen_sync_compare_and_swapsi (target
, mem
, cmp
, src
);
3649 pat
= gen_sync_compare_and_swapdi (target
, mem
, cmp
, src
);
3656 /* Expand many different builtins.
3658 Intended for use in gcn-builtins.def. */
3661 gcn_expand_builtin_1 (tree exp
, rtx target
, rtx
/*subtarget */ ,
3662 machine_mode
/*mode */ , int ignore
,
3663 struct gcn_builtin_description
*)
3665 tree fndecl
= TREE_OPERAND (CALL_EXPR_FN (exp
), 0);
3666 switch (DECL_MD_FUNCTION_CODE (fndecl
))
3668 case GCN_BUILTIN_FLAT_LOAD_INT32
:
3674 expand_expr (CALL_EXPR_ARG (exp
, 0), NULL_RTX
, DImode
,
3677 force_reg (V64DImode
,
3678 expand_expr (CALL_EXPR_ARG (exp
, 1), NULL_RTX
, V64DImode
,
3680 /*emit_insn (gen_vector_flat_loadv64si
3681 (target, gcn_gen_undef (V64SImode), ptr, exec)); */
3684 case GCN_BUILTIN_FLAT_LOAD_PTR_INT32
:
3685 case GCN_BUILTIN_FLAT_LOAD_PTR_FLOAT
:
3689 rtx exec
= force_reg (DImode
,
3690 expand_expr (CALL_EXPR_ARG (exp
, 0), NULL_RTX
,
3693 rtx ptr
= force_reg (DImode
,
3694 expand_expr (CALL_EXPR_ARG (exp
, 1), NULL_RTX
,
3697 rtx offsets
= force_reg (V64SImode
,
3698 expand_expr (CALL_EXPR_ARG (exp
, 2),
3699 NULL_RTX
, V64DImode
,
3701 rtx addrs
= gen_reg_rtx (V64DImode
);
3702 rtx tmp
= gen_reg_rtx (V64SImode
);
3703 emit_insn (gen_ashlv64si3_exec (tmp
, offsets
,
3705 gcn_gen_undef (V64SImode
), exec
));
3706 emit_insn (gen_addv64di3_zext_dup2_exec (addrs
, tmp
, ptr
,
3707 gcn_gen_undef (V64DImode
),
3709 rtx mem
= gen_rtx_MEM (GET_MODE (target
), addrs
);
3710 /*set_mem_addr_space (mem, ADDR_SPACE_FLAT); */
3711 /* FIXME: set attributes. */
3712 emit_insn (gen_mov_with_exec (target
, mem
, exec
));
3715 case GCN_BUILTIN_FLAT_STORE_PTR_INT32
:
3716 case GCN_BUILTIN_FLAT_STORE_PTR_FLOAT
:
3718 rtx exec
= force_reg (DImode
,
3719 expand_expr (CALL_EXPR_ARG (exp
, 0), NULL_RTX
,
3722 rtx ptr
= force_reg (DImode
,
3723 expand_expr (CALL_EXPR_ARG (exp
, 1), NULL_RTX
,
3726 rtx offsets
= force_reg (V64SImode
,
3727 expand_expr (CALL_EXPR_ARG (exp
, 2),
3728 NULL_RTX
, V64DImode
,
3730 machine_mode vmode
= TYPE_MODE (TREE_TYPE (CALL_EXPR_ARG (exp
,
3732 rtx val
= force_reg (vmode
,
3733 expand_expr (CALL_EXPR_ARG (exp
, 3), NULL_RTX
,
3736 rtx addrs
= gen_reg_rtx (V64DImode
);
3737 rtx tmp
= gen_reg_rtx (V64SImode
);
3738 emit_insn (gen_ashlv64si3_exec (tmp
, offsets
,
3740 gcn_gen_undef (V64SImode
), exec
));
3741 emit_insn (gen_addv64di3_zext_dup2_exec (addrs
, tmp
, ptr
,
3742 gcn_gen_undef (V64DImode
),
3744 rtx mem
= gen_rtx_MEM (vmode
, addrs
);
3745 /*set_mem_addr_space (mem, ADDR_SPACE_FLAT); */
3746 /* FIXME: set attributes. */
3747 emit_insn (gen_mov_with_exec (mem
, val
, exec
));
3750 case GCN_BUILTIN_SQRTVF
:
3754 rtx exec
= gcn_full_exec_reg ();
3755 rtx arg
= force_reg (V64SFmode
,
3756 expand_expr (CALL_EXPR_ARG (exp
, 0), NULL_RTX
,
3759 emit_insn (gen_sqrtv64sf2_exec
3760 (target
, arg
, gcn_gen_undef (V64SFmode
), exec
));
3763 case GCN_BUILTIN_SQRTF
:
3767 rtx arg
= force_reg (SFmode
,
3768 expand_expr (CALL_EXPR_ARG (exp
, 0), NULL_RTX
,
3771 emit_insn (gen_sqrtsf2 (target
, arg
));
3774 case GCN_BUILTIN_OMP_DIM_SIZE
:
3778 emit_insn (gen_oacc_dim_size (target
,
3779 expand_expr (CALL_EXPR_ARG (exp
, 0),
3784 case GCN_BUILTIN_OMP_DIM_POS
:
3788 emit_insn (gen_oacc_dim_pos (target
,
3789 expand_expr (CALL_EXPR_ARG (exp
, 0),
3794 case GCN_BUILTIN_CMP_SWAP
:
3795 case GCN_BUILTIN_CMP_SWAPLL
:
3796 return gcn_expand_cmp_swap (exp
, target
);
3798 case GCN_BUILTIN_ACC_SINGLE_START
:
3803 rtx wavefront
= gcn_oacc_dim_pos (1);
3804 rtx cond
= gen_rtx_EQ (VOIDmode
, wavefront
, const0_rtx
);
3805 rtx cc
= (target
&& REG_P (target
)) ? target
: gen_reg_rtx (BImode
);
3806 emit_insn (gen_cstoresi4 (cc
, cond
, wavefront
, const0_rtx
));
3810 case GCN_BUILTIN_ACC_SINGLE_COPY_START
:
3812 rtx blk
= force_reg (SImode
,
3813 expand_expr (CALL_EXPR_ARG (exp
, 0), NULL_RTX
,
3814 SImode
, EXPAND_NORMAL
));
3815 rtx wavefront
= gcn_oacc_dim_pos (1);
3816 rtx cond
= gen_rtx_NE (VOIDmode
, wavefront
, const0_rtx
);
3817 rtx not_zero
= gen_label_rtx ();
3818 emit_insn (gen_cbranchsi4 (cond
, wavefront
, const0_rtx
, not_zero
));
3819 emit_move_insn (blk
, const0_rtx
);
3820 emit_label (not_zero
);
3824 case GCN_BUILTIN_ACC_SINGLE_COPY_END
:
3827 case GCN_BUILTIN_ACC_BARRIER
:
3828 emit_insn (gen_gcn_wavefront_barrier ());
3836 /* Expansion of simple arithmetic and bit binary operation builtins.
3838 Intended for use with gcn_builtins table. */
3841 gcn_expand_builtin_binop (tree exp
, rtx target
, rtx
/*subtarget */ ,
3842 machine_mode
/*mode */ , int ignore
,
3843 struct gcn_builtin_description
*d
)
3845 int icode
= d
->icode
;
3849 rtx exec
= force_reg (DImode
,
3850 expand_expr (CALL_EXPR_ARG (exp
, 0), NULL_RTX
, DImode
,
3853 machine_mode m1
= insn_data
[icode
].operand
[1].mode
;
3854 rtx arg1
= expand_expr (CALL_EXPR_ARG (exp
, 1), NULL_RTX
, m1
,
3856 if (!insn_data
[icode
].operand
[1].predicate (arg1
, m1
))
3857 arg1
= force_reg (m1
, arg1
);
3859 machine_mode m2
= insn_data
[icode
].operand
[2].mode
;
3860 rtx arg2
= expand_expr (CALL_EXPR_ARG (exp
, 2), NULL_RTX
, m2
,
3862 if (!insn_data
[icode
].operand
[2].predicate (arg2
, m2
))
3863 arg2
= force_reg (m2
, arg2
);
3866 if (call_expr_nargs (exp
) == 4)
3868 machine_mode m_prev
= insn_data
[icode
].operand
[4].mode
;
3869 arg_prev
= force_reg (m_prev
,
3870 expand_expr (CALL_EXPR_ARG (exp
, 3), NULL_RTX
,
3871 m_prev
, EXPAND_NORMAL
));
3874 arg_prev
= gcn_gen_undef (GET_MODE (target
));
3876 rtx pat
= GEN_FCN (icode
) (target
, arg1
, arg2
, exec
, arg_prev
);
3881 /* Implement TARGET_EXPAND_BUILTIN.
3883 Expand an expression EXP that calls a built-in function, with result going
3884 to TARGET if that's convenient (and in mode MODE if that's convenient).
3885 SUBTARGET may be used as the target for computing one of EXP's operands.
3886 IGNORE is nonzero if the value is to be ignored. */
3889 gcn_expand_builtin (tree exp
, rtx target
, rtx subtarget
, machine_mode mode
,
3892 tree fndecl
= TREE_OPERAND (CALL_EXPR_FN (exp
), 0);
3893 unsigned int fcode
= DECL_MD_FUNCTION_CODE (fndecl
);
3894 struct gcn_builtin_description
*d
;
3896 gcc_assert (fcode
< GCN_BUILTIN_MAX
);
3897 d
= &gcn_builtins
[fcode
];
3899 if (d
->type
== B_UNIMPLEMENTED
)
3900 sorry ("Builtin not implemented");
3902 return d
->expander (exp
, target
, subtarget
, mode
, ignore
, d
);
3906 /* {{{ Vectorization. */
3908 /* Implement TARGET_VECTORIZE_GET_MASK_MODE.
3910 A vector mask is a value that holds one boolean result for every element in
3914 gcn_vectorize_get_mask_mode (machine_mode
)
3916 /* GCN uses a DImode bit-mask. */
3920 /* Return an RTX that references a vector with the i-th lane containing
3923 Helper function for gcn_vectorize_vec_perm_const. */
3926 gcn_make_vec_perm_address (unsigned int *perm
)
3928 rtx x
= gen_reg_rtx (V64SImode
);
3929 emit_move_insn (x
, gcn_vec_constant (V64SImode
, 0));
3931 /* Permutation addresses use byte addressing. With each vector lane being
3932 4 bytes wide, and with 64 lanes in total, only bits 2..7 are significant,
3935 The permutation given to the vec_perm* patterns range from 0 to 2N-1 to
3936 select between lanes in two vectors, but as the DS_BPERMUTE* instructions
3937 only take one source vector, the most-significant bit can be ignored
3938 here. Instead, we can use EXEC masking to select the relevant part of
3939 each source vector after they are permuted separately. */
3940 uint64_t bit_mask
= 1 << 2;
3941 for (int i
= 2; i
< 8; i
++, bit_mask
<<= 1)
3943 uint64_t exec_mask
= 0;
3944 uint64_t lane_mask
= 1;
3945 for (int j
= 0; j
< 64; j
++, lane_mask
<<= 1)
3946 if ((perm
[j
] * 4) & bit_mask
)
3947 exec_mask
|= lane_mask
;
3950 emit_insn (gen_addv64si3_exec (x
, x
,
3951 gcn_vec_constant (V64SImode
,
3953 x
, get_exec (exec_mask
)));
3959 /* Implement TARGET_VECTORIZE_VEC_PERM_CONST.
3961 Return true if permutation with SEL is possible.
3963 If DST/SRC0/SRC1 are non-null, emit the instructions to perform the
3967 gcn_vectorize_vec_perm_const (machine_mode vmode
, rtx dst
,
3969 const vec_perm_indices
& sel
)
3971 unsigned int nelt
= GET_MODE_NUNITS (vmode
);
3973 gcc_assert (VECTOR_MODE_P (vmode
));
3974 gcc_assert (nelt
<= 64);
3975 gcc_assert (sel
.length () == nelt
);
3979 /* All vector permutations are possible on this architecture,
3980 with varying degrees of efficiency depending on the permutation. */
3984 unsigned int perm
[64];
3985 for (unsigned int i
= 0; i
< nelt
; ++i
)
3986 perm
[i
] = sel
[i
] & (2 * nelt
- 1);
3987 for (unsigned int i
= nelt
; i
< 64; ++i
)
3990 src0
= force_reg (vmode
, src0
);
3991 src1
= force_reg (vmode
, src1
);
3993 /* Make life a bit easier by swapping operands if necessary so that
3994 the first element always comes from src0. */
3995 if (perm
[0] >= nelt
)
3997 std::swap (src0
, src1
);
3999 for (unsigned int i
= 0; i
< nelt
; ++i
)
4006 /* TODO: There are more efficient ways to implement certain permutations
4007 using ds_swizzle_b32 and/or DPP. Test for and expand them here, before
4008 this more inefficient generic approach is used. */
4010 int64_t src1_lanes
= 0;
4011 int64_t lane_bit
= 1;
4013 for (unsigned int i
= 0; i
< nelt
; ++i
, lane_bit
<<= 1)
4015 /* Set the bits for lanes from src1. */
4016 if (perm
[i
] >= nelt
)
4017 src1_lanes
|= lane_bit
;
4020 rtx addr
= gcn_make_vec_perm_address (perm
);
4021 rtx (*ds_bpermute
) (rtx
, rtx
, rtx
, rtx
);
4026 ds_bpermute
= gen_ds_bpermutev64qi
;
4029 ds_bpermute
= gen_ds_bpermutev64hi
;
4032 ds_bpermute
= gen_ds_bpermutev64si
;
4035 ds_bpermute
= gen_ds_bpermutev64hf
;
4038 ds_bpermute
= gen_ds_bpermutev64sf
;
4041 ds_bpermute
= gen_ds_bpermutev64di
;
4044 ds_bpermute
= gen_ds_bpermutev64df
;
4050 /* Load elements from src0 to dst. */
4051 gcc_assert (~src1_lanes
);
4052 emit_insn (ds_bpermute (dst
, addr
, src0
, gcn_full_exec_reg ()));
4054 /* Load elements from src1 to dst. */
4057 /* Masking a lane masks both the destination and source lanes for
4058 DS_BPERMUTE, so we need to have all lanes enabled for the permute,
4059 then add an extra masked move to merge the results of permuting
4060 the two source vectors together.
4062 rtx tmp
= gen_reg_rtx (vmode
);
4063 emit_insn (ds_bpermute (tmp
, addr
, src1
, gcn_full_exec_reg ()));
4064 emit_insn (gen_mov_with_exec (dst
, tmp
, get_exec (src1_lanes
)));
4070 /* Implements TARGET_VECTOR_MODE_SUPPORTED_P.
4072 Return nonzero if vector MODE is supported with at least move
4076 gcn_vector_mode_supported_p (machine_mode mode
)
4078 return (mode
== V64QImode
|| mode
== V64HImode
4079 || mode
== V64SImode
|| mode
== V64DImode
4080 || mode
== V64SFmode
|| mode
== V64DFmode
);
4083 /* Implement TARGET_VECTORIZE_PREFERRED_SIMD_MODE.
4085 Enables autovectorization for all supported modes. */
4088 gcn_vectorize_preferred_simd_mode (scalar_mode mode
)
4109 /* Implement TARGET_VECTORIZE_RELATED_MODE.
4111 All GCN vectors are 64-lane, so this is simpler than other architectures.
4112 In particular, we do *not* want to match vector bit-size. */
4114 static opt_machine_mode
4115 gcn_related_vector_mode (machine_mode
ARG_UNUSED (vector_mode
),
4116 scalar_mode element_mode
, poly_uint64 nunits
)
4118 if (known_ne (nunits
, 0U) && known_ne (nunits
, 64U))
4121 machine_mode pref_mode
= gcn_vectorize_preferred_simd_mode (element_mode
);
4122 if (!VECTOR_MODE_P (pref_mode
))
4128 /* Implement TARGET_VECTORIZE_PREFERRED_VECTOR_ALIGNMENT.
4130 Returns the preferred alignment in bits for accesses to vectors of type type
4131 in vectorized code. This might be less than or greater than the ABI-defined
4132 value returned by TARGET_VECTOR_ALIGNMENT. It can be equal to the alignment
4133 of a single element, in which case the vectorizer will not try to optimize
4137 gcn_preferred_vector_alignment (const_tree type
)
4139 return TYPE_ALIGN (TREE_TYPE (type
));
4142 /* Implement TARGET_VECTORIZE_SUPPORT_VECTOR_MISALIGNMENT.
4144 Return true if the target supports misaligned vector store/load of a
4145 specific factor denoted in the misalignment parameter. */
4148 gcn_vectorize_support_vector_misalignment (machine_mode
ARG_UNUSED (mode
),
4149 const_tree type
, int misalignment
,
4155 /* If the misalignment is unknown, we should be able to handle the access
4156 so long as it is not to a member of a packed data structure. */
4157 if (misalignment
== -1)
4160 /* Return true if the misalignment is a multiple of the natural alignment
4161 of the vector's element type. This is probably always going to be
4162 true in practice, since we've already established that this isn't a
4164 return misalignment
% TYPE_ALIGN_UNIT (type
) == 0;
4167 /* Implement TARGET_VECTORIZE_VECTOR_ALIGNMENT_REACHABLE.
4169 Return true if vector alignment is reachable (by peeling N iterations) for
4170 the given scalar type TYPE. */
4173 gcn_vector_alignment_reachable (const_tree
ARG_UNUSED (type
), bool is_packed
)
4175 /* Vectors which aren't in packed structures will not be less aligned than
4176 the natural alignment of their element type, so this is safe. */
4180 /* Generate DPP instructions used for vector reductions.
4182 The opcode is given by INSN.
4183 The first operand of the operation is shifted right by SHIFT vector lanes.
4184 SHIFT must be a power of 2. If SHIFT is 16, the 15th lane of each row is
4185 broadcast the next row (thereby acting like a shift of 16 for the end of
4186 each row). If SHIFT is 32, lane 31 is broadcast to all the
4187 following lanes (thereby acting like a shift of 32 for lane 63). */
4190 gcn_expand_dpp_shr_insn (machine_mode mode
, const char *insn
,
4191 int unspec
, int shift
)
4193 static char buf
[128];
4195 const char *vcc_in
= "";
4196 const char *vcc_out
= "";
4198 /* Add the vcc operand if needed. */
4199 if (GET_MODE_CLASS (mode
) == MODE_VECTOR_INT
)
4201 if (unspec
== UNSPEC_PLUS_CARRY_IN_DPP_SHR
)
4204 if (unspec
== UNSPEC_PLUS_CARRY_DPP_SHR
4205 || unspec
== UNSPEC_PLUS_CARRY_IN_DPP_SHR
)
4209 /* Add the DPP modifiers. */
4213 dpp
= "row_shr:1 bound_ctrl:0";
4216 dpp
= "row_shr:2 bound_ctrl:0";
4219 dpp
= "row_shr:4 bank_mask:0xe";
4222 dpp
= "row_shr:8 bank_mask:0xc";
4225 dpp
= "row_bcast:15 row_mask:0xa";
4228 dpp
= "row_bcast:31 row_mask:0xc";
4234 if (unspec
== UNSPEC_MOV_DPP_SHR
&& vgpr_2reg_mode_p (mode
))
4235 sprintf (buf
, "%s\t%%L0, %%L1 %s\n\t%s\t%%H0, %%H1 %s",
4236 insn
, dpp
, insn
, dpp
);
4237 else if (unspec
== UNSPEC_MOV_DPP_SHR
)
4238 sprintf (buf
, "%s\t%%0, %%1 %s", insn
, dpp
);
4240 sprintf (buf
, "%s\t%%0%s, %%1, %%2%s %s", insn
, vcc_out
, vcc_in
, dpp
);
4245 /* Generate vector reductions in terms of DPP instructions.
4247 The vector register SRC of mode MODE is reduced using the operation given
4248 by UNSPEC, and the scalar result is returned in lane 63 of a vector
4252 gcn_expand_reduc_scalar (machine_mode mode
, rtx src
, int unspec
)
4254 machine_mode orig_mode
= mode
;
4255 bool use_moves
= (((unspec
== UNSPEC_SMIN_DPP_SHR
4256 || unspec
== UNSPEC_SMAX_DPP_SHR
4257 || unspec
== UNSPEC_UMIN_DPP_SHR
4258 || unspec
== UNSPEC_UMAX_DPP_SHR
)
4259 && (mode
== V64DImode
4260 || mode
== V64DFmode
))
4261 || (unspec
== UNSPEC_PLUS_DPP_SHR
4262 && mode
== V64DFmode
));
4263 rtx_code code
= (unspec
== UNSPEC_SMIN_DPP_SHR
? SMIN
4264 : unspec
== UNSPEC_SMAX_DPP_SHR
? SMAX
4265 : unspec
== UNSPEC_UMIN_DPP_SHR
? UMIN
4266 : unspec
== UNSPEC_UMAX_DPP_SHR
? UMAX
4267 : unspec
== UNSPEC_PLUS_DPP_SHR
? PLUS
4269 bool use_extends
= ((unspec
== UNSPEC_SMIN_DPP_SHR
4270 || unspec
== UNSPEC_SMAX_DPP_SHR
4271 || unspec
== UNSPEC_UMIN_DPP_SHR
4272 || unspec
== UNSPEC_UMAX_DPP_SHR
)
4273 && (mode
== V64QImode
4274 || mode
== V64HImode
));
4275 bool unsignedp
= (unspec
== UNSPEC_UMIN_DPP_SHR
4276 || unspec
== UNSPEC_UMAX_DPP_SHR
);
4277 bool use_plus_carry
= unspec
== UNSPEC_PLUS_DPP_SHR
4278 && GET_MODE_CLASS (mode
) == MODE_VECTOR_INT
4279 && (TARGET_GCN3
|| mode
== V64DImode
);
4282 unspec
= UNSPEC_PLUS_CARRY_DPP_SHR
;
4286 rtx tmp
= gen_reg_rtx (V64SImode
);
4287 convert_move (tmp
, src
, unsignedp
);
4292 /* Perform reduction by first performing the reduction operation on every
4293 pair of lanes, then on every pair of results from the previous
4294 iteration (thereby effectively reducing every 4 lanes) and so on until
4295 all lanes are reduced. */
4297 for (int i
= 0, shift
= 1; i
< 6; i
++, shift
<<= 1)
4299 rtx shift_val
= gen_rtx_CONST_INT (VOIDmode
, shift
);
4301 out
= gen_reg_rtx (mode
);
4305 rtx tmp
= gen_reg_rtx (mode
);
4306 emit_insn (gen_dpp_move (mode
, tmp
, in
, shift_val
));
4307 emit_insn (gen_rtx_SET (out
, gen_rtx_fmt_ee (code
, mode
, tmp
, in
)));
4311 rtx insn
= gen_rtx_SET (out
,
4312 gen_rtx_UNSPEC (mode
,
4313 gen_rtvec (3, in
, in
,
4317 /* Add clobber for instructions that set the carry flags. */
4320 rtx clobber
= gen_rtx_CLOBBER (VOIDmode
,
4321 gen_rtx_REG (DImode
, VCC_REG
));
4322 insn
= gen_rtx_PARALLEL (VOIDmode
,
4323 gen_rtvec (2, insn
, clobber
));
4332 rtx tmp
= gen_reg_rtx (orig_mode
);
4333 convert_move (tmp
, out
, unsignedp
);
4340 /* Implement TARGET_VECTORIZE_BUILTIN_VECTORIZATION_COST. */
4343 gcn_vectorization_cost (enum vect_cost_for_stmt
ARG_UNUSED (type_of_cost
),
4344 tree
ARG_UNUSED (vectype
), int ARG_UNUSED (misalign
))
4346 /* Always vectorize. */
4351 /* {{{ md_reorg pass. */
4353 /* Identify VMEM instructions from their "type" attribute. */
4356 gcn_vmem_insn_p (attr_type type
)
4387 /* If INSN sets the EXEC register to a constant value, return the value,
4388 otherwise return zero. */
4391 gcn_insn_exec_value (rtx_insn
*insn
)
4393 if (!NONDEBUG_INSN_P (insn
))
4396 rtx pattern
= PATTERN (insn
);
4398 if (GET_CODE (pattern
) == SET
)
4400 rtx dest
= XEXP (pattern
, 0);
4401 rtx src
= XEXP (pattern
, 1);
4403 if (GET_MODE (dest
) == DImode
4404 && REG_P (dest
) && REGNO (dest
) == EXEC_REG
4405 && CONST_INT_P (src
))
4406 return INTVAL (src
);
4412 /* Sets the EXEC register before INSN to the value that it had after
4413 LAST_EXEC_DEF. The constant value of the EXEC register is returned if
4414 known, otherwise it returns zero. */
4417 gcn_restore_exec (rtx_insn
*insn
, rtx_insn
*last_exec_def
, int64_t curr_exec
,
4418 bool curr_exec_known
, bool &last_exec_def_saved
)
4420 rtx exec_reg
= gen_rtx_REG (DImode
, EXEC_REG
);
4423 int64_t exec_value
= gcn_insn_exec_value (last_exec_def
);
4427 /* If the EXEC value is a constant and it happens to be the same as the
4428 current EXEC value, the restore can be skipped. */
4429 if (curr_exec_known
&& exec_value
== curr_exec
)
4432 exec
= GEN_INT (exec_value
);
4436 /* If the EXEC value is not a constant, save it in a register after the
4437 point of definition. */
4438 rtx exec_save_reg
= gen_rtx_REG (DImode
, EXEC_SAVE_REG
);
4440 if (!last_exec_def_saved
)
4443 emit_move_insn (exec_save_reg
, exec_reg
);
4444 rtx_insn
*seq
= get_insns ();
4447 emit_insn_after (seq
, last_exec_def
);
4448 if (dump_file
&& (dump_flags
& TDF_DETAILS
))
4449 fprintf (dump_file
, "Saving EXEC after insn %d.\n",
4450 INSN_UID (last_exec_def
));
4452 last_exec_def_saved
= true;
4455 exec
= exec_save_reg
;
4458 /* Restore EXEC register before the usage. */
4460 emit_move_insn (exec_reg
, exec
);
4461 rtx_insn
*seq
= get_insns ();
4463 emit_insn_before (seq
, insn
);
4465 if (dump_file
&& (dump_flags
& TDF_DETAILS
))
4468 fprintf (dump_file
, "Restoring EXEC to %ld before insn %d.\n",
4469 exec_value
, INSN_UID (insn
));
4472 "Restoring EXEC from saved value before insn %d.\n",
4479 /* Implement TARGET_MACHINE_DEPENDENT_REORG.
4481 Ensure that pipeline dependencies and lane masking are set correctly. */
4487 rtx exec_reg
= gen_rtx_REG (DImode
, EXEC_REG
);
4490 INIT_REG_SET (&live
);
4492 compute_bb_for_insn ();
4497 if (dump_file
&& (dump_flags
& TDF_DETAILS
))
4499 fprintf (dump_file
, "After split:\n");
4500 print_rtl_with_bb (dump_file
, get_insns (), dump_flags
);
4503 /* Update data-flow information for split instructions. */
4504 df_insn_rescan_all ();
4507 df_live_add_problem ();
4508 df_live_set_all_dirty ();
4511 /* This pass ensures that the EXEC register is set correctly, according
4512 to the "exec" attribute. However, care must be taken so that the
4513 value that reaches explicit uses of the EXEC register remains the
4517 FOR_EACH_BB_FN (bb
, cfun
)
4519 if (dump_file
&& (dump_flags
& TDF_DETAILS
))
4520 fprintf (dump_file
, "BB %d:\n", bb
->index
);
4522 rtx_insn
*insn
, *curr
;
4523 rtx_insn
*last_exec_def
= BB_HEAD (bb
);
4524 bool last_exec_def_saved
= false;
4525 bool curr_exec_explicit
= true;
4526 bool curr_exec_known
= true;
4527 int64_t curr_exec
= 0; /* 0 here means 'the value is that of EXEC
4528 after last_exec_def is executed'. */
4530 bitmap live_in
= DF_LR_IN (bb
);
4531 bool exec_live_on_entry
= false;
4532 if (bitmap_bit_p (live_in
, EXEC_LO_REG
)
4533 || bitmap_bit_p (live_in
, EXEC_HI_REG
))
4536 fprintf (dump_file
, "EXEC reg is live on entry to block %d\n",
4538 exec_live_on_entry
= true;
4541 FOR_BB_INSNS_SAFE (bb
, insn
, curr
)
4543 if (!NONDEBUG_INSN_P (insn
))
4546 if (GET_CODE (PATTERN (insn
)) == USE
4547 || GET_CODE (PATTERN (insn
)) == CLOBBER
)
4550 HARD_REG_SET defs
, uses
;
4551 CLEAR_HARD_REG_SET (defs
);
4552 CLEAR_HARD_REG_SET (uses
);
4553 note_stores (insn
, record_hard_reg_sets
, &defs
);
4554 note_uses (&PATTERN (insn
), record_hard_reg_uses
, &uses
);
4556 bool exec_lo_def_p
= TEST_HARD_REG_BIT (defs
, EXEC_LO_REG
);
4557 bool exec_hi_def_p
= TEST_HARD_REG_BIT (defs
, EXEC_HI_REG
);
4558 bool exec_used
= (hard_reg_set_intersect_p
4559 (uses
, reg_class_contents
[(int) EXEC_MASK_REG
])
4560 || TEST_HARD_REG_BIT (uses
, EXECZ_REG
));
4562 /* Check the instruction for implicit setting of EXEC via an
4564 attr_exec exec_attr
= get_attr_exec (insn
);
4574 /* Instructions that do not involve memory accesses only require
4575 bit 0 of EXEC to be set. */
4576 if (gcn_vmem_insn_p (get_attr_type (insn
))
4577 || get_attr_type (insn
) == TYPE_DS
)
4580 new_exec
= curr_exec
| 1;
4587 default: /* Auto-detect what setting is appropriate. */
4591 /* If EXEC is referenced explicitly then we don't need to do
4592 anything to set it, so we're done. */
4596 /* Scan the insn for VGPRs defs or uses. The mode determines
4597 what kind of exec is needed. */
4598 subrtx_iterator::array_type array
;
4599 FOR_EACH_SUBRTX (iter
, array
, PATTERN (insn
), NONCONST
)
4601 const_rtx x
= *iter
;
4602 if (REG_P (x
) && VGPR_REGNO_P (REGNO (x
)))
4604 if (VECTOR_MODE_P (GET_MODE (x
)))
4617 if (new_exec
&& (!curr_exec_known
|| new_exec
!= curr_exec
))
4620 emit_move_insn (exec_reg
, GEN_INT (new_exec
));
4621 rtx_insn
*seq
= get_insns ();
4623 emit_insn_before (seq
, insn
);
4625 if (dump_file
&& (dump_flags
& TDF_DETAILS
))
4626 fprintf (dump_file
, "Setting EXEC to %ld before insn %d.\n",
4627 new_exec
, INSN_UID (insn
));
4629 curr_exec
= new_exec
;
4630 curr_exec_explicit
= false;
4631 curr_exec_known
= true;
4633 else if (new_exec
&& dump_file
&& (dump_flags
& TDF_DETAILS
))
4635 fprintf (dump_file
, "Exec already is %ld before insn %d.\n",
4636 new_exec
, INSN_UID (insn
));
4639 /* The state of the EXEC register is unknown after a
4642 curr_exec_known
= false;
4644 /* Handle explicit uses of EXEC. If the instruction is a partial
4645 explicit definition of EXEC, then treat it as an explicit use of
4647 if (exec_used
|| exec_lo_def_p
!= exec_hi_def_p
)
4649 /* An instruction that explicitly uses EXEC should not also
4650 implicitly define it. */
4651 gcc_assert (!exec_used
|| !new_exec
);
4653 if (!curr_exec_known
|| !curr_exec_explicit
)
4655 /* Restore the previous explicitly defined value. */
4656 curr_exec
= gcn_restore_exec (insn
, last_exec_def
,
4657 curr_exec
, curr_exec_known
,
4658 last_exec_def_saved
);
4659 curr_exec_explicit
= true;
4660 curr_exec_known
= true;
4664 /* Handle explicit definitions of EXEC. */
4665 if (exec_lo_def_p
|| exec_hi_def_p
)
4667 last_exec_def
= insn
;
4668 last_exec_def_saved
= false;
4669 curr_exec
= gcn_insn_exec_value (insn
);
4670 curr_exec_explicit
= true;
4671 curr_exec_known
= true;
4673 if (dump_file
&& (dump_flags
& TDF_DETAILS
))
4675 "Found %s definition of EXEC at insn %d.\n",
4676 exec_lo_def_p
== exec_hi_def_p
? "full" : "partial",
4680 exec_live_on_entry
= false;
4683 COPY_REG_SET (&live
, DF_LR_OUT (bb
));
4684 df_simulate_initialize_backwards (bb
, &live
);
4686 /* If EXEC is live after the basic block, restore the value of EXEC
4687 at the end of the block. */
4688 if ((REGNO_REG_SET_P (&live
, EXEC_LO_REG
)
4689 || REGNO_REG_SET_P (&live
, EXEC_HI_REG
))
4690 && (!curr_exec_known
|| !curr_exec_explicit
|| exec_live_on_entry
))
4692 rtx_insn
*end_insn
= BB_END (bb
);
4694 /* If the instruction is not a jump instruction, do the restore
4695 after the last instruction in the basic block. */
4696 if (NONJUMP_INSN_P (end_insn
))
4697 end_insn
= NEXT_INSN (end_insn
);
4699 gcn_restore_exec (end_insn
, last_exec_def
, curr_exec
,
4700 curr_exec_known
, last_exec_def_saved
);
4704 CLEAR_REG_SET (&live
);
4706 /* "Manually Inserted Wait States (NOPs)."
4708 GCN hardware detects most kinds of register dependencies, but there
4709 are some exceptions documented in the ISA manual. This pass
4710 detects the missed cases, and inserts the documented number of NOPs
4711 required for correct execution. */
4713 const int max_waits
= 5;
4718 attr_delayeduse delayeduse
;
4719 HARD_REG_SET writes
;
4724 for (int i
= 0; i
< max_waits
; i
++)
4725 back
[i
].insn
= NULL
;
4727 rtx_insn
*insn
, *last_insn
= NULL
;
4728 for (insn
= get_insns (); insn
!= 0; insn
= NEXT_INSN (insn
))
4730 if (!NONDEBUG_INSN_P (insn
))
4733 if (GET_CODE (PATTERN (insn
)) == USE
4734 || GET_CODE (PATTERN (insn
)) == CLOBBER
)
4737 attr_type itype
= get_attr_type (insn
);
4738 attr_unit iunit
= get_attr_unit (insn
);
4739 attr_delayeduse idelayeduse
= get_attr_delayeduse (insn
);
4740 HARD_REG_SET ireads
, iwrites
;
4741 CLEAR_HARD_REG_SET (ireads
);
4742 CLEAR_HARD_REG_SET (iwrites
);
4743 note_stores (insn
, record_hard_reg_sets
, &iwrites
);
4744 note_uses (&PATTERN (insn
), record_hard_reg_uses
, &ireads
);
4746 /* Scan recent previous instructions for dependencies not handled in
4749 for (int i
= oldest
; i
< oldest
+ max_waits
; i
++)
4751 struct ilist
*prev_insn
= &back
[i
% max_waits
];
4753 if (!prev_insn
->insn
)
4756 /* VALU writes SGPR followed by VMEM reading the same SGPR
4757 requires 5 wait states. */
4758 if ((prev_insn
->age
+ nops_rqd
) < 5
4759 && prev_insn
->unit
== UNIT_VECTOR
4760 && gcn_vmem_insn_p (itype
))
4762 HARD_REG_SET regs
= prev_insn
->writes
& ireads
;
4763 if (hard_reg_set_intersect_p
4764 (regs
, reg_class_contents
[(int) SGPR_REGS
]))
4765 nops_rqd
= 5 - prev_insn
->age
;
4768 /* VALU sets VCC/EXEC followed by VALU uses VCCZ/EXECZ
4769 requires 5 wait states. */
4770 if ((prev_insn
->age
+ nops_rqd
) < 5
4771 && prev_insn
->unit
== UNIT_VECTOR
4772 && iunit
== UNIT_VECTOR
4773 && ((hard_reg_set_intersect_p
4775 reg_class_contents
[(int) EXEC_MASK_REG
])
4776 && TEST_HARD_REG_BIT (ireads
, EXECZ_REG
))
4778 (hard_reg_set_intersect_p
4780 reg_class_contents
[(int) VCC_CONDITIONAL_REG
])
4781 && TEST_HARD_REG_BIT (ireads
, VCCZ_REG
))))
4782 nops_rqd
= 5 - prev_insn
->age
;
4784 /* VALU writes SGPR/VCC followed by v_{read,write}lane using
4785 SGPR/VCC as lane select requires 4 wait states. */
4786 if ((prev_insn
->age
+ nops_rqd
) < 4
4787 && prev_insn
->unit
== UNIT_VECTOR
4788 && get_attr_laneselect (insn
) == LANESELECT_YES
)
4790 HARD_REG_SET regs
= prev_insn
->writes
& ireads
;
4791 if (hard_reg_set_intersect_p
4792 (regs
, reg_class_contents
[(int) SGPR_REGS
])
4793 || hard_reg_set_intersect_p
4794 (regs
, reg_class_contents
[(int) VCC_CONDITIONAL_REG
]))
4795 nops_rqd
= 4 - prev_insn
->age
;
4798 /* VALU writes VGPR followed by VALU_DPP reading that VGPR
4799 requires 2 wait states. */
4800 if ((prev_insn
->age
+ nops_rqd
) < 2
4801 && prev_insn
->unit
== UNIT_VECTOR
4802 && itype
== TYPE_VOP_DPP
)
4804 HARD_REG_SET regs
= prev_insn
->writes
& ireads
;
4805 if (hard_reg_set_intersect_p
4806 (regs
, reg_class_contents
[(int) VGPR_REGS
]))
4807 nops_rqd
= 2 - prev_insn
->age
;
4810 /* Store that requires input registers are not overwritten by
4811 following instruction. */
4812 if ((prev_insn
->age
+ nops_rqd
) < 1
4813 && prev_insn
->delayeduse
== DELAYEDUSE_YES
4814 && ((hard_reg_set_intersect_p
4815 (prev_insn
->reads
, iwrites
))))
4816 nops_rqd
= 1 - prev_insn
->age
;
4819 /* Insert the required number of NOPs. */
4820 for (int i
= nops_rqd
; i
> 0; i
--)
4821 emit_insn_after (gen_nop (), last_insn
);
4823 /* Age the previous instructions. We can also ignore writes to
4824 registers subsequently overwritten. */
4825 HARD_REG_SET written
;
4826 CLEAR_HARD_REG_SET (written
);
4827 for (int i
= oldest
+ max_waits
- 1; i
> oldest
; i
--)
4829 struct ilist
*prev_insn
= &back
[i
% max_waits
];
4831 /* Assume all instructions are equivalent to one "wait", the same
4832 as s_nop. This is probably true for SALU, but not VALU (which
4833 may take longer), so this is not optimal. However, AMD do
4834 not publish the cycle times for instructions. */
4835 prev_insn
->age
+= 1 + nops_rqd
;
4838 prev_insn
->writes
&= ~written
;
4841 /* Track the current instruction as a previous instruction. */
4842 back
[oldest
].insn
= insn
;
4843 back
[oldest
].unit
= iunit
;
4844 back
[oldest
].delayeduse
= idelayeduse
;
4845 back
[oldest
].writes
= iwrites
;
4846 back
[oldest
].reads
= ireads
;
4847 back
[oldest
].age
= 0;
4848 oldest
= (oldest
+ 1) % max_waits
;
4855 /* {{{ OpenACC / OpenMP. */
4857 #define GCN_DEFAULT_GANGS 0 /* Choose at runtime. */
4858 #define GCN_DEFAULT_WORKERS 0 /* Choose at runtime. */
4859 #define GCN_DEFAULT_VECTORS 1 /* Use autovectorization only, for now. */
4861 /* Implement TARGET_GOACC_VALIDATE_DIMS.
4863 Check the launch dimensions provided for an OpenACC compute
4864 region, or routine. */
4867 gcn_goacc_validate_dims (tree decl
, int dims
[], int fn_level
,
4870 bool changed
= false;
4872 /* FIXME: remove -facc-experimental-workers when they're ready. */
4873 int max_workers
= flag_worker_partitioning
? 16 : 1;
4875 gcc_assert (!flag_worker_partitioning
);
4877 /* The vector size must appear to be 64, to the user, unless this is a
4878 SEQ routine. The real, internal value is always 1, which means use
4879 autovectorization, but the user should not see that. */
4880 if (fn_level
<= GOMP_DIM_VECTOR
&& fn_level
>= -1
4881 && dims
[GOMP_DIM_VECTOR
] >= 0)
4883 if (fn_level
< 0 && dims
[GOMP_DIM_VECTOR
] >= 0
4884 && dims
[GOMP_DIM_VECTOR
] != 64)
4885 warning_at (decl
? DECL_SOURCE_LOCATION (decl
) : UNKNOWN_LOCATION
,
4887 (dims
[GOMP_DIM_VECTOR
]
4888 ? G_("using %<vector_length (64)%>, ignoring %d")
4889 : G_("using %<vector_length (64)%>, "
4890 "ignoring runtime setting")),
4891 dims
[GOMP_DIM_VECTOR
]);
4892 dims
[GOMP_DIM_VECTOR
] = 1;
4896 /* Check the num workers is not too large. */
4897 if (dims
[GOMP_DIM_WORKER
] > max_workers
)
4899 warning_at (decl
? DECL_SOURCE_LOCATION (decl
) : UNKNOWN_LOCATION
,
4901 "using %<num_workers (%d)%>, ignoring %d",
4902 max_workers
, dims
[GOMP_DIM_WORKER
]);
4903 dims
[GOMP_DIM_WORKER
] = max_workers
;
4907 /* Set global defaults. */
4910 dims
[GOMP_DIM_VECTOR
] = GCN_DEFAULT_VECTORS
;
4911 if (dims
[GOMP_DIM_WORKER
] < 0)
4912 dims
[GOMP_DIM_WORKER
] = (flag_worker_partitioning
4913 ? GCN_DEFAULT_WORKERS
: 1);
4914 if (dims
[GOMP_DIM_GANG
] < 0)
4915 dims
[GOMP_DIM_GANG
] = GCN_DEFAULT_GANGS
;
4922 /* Helper function for oacc_dim_size instruction.
4923 Also used for OpenMP, via builtin_gcn_dim_size, and the omp_gcn pass. */
4926 gcn_oacc_dim_size (int dim
)
4928 if (dim
< 0 || dim
> 2)
4929 error ("offload dimension out of range (%d)", dim
);
4931 /* Vectors are a special case. */
4933 return const1_rtx
; /* Think of this as 1 times 64. */
4935 static int offset
[] = {
4936 /* Offsets into dispatch packet. */
4937 12, /* X dim = Gang / Team / Work-group. */
4938 20, /* Z dim = Worker / Thread / Wavefront. */
4939 16 /* Y dim = Vector / SIMD / Work-item. */
4941 rtx addr
= gen_rtx_PLUS (DImode
,
4942 gen_rtx_REG (DImode
,
4943 cfun
->machine
->args
.
4944 reg
[DISPATCH_PTR_ARG
]),
4945 GEN_INT (offset
[dim
]));
4946 return gen_rtx_MEM (SImode
, addr
);
4949 /* Helper function for oacc_dim_pos instruction.
4950 Also used for OpenMP, via builtin_gcn_dim_pos, and the omp_gcn pass. */
4953 gcn_oacc_dim_pos (int dim
)
4955 if (dim
< 0 || dim
> 2)
4956 error ("offload dimension out of range (%d)", dim
);
4958 static const int reg
[] = {
4959 WORKGROUP_ID_X_ARG
, /* Gang / Team / Work-group. */
4960 WORK_ITEM_ID_Z_ARG
, /* Worker / Thread / Wavefront. */
4961 WORK_ITEM_ID_Y_ARG
/* Vector / SIMD / Work-item. */
4964 int reg_num
= cfun
->machine
->args
.reg
[reg
[dim
]];
4966 /* The information must have been requested by the kernel. */
4967 gcc_assert (reg_num
>= 0);
4969 return gen_rtx_REG (SImode
, reg_num
);
4972 /* Implement TARGET_GOACC_FORK_JOIN. */
4975 gcn_fork_join (gcall
*ARG_UNUSED (call
), const int *ARG_UNUSED (dims
),
4976 bool ARG_UNUSED (is_fork
))
4978 /* GCN does not use the fork/join concept invented for NVPTX.
4979 Instead we use standard autovectorization. */
4983 /* Implement ???????
4984 FIXME make this a real hook.
4986 Adjust FNDECL such that options inherited from the host compiler
4987 are made appropriate for the accelerator compiler. */
4990 gcn_fixup_accel_lto_options (tree fndecl
)
4992 tree func_optimize
= DECL_FUNCTION_SPECIFIC_OPTIMIZATION (fndecl
);
4997 = build_optimization_node (&global_options
, &global_options_set
);
5000 /* If the function changed the optimization levels as well as
5001 setting target options, start with the optimizations
5003 if (func_optimize
!= old_optimize
)
5004 cl_optimization_restore (&global_options
, &global_options_set
,
5005 TREE_OPTIMIZATION (func_optimize
));
5007 gcn_option_override ();
5009 /* The target attributes may also change some optimization flags,
5010 so update the optimization options if necessary. */
5011 new_optimize
= build_optimization_node (&global_options
,
5012 &global_options_set
);
5014 if (old_optimize
!= new_optimize
)
5016 DECL_FUNCTION_SPECIFIC_OPTIMIZATION (fndecl
) = new_optimize
;
5017 cl_optimization_restore (&global_options
, &global_options_set
,
5018 TREE_OPTIMIZATION (old_optimize
));
5023 /* {{{ ASM Output. */
5025 /* Implement TARGET_ASM_FILE_START.
5027 Print assembler file header text. */
5030 output_file_start (void)
5035 case PROCESSOR_FIJI
: cpu
= "gfx803"; break;
5036 case PROCESSOR_VEGA10
: cpu
= "gfx900"; break;
5037 case PROCESSOR_VEGA20
: cpu
= "gfx906"; break;
5038 case PROCESSOR_GFX908
: cpu
= "gfx908+sram-ecc"; break;
5039 default: gcc_unreachable ();
5042 fprintf(asm_out_file
, "\t.amdgcn_target \"amdgcn-unknown-amdhsa--%s\"\n", cpu
);
5045 /* Implement ASM_DECLARE_FUNCTION_NAME via gcn-hsa.h.
5047 Print the initial definition of a function name.
5049 For GCN kernel entry points this includes all the HSA meta-data, special
5050 alignment constraints that don't apply to regular functions, and magic
5051 comments that pass information to mkoffload. */
5054 gcn_hsa_declare_function_name (FILE *file
, const char *name
, tree
)
5057 bool xnack_enabled
= false;
5059 fputs ("\n\n", file
);
5061 if (cfun
&& cfun
->machine
&& cfun
->machine
->normal_function
)
5063 fputs ("\t.type\t", file
);
5064 assemble_name (file
, name
);
5065 fputs (",@function\n", file
);
5066 assemble_name (file
, name
);
5067 fputs (":\n", file
);
5071 /* Determine count of sgpr/vgpr registers by looking for last
5073 for (sgpr
= 101; sgpr
>= 0; sgpr
--)
5074 if (df_regs_ever_live_p (FIRST_SGPR_REG
+ sgpr
))
5077 for (vgpr
= 255; vgpr
>= 0; vgpr
--)
5078 if (df_regs_ever_live_p (FIRST_VGPR_REG
+ vgpr
))
5082 if (!leaf_function_p ())
5084 /* We can't know how many registers function calls might use. */
5085 if (vgpr
< MAX_NORMAL_VGPR_COUNT
)
5086 vgpr
= MAX_NORMAL_VGPR_COUNT
;
5087 if (sgpr
< MAX_NORMAL_SGPR_COUNT
)
5088 sgpr
= MAX_NORMAL_SGPR_COUNT
;
5091 fputs ("\t.rodata\n"
5093 "\t.amdhsa_kernel\t", file
);
5094 assemble_name (file
, name
);
5096 int reg
= FIRST_SGPR_REG
;
5097 for (int a
= 0; a
< GCN_KERNEL_ARG_TYPES
; a
++)
5101 if ((cfun
->machine
->args
.requested
& (1 << a
))
5102 && (gcn_kernel_arg_types
[a
].fixed_regno
< 0))
5105 reg_last
= (reg_first
5106 + (GET_MODE_SIZE (gcn_kernel_arg_types
[a
].mode
)
5107 / UNITS_PER_WORD
) - 1);
5111 if (gcn_kernel_arg_types
[a
].header_pseudo
)
5113 fprintf (file
, "\t %s%s\t%i",
5114 (cfun
->machine
->args
.requested
& (1 << a
)) != 0 ? "" : ";",
5115 gcn_kernel_arg_types
[a
].header_pseudo
,
5116 (cfun
->machine
->args
.requested
& (1 << a
)) != 0);
5117 if (reg_first
!= -1)
5119 fprintf (file
, " ; (");
5120 for (int i
= reg_first
; i
<= reg_last
; ++i
)
5123 fprintf (file
, ", ");
5124 fprintf (file
, "%s", reg_names
[i
]);
5126 fprintf (file
, ")");
5128 fprintf (file
, "\n");
5130 else if (gcn_kernel_arg_types
[a
].fixed_regno
>= 0
5131 && cfun
->machine
->args
.requested
& (1 << a
))
5132 fprintf (file
, "\t ; %s\t%i (%s)\n",
5133 gcn_kernel_arg_types
[a
].name
,
5134 (cfun
->machine
->args
.requested
& (1 << a
)) != 0,
5135 reg_names
[gcn_kernel_arg_types
[a
].fixed_regno
]);
5137 fprintf (file
, "\t .amdhsa_system_vgpr_workitem_id\t%i\n",
5138 (cfun
->machine
->args
.requested
& (1 << WORK_ITEM_ID_Z_ARG
))
5140 : cfun
->machine
->args
.requested
& (1 << WORK_ITEM_ID_Y_ARG
)
5143 "\t .amdhsa_next_free_vgpr\t%i\n"
5144 "\t .amdhsa_next_free_sgpr\t%i\n"
5145 "\t .amdhsa_reserve_vcc\t1\n"
5146 "\t .amdhsa_reserve_flat_scratch\t0\n"
5147 "\t .amdhsa_reserve_xnack_mask\t%i\n"
5148 "\t .amdhsa_private_segment_fixed_size\t%i\n"
5149 "\t .amdhsa_group_segment_fixed_size\t%u\n"
5150 "\t .amdhsa_float_denorm_mode_32\t3\n"
5151 "\t .amdhsa_float_denorm_mode_16_64\t3\n",
5155 /* workitem_private_segment_bytes_size needs to be
5156 one 64th the wave-front stack size. */
5157 stack_size_opt
/ 64,
5159 fputs ("\t.end_amdhsa_kernel\n", file
);
5162 /* The following is YAML embedded in assembler; tabs are not allowed. */
5163 fputs (" .amdgpu_metadata\n"
5164 " amdhsa.version:\n"
5167 " amdhsa.kernels:\n"
5168 " - .name: ", file
);
5169 assemble_name (file
, name
);
5170 fputs ("\n .symbol: ", file
);
5171 assemble_name (file
, name
);
5174 " .kernarg_segment_size: %i\n"
5175 " .kernarg_segment_align: %i\n"
5176 " .group_segment_fixed_size: %u\n"
5177 " .private_segment_fixed_size: %i\n"
5178 " .wavefront_size: 64\n"
5179 " .sgpr_count: %i\n"
5180 " .vgpr_count: %i\n"
5181 " .max_flat_workgroup_size: 1024\n",
5182 cfun
->machine
->kernarg_segment_byte_size
,
5183 cfun
->machine
->kernarg_segment_alignment
,
5185 stack_size_opt
/ 64,
5187 fputs (" .end_amdgpu_metadata\n", file
);
5190 fputs ("\t.text\n", file
);
5191 fputs ("\t.align\t256\n", file
);
5192 fputs ("\t.type\t", file
);
5193 assemble_name (file
, name
);
5194 fputs (",@function\n", file
);
5195 assemble_name (file
, name
);
5196 fputs (":\n", file
);
5198 /* This comment is read by mkoffload. */
5200 fprintf (file
, "\t;; OPENACC-DIMS: %d, %d, %d : %s\n",
5201 oacc_get_fn_dim_size (cfun
->decl
, GOMP_DIM_GANG
),
5202 oacc_get_fn_dim_size (cfun
->decl
, GOMP_DIM_WORKER
),
5203 oacc_get_fn_dim_size (cfun
->decl
, GOMP_DIM_VECTOR
), name
);
5206 /* Implement TARGET_ASM_SELECT_SECTION.
5208 Return the section into which EXP should be placed. */
5211 gcn_asm_select_section (tree exp
, int reloc
, unsigned HOST_WIDE_INT align
)
5213 if (TREE_TYPE (exp
) != error_mark_node
5214 && TYPE_ADDR_SPACE (TREE_TYPE (exp
)) == ADDR_SPACE_LDS
)
5217 return get_section (".lds_bss",
5218 SECTION_WRITE
| SECTION_BSS
| SECTION_DEBUG
,
5221 return get_named_section (exp
, ".lds_bss", reloc
);
5224 return default_elf_select_section (exp
, reloc
, align
);
5227 /* Implement TARGET_ASM_FUNCTION_PROLOGUE.
5229 Emits custom text into the assembler file at the head of each function. */
5232 gcn_target_asm_function_prologue (FILE *file
)
5234 machine_function
*offsets
= gcn_compute_frame_offsets ();
5236 asm_fprintf (file
, "\t; using %s addressing in function\n",
5237 offsets
->use_flat_addressing
? "flat" : "global");
5239 if (offsets
->normal_function
)
5241 asm_fprintf (file
, "\t; frame pointer needed: %s\n",
5242 offsets
->need_frame_pointer
? "true" : "false");
5243 asm_fprintf (file
, "\t; lr needs saving: %s\n",
5244 offsets
->lr_needs_saving
? "true" : "false");
5245 asm_fprintf (file
, "\t; outgoing args size: %wd\n",
5246 offsets
->outgoing_args_size
);
5247 asm_fprintf (file
, "\t; pretend size: %wd\n", offsets
->pretend_size
);
5248 asm_fprintf (file
, "\t; local vars size: %wd\n", offsets
->local_vars
);
5249 asm_fprintf (file
, "\t; callee save size: %wd\n",
5250 offsets
->callee_saves
);
5254 asm_fprintf (file
, "\t; HSA kernel entry point\n");
5255 asm_fprintf (file
, "\t; local vars size: %wd\n", offsets
->local_vars
);
5256 asm_fprintf (file
, "\t; outgoing args size: %wd\n",
5257 offsets
->outgoing_args_size
);
5261 /* Helper function for print_operand and print_operand_address.
5263 Print a register as the assembler requires, according to mode and name. */
5266 print_reg (FILE *file
, rtx x
)
5268 machine_mode mode
= GET_MODE (x
);
5269 if (mode
== BImode
|| mode
== QImode
|| mode
== HImode
|| mode
== SImode
5270 || mode
== HFmode
|| mode
== SFmode
5271 || mode
== V64SFmode
|| mode
== V64SImode
5272 || mode
== V64QImode
|| mode
== V64HImode
)
5273 fprintf (file
, "%s", reg_names
[REGNO (x
)]);
5274 else if (mode
== DImode
|| mode
== V64DImode
5275 || mode
== DFmode
|| mode
== V64DFmode
)
5277 if (SGPR_REGNO_P (REGNO (x
)))
5278 fprintf (file
, "s[%i:%i]", REGNO (x
) - FIRST_SGPR_REG
,
5279 REGNO (x
) - FIRST_SGPR_REG
+ 1);
5280 else if (VGPR_REGNO_P (REGNO (x
)))
5281 fprintf (file
, "v[%i:%i]", REGNO (x
) - FIRST_VGPR_REG
,
5282 REGNO (x
) - FIRST_VGPR_REG
+ 1);
5283 else if (REGNO (x
) == FLAT_SCRATCH_REG
)
5284 fprintf (file
, "flat_scratch");
5285 else if (REGNO (x
) == EXEC_REG
)
5286 fprintf (file
, "exec");
5287 else if (REGNO (x
) == VCC_LO_REG
)
5288 fprintf (file
, "vcc");
5290 fprintf (file
, "[%s:%s]",
5291 reg_names
[REGNO (x
)], reg_names
[REGNO (x
) + 1]);
5293 else if (mode
== TImode
)
5295 if (SGPR_REGNO_P (REGNO (x
)))
5296 fprintf (file
, "s[%i:%i]", REGNO (x
) - FIRST_SGPR_REG
,
5297 REGNO (x
) - FIRST_SGPR_REG
+ 3);
5298 else if (VGPR_REGNO_P (REGNO (x
)))
5299 fprintf (file
, "v[%i:%i]", REGNO (x
) - FIRST_VGPR_REG
,
5300 REGNO (x
) - FIRST_VGPR_REG
+ 3);
5308 /* Implement TARGET_SECTION_TYPE_FLAGS.
5310 Return a set of section attributes for use by TARGET_ASM_NAMED_SECTION. */
5313 gcn_section_type_flags (tree decl
, const char *name
, int reloc
)
5315 if (strcmp (name
, ".lds_bss") == 0)
5316 return SECTION_WRITE
| SECTION_BSS
| SECTION_DEBUG
;
5318 return default_section_type_flags (decl
, name
, reloc
);
5321 /* Helper function for gcn_asm_output_symbol_ref.
5323 FIXME: If we want to have propagation blocks allocated separately and
5324 statically like this, it would be better done via symbol refs and the
5325 assembler/linker. This is a temporary hack. */
5328 gcn_print_lds_decl (FILE *f
, tree var
)
5331 machine_function
*machfun
= cfun
->machine
;
5333 if ((offset
= machfun
->lds_allocs
->get (var
)))
5334 fprintf (f
, "%u", (unsigned) *offset
);
5337 unsigned HOST_WIDE_INT align
= DECL_ALIGN_UNIT (var
);
5338 tree type
= TREE_TYPE (var
);
5339 unsigned HOST_WIDE_INT size
= tree_to_uhwi (TYPE_SIZE_UNIT (type
));
5340 if (size
> align
&& size
> 4 && align
< 8)
5343 machfun
->lds_allocated
= ((machfun
->lds_allocated
+ align
- 1)
5346 machfun
->lds_allocs
->put (var
, machfun
->lds_allocated
);
5347 fprintf (f
, "%u", machfun
->lds_allocated
);
5348 machfun
->lds_allocated
+= size
;
5349 if (machfun
->lds_allocated
> LDS_SIZE
)
5350 error ("local data-share memory exhausted");
5354 /* Implement ASM_OUTPUT_SYMBOL_REF via gcn-hsa.h. */
5357 gcn_asm_output_symbol_ref (FILE *file
, rtx x
)
5361 && (decl
= SYMBOL_REF_DECL (x
)) != 0
5362 && TREE_CODE (decl
) == VAR_DECL
5363 && AS_LDS_P (TYPE_ADDR_SPACE (TREE_TYPE (decl
))))
5365 /* LDS symbols (emitted using this hook) are only used at present
5366 to propagate worker values from an active thread to neutered
5367 threads. Use the same offset for each such block, but don't
5368 use zero because null pointers are used to identify the active
5369 thread in GOACC_single_copy_start calls. */
5370 gcn_print_lds_decl (file
, decl
);
5374 assemble_name (file
, XSTR (x
, 0));
5375 /* FIXME: See above -- this condition is unreachable. */
5377 && (decl
= SYMBOL_REF_DECL (x
)) != 0
5378 && TREE_CODE (decl
) == VAR_DECL
5379 && AS_LDS_P (TYPE_ADDR_SPACE (TREE_TYPE (decl
))))
5380 fputs ("@abs32", file
);
5384 /* Implement TARGET_CONSTANT_ALIGNMENT.
5386 Returns the alignment in bits of a constant that is being placed in memory.
5387 CONSTANT is the constant and BASIC_ALIGN is the alignment that the object
5388 would ordinarily have. */
5390 static HOST_WIDE_INT
5391 gcn_constant_alignment (const_tree
ARG_UNUSED (constant
),
5392 HOST_WIDE_INT basic_align
)
5394 return basic_align
> 128 ? basic_align
: 128;
5397 /* Implement PRINT_OPERAND_ADDRESS via gcn.h. */
5400 print_operand_address (FILE *file
, rtx mem
)
5402 gcc_assert (MEM_P (mem
));
5406 addr_space_t as
= MEM_ADDR_SPACE (mem
);
5407 rtx addr
= XEXP (mem
, 0);
5408 gcc_assert (REG_P (addr
) || GET_CODE (addr
) == PLUS
);
5410 if (AS_SCRATCH_P (as
))
5411 switch (GET_CODE (addr
))
5414 print_reg (file
, addr
);
5418 reg
= XEXP (addr
, 0);
5419 offset
= XEXP (addr
, 1);
5420 print_reg (file
, reg
);
5421 if (GET_CODE (offset
) == CONST_INT
)
5422 fprintf (file
, " offset:" HOST_WIDE_INT_PRINT_DEC
, INTVAL (offset
));
5431 else if (AS_ANY_FLAT_P (as
))
5433 if (GET_CODE (addr
) == REG
)
5434 print_reg (file
, addr
);
5437 gcc_assert (TARGET_GCN5_PLUS
);
5438 print_reg (file
, XEXP (addr
, 0));
5441 else if (AS_GLOBAL_P (as
))
5443 gcc_assert (TARGET_GCN5_PLUS
);
5446 rtx vgpr_offset
= NULL_RTX
;
5448 if (GET_CODE (addr
) == PLUS
)
5450 base
= XEXP (addr
, 0);
5452 if (GET_CODE (base
) == PLUS
)
5454 /* (SGPR + VGPR) + CONST */
5455 vgpr_offset
= XEXP (base
, 1);
5456 base
= XEXP (base
, 0);
5460 rtx offset
= XEXP (addr
, 1);
5464 vgpr_offset
= offset
;
5465 else if (CONST_INT_P (offset
))
5466 /* VGPR + CONST or SGPR + CONST */
5469 output_operand_lossage ("bad ADDR_SPACE_GLOBAL address");
5475 if (VGPR_REGNO_P (REGNO (base
)))
5476 print_reg (file
, base
);
5477 else if (SGPR_REGNO_P (REGNO (base
)))
5479 /* The assembler requires a 64-bit VGPR pair here, even though
5480 the offset should be only 32-bit. */
5481 if (vgpr_offset
== NULL_RTX
)
5482 /* In this case, the vector offset is zero, so we use the first
5483 lane of v1, which is initialized to zero. */
5484 fprintf (file
, "v[1:2]");
5485 else if (REG_P (vgpr_offset
)
5486 && VGPR_REGNO_P (REGNO (vgpr_offset
)))
5488 fprintf (file
, "v[%d:%d]",
5489 REGNO (vgpr_offset
) - FIRST_VGPR_REG
,
5490 REGNO (vgpr_offset
) - FIRST_VGPR_REG
+ 1);
5493 output_operand_lossage ("bad ADDR_SPACE_GLOBAL address");
5497 output_operand_lossage ("bad ADDR_SPACE_GLOBAL address");
5499 else if (AS_ANY_DS_P (as
))
5500 switch (GET_CODE (addr
))
5503 print_reg (file
, addr
);
5507 reg
= XEXP (addr
, 0);
5508 print_reg (file
, reg
);
5516 switch (GET_CODE (addr
))
5519 print_reg (file
, addr
);
5520 fprintf (file
, ", 0");
5524 reg
= XEXP (addr
, 0);
5525 offset
= XEXP (addr
, 1);
5526 print_reg (file
, reg
);
5527 fprintf (file
, ", ");
5528 if (GET_CODE (offset
) == REG
)
5529 print_reg (file
, reg
);
5530 else if (GET_CODE (offset
) == CONST_INT
)
5531 fprintf (file
, HOST_WIDE_INT_PRINT_DEC
, INTVAL (offset
));
5542 /* Implement PRINT_OPERAND via gcn.h.
5544 b - print operand size as untyped operand (b8/b16/b32/b64)
5545 B - print operand size as SI/DI untyped operand (b32/b32/b32/b64)
5546 i - print operand size as untyped operand (i16/b32/i64)
5547 I - print operand size as SI/DI untyped operand(i32/b32/i64)
5548 u - print operand size as untyped operand (u16/u32/u64)
5549 U - print operand size as SI/DI untyped operand(u32/u64)
5550 o - print operand size as memory access size for loads
5551 (ubyte/ushort/dword/dwordx2/wordx3/dwordx4)
5552 s - print operand size as memory access size for stores
5553 (byte/short/dword/dwordx2/wordx3/dwordx4)
5554 C - print conditional code for s_cbranch (_sccz/_sccnz/_vccz/_vccnz...)
5555 c - print inverse conditional code for s_cbranch
5556 D - print conditional code for s_cmp (eq_u64/lg_u64...)
5557 E - print conditional code for v_cmp (eq_u64/ne_u64...)
5558 A - print address in formatting suitable for given address space.
5559 O - print offset:n for data share operations.
5560 ^ - print "_co" suffix for GCN5 mnemonics
5561 g - print "glc", if appropriate for given MEM
5565 print_operand (FILE *file
, rtx x
, int code
)
5567 int xcode
= x
? GET_CODE (x
) : 0;
5568 bool invert
= false;
5571 /* Instructions have the following suffixes.
5572 If there are two suffixes, the first is the destination type,
5573 and the second is the source type.
5575 B32 Bitfield (untyped data) 32-bit
5576 B64 Bitfield (untyped data) 64-bit
5577 F16 floating-point 16-bit
5578 F32 floating-point 32-bit (IEEE 754 single-precision float)
5579 F64 floating-point 64-bit (IEEE 754 double-precision float)
5580 I16 signed 32-bit integer
5581 I32 signed 32-bit integer
5582 I64 signed 64-bit integer
5583 U16 unsigned 32-bit integer
5584 U32 unsigned 32-bit integer
5585 U64 unsigned 64-bit integer */
5587 /* Print operand size as untyped suffix. */
5591 machine_mode mode
= GET_MODE (x
);
5592 if (VECTOR_MODE_P (mode
))
5593 mode
= GET_MODE_INNER (mode
);
5594 switch (GET_MODE_SIZE (mode
))
5609 output_operand_lossage ("invalid operand %%xn code");
5618 machine_mode mode
= GET_MODE (x
);
5619 if (VECTOR_MODE_P (mode
))
5620 mode
= GET_MODE_INNER (mode
);
5621 switch (GET_MODE_SIZE (mode
))
5632 output_operand_lossage ("invalid operand %%xn code");
5639 fputs ("sext(", file
);
5640 print_operand (file
, x
, 0);
5648 bool signed_p
= code
== 'i';
5649 bool min32_p
= code
== 'I' || code
== 'U';
5651 machine_mode mode
= GET_MODE (x
);
5652 if (VECTOR_MODE_P (mode
))
5653 mode
= GET_MODE_INNER (mode
);
5654 if (mode
== VOIDmode
)
5655 switch (GET_CODE (x
))
5658 s
= signed_p
? "_i32" : "_u32";
5664 output_operand_lossage ("invalid operand %%xn code");
5667 else if (FLOAT_MODE_P (mode
))
5668 switch (GET_MODE_SIZE (mode
))
5680 output_operand_lossage ("invalid operand %%xn code");
5684 switch (GET_MODE_SIZE (mode
))
5689 s
= signed_p
? "_i32" : "_u32";
5692 s
= signed_p
? "_i64" : "_u64";
5695 output_operand_lossage ("invalid operand %%xn code");
5699 switch (GET_MODE_SIZE (mode
))
5702 s
= signed_p
? "_i8" : "_u8";
5705 s
= signed_p
? "_i16" : "_u16";
5708 s
= signed_p
? "_i32" : "_u32";
5711 s
= signed_p
? "_i64" : "_u64";
5714 output_operand_lossage ("invalid operand %%xn code");
5720 /* Print operand size as untyped suffix. */
5724 switch (GET_MODE_SIZE (GET_MODE (x
)))
5732 /* The following are full-vector variants. */
5747 /* Fall-through - the other cases for 'o' are the same as for 's'. */
5753 switch (GET_MODE_SIZE (GET_MODE (x
)))
5777 s
= VECTOR_MODE_P (GET_MODE (x
)) ? "_byte" : "_dwordx16";
5779 /* The following are full-vector variants. */
5790 output_operand_lossage ("invalid operand %%xn code");
5799 output_operand_lossage ("invalid %%xn code");
5802 print_operand_address (file
, x
);
5808 output_operand_lossage ("invalid %%xn code");
5811 if (AS_GDS_P (MEM_ADDR_SPACE (x
)))
5812 fprintf (file
, " gds");
5814 rtx x0
= XEXP (x
, 0);
5815 if (AS_GLOBAL_P (MEM_ADDR_SPACE (x
)))
5817 gcc_assert (TARGET_GCN5_PLUS
);
5819 fprintf (file
, ", ");
5822 rtx const_offset
= NULL_RTX
;
5824 if (GET_CODE (base
) == PLUS
)
5826 rtx offset
= XEXP (x0
, 1);
5827 base
= XEXP (x0
, 0);
5829 if (GET_CODE (base
) == PLUS
)
5830 /* (SGPR + VGPR) + CONST */
5831 /* Ignore the VGPR offset for this operand. */
5832 base
= XEXP (base
, 0);
5834 if (CONST_INT_P (offset
))
5835 const_offset
= XEXP (x0
, 1);
5836 else if (REG_P (offset
))
5838 /* Ignore the VGPR offset for this operand. */
5841 output_operand_lossage ("bad ADDR_SPACE_GLOBAL address");
5846 if (VGPR_REGNO_P (REGNO (base
)))
5847 /* The VGPR address is specified in the %A operand. */
5848 fprintf (file
, "off");
5849 else if (SGPR_REGNO_P (REGNO (base
)))
5850 print_reg (file
, base
);
5852 output_operand_lossage ("bad ADDR_SPACE_GLOBAL address");
5855 output_operand_lossage ("bad ADDR_SPACE_GLOBAL address");
5857 if (const_offset
!= NULL_RTX
)
5858 fprintf (file
, " offset:" HOST_WIDE_INT_PRINT_DEC
,
5859 INTVAL (const_offset
));
5864 if (GET_CODE (x0
) == REG
)
5866 if (GET_CODE (x0
) != PLUS
)
5868 output_operand_lossage ("invalid %%xn code");
5871 rtx val
= XEXP (x0
, 1);
5872 if (GET_CODE (val
) == CONST_VECTOR
)
5873 val
= CONST_VECTOR_ELT (val
, 0);
5874 if (GET_CODE (val
) != CONST_INT
)
5876 output_operand_lossage ("invalid %%xn code");
5879 fprintf (file
, " offset:" HOST_WIDE_INT_PRINT_DEC
, INTVAL (val
));
5890 if ((xcode
!= EQ
&& xcode
!= NE
) || !REG_P (XEXP (x
, 0)))
5892 output_operand_lossage ("invalid %%xn code");
5895 switch (REGNO (XEXP (x
, 0)))
5902 /* For some reason llvm-mc insists on scc0 instead of sccz. */
5910 output_operand_lossage ("invalid %%xn code");
5914 if (xcode
== (invert
? NE
: EQ
))
5915 fputc (num
? '0' : 'z', file
);
5917 fputs (num
? "1" : "nz", file
);
5923 bool cmp_signed
= false;
5961 output_operand_lossage ("invalid %%xn code");
5965 fputc (cmp_signed
? 'i' : 'u', file
);
5967 machine_mode mode
= GET_MODE (XEXP (x
, 0));
5969 if (mode
== VOIDmode
)
5970 mode
= GET_MODE (XEXP (x
, 1));
5972 /* If both sides are constants, then assume the instruction is in
5973 SImode since s_cmp can only do integer compares. */
5974 if (mode
== VOIDmode
)
5977 switch (GET_MODE_SIZE (mode
))
5986 output_operand_lossage ("invalid operand %%xn code");
5995 bool cmp_signed
= false;
5996 machine_mode mode
= GET_MODE (XEXP (x
, 0));
5998 if (mode
== VOIDmode
)
5999 mode
= GET_MODE (XEXP (x
, 1));
6001 /* If both sides are constants, assume the instruction is in SFmode
6002 if either operand is floating point, otherwise assume SImode. */
6003 if (mode
== VOIDmode
)
6005 if (GET_CODE (XEXP (x
, 0)) == CONST_DOUBLE
6006 || GET_CODE (XEXP (x
, 1)) == CONST_DOUBLE
)
6012 /* Use the same format code for vector comparisons. */
6013 if (GET_MODE_CLASS (mode
) == MODE_VECTOR_FLOAT
6014 || GET_MODE_CLASS (mode
) == MODE_VECTOR_INT
)
6015 mode
= GET_MODE_INNER (mode
);
6017 bool float_p
= GET_MODE_CLASS (mode
) == MODE_FLOAT
;
6025 s
= float_p
? "_neq_" : "_ne_";
6080 output_operand_lossage ("invalid %%xn code");
6084 fputc (float_p
? 'f' : cmp_signed
? 'i' : 'u', file
);
6086 switch (GET_MODE_SIZE (mode
))
6089 output_operand_lossage ("operand %%xn code invalid for QImode");
6101 output_operand_lossage ("invalid operand %%xn code");
6108 print_operand (file
, gcn_operand_part (GET_MODE (x
), x
, 0), 0);
6111 print_operand (file
, gcn_operand_part (GET_MODE (x
), x
, 1), 0);
6114 /* Print a scalar register number as an integer. Temporary hack. */
6115 gcc_assert (REG_P (x
));
6116 fprintf (file
, "%u", (int) REGNO (x
));
6119 /* Print a vector register number as an integer. Temporary hack. */
6120 gcc_assert (REG_P (x
));
6121 fprintf (file
, "%u", (int) REGNO (x
) - FIRST_VGPR_REG
);
6125 print_reg (file
, x
);
6126 else if (xcode
== MEM
)
6127 output_address (GET_MODE (x
), x
);
6128 else if (xcode
== CONST_INT
)
6129 fprintf (file
, "%i", (int) INTVAL (x
));
6130 else if (xcode
== CONST_VECTOR
)
6131 print_operand (file
, CONST_VECTOR_ELT (x
, 0), code
);
6132 else if (xcode
== CONST_DOUBLE
)
6135 switch (gcn_inline_fp_constant_p (x
, false))
6165 rtx ix
= simplify_gen_subreg (GET_MODE (x
) == DFmode
6167 x
, GET_MODE (x
), 0);
6169 print_operand (file
, ix
, code
);
6171 output_operand_lossage ("invalid fp constant");
6175 fprintf (file
, str
);
6179 output_addr_const (file
, x
);
6182 if (TARGET_GCN5_PLUS
)
6183 fputs ("_co", file
);
6186 gcc_assert (xcode
== MEM
);
6187 if (MEM_VOLATILE_P (x
))
6188 fputs (" glc", file
);
6191 output_operand_lossage ("invalid %%xn code");
6196 /* Implement DBX_REGISTER_NUMBER macro.
6198 Return the DWARF register number that corresponds to the GCC internal
6202 gcn_dwarf_register_number (unsigned int regno
)
6204 /* Registers defined in DWARF. */
6205 if (regno
== EXEC_LO_REG
)
6207 /* We need to use a more complex DWARF expression for this
6208 else if (regno == EXEC_HI_REG)
6210 else if (regno
== VCC_LO_REG
)
6212 /* We need to use a more complex DWARF expression for this
6213 else if (regno == VCC_HI_REG)
6215 else if (regno
== SCC_REG
)
6217 else if (SGPR_REGNO_P (regno
))
6219 if (regno
- FIRST_SGPR_REG
< 64)
6220 return (regno
- FIRST_SGPR_REG
+ 32);
6222 return (regno
- FIRST_SGPR_REG
+ 1024);
6224 else if (VGPR_REGNO_P (regno
))
6225 return (regno
- FIRST_VGPR_REG
+ 2560);
6227 /* Otherwise, there's nothing sensible to do. */
6228 return regno
+ 100000;
6231 /* Implement TARGET_DWARF_REGISTER_SPAN.
6233 DImode and Vector DImode require additional registers. */
6236 gcn_dwarf_register_span (rtx rtl
)
6238 machine_mode mode
= GET_MODE (rtl
);
6240 if (VECTOR_MODE_P (mode
))
6241 mode
= GET_MODE_INNER (mode
);
6243 if (GET_MODE_SIZE (mode
) != 8)
6246 rtx p
= gen_rtx_PARALLEL (VOIDmode
, rtvec_alloc (2));
6247 unsigned regno
= REGNO (rtl
);
6248 XVECEXP (p
, 0, 0) = gen_rtx_REG (SImode
, regno
);
6249 XVECEXP (p
, 0, 1) = gen_rtx_REG (SImode
, regno
+ 1);
6255 /* {{{ TARGET hook overrides. */
6257 #undef TARGET_ADDR_SPACE_ADDRESS_MODE
6258 #define TARGET_ADDR_SPACE_ADDRESS_MODE gcn_addr_space_address_mode
6259 #undef TARGET_ADDR_SPACE_LEGITIMATE_ADDRESS_P
6260 #define TARGET_ADDR_SPACE_LEGITIMATE_ADDRESS_P \
6261 gcn_addr_space_legitimate_address_p
6262 #undef TARGET_ADDR_SPACE_LEGITIMIZE_ADDRESS
6263 #define TARGET_ADDR_SPACE_LEGITIMIZE_ADDRESS gcn_addr_space_legitimize_address
6264 #undef TARGET_ADDR_SPACE_POINTER_MODE
6265 #define TARGET_ADDR_SPACE_POINTER_MODE gcn_addr_space_pointer_mode
6266 #undef TARGET_ADDR_SPACE_SUBSET_P
6267 #define TARGET_ADDR_SPACE_SUBSET_P gcn_addr_space_subset_p
6268 #undef TARGET_ADDR_SPACE_CONVERT
6269 #define TARGET_ADDR_SPACE_CONVERT gcn_addr_space_convert
6270 #undef TARGET_ARG_PARTIAL_BYTES
6271 #define TARGET_ARG_PARTIAL_BYTES gcn_arg_partial_bytes
6272 #undef TARGET_ASM_ALIGNED_DI_OP
6273 #define TARGET_ASM_ALIGNED_DI_OP "\t.8byte\t"
6274 #undef TARGET_ASM_FILE_START
6275 #define TARGET_ASM_FILE_START output_file_start
6276 #undef TARGET_ASM_FUNCTION_PROLOGUE
6277 #define TARGET_ASM_FUNCTION_PROLOGUE gcn_target_asm_function_prologue
6278 #undef TARGET_ASM_SELECT_SECTION
6279 #define TARGET_ASM_SELECT_SECTION gcn_asm_select_section
6280 #undef TARGET_ASM_TRAMPOLINE_TEMPLATE
6281 #define TARGET_ASM_TRAMPOLINE_TEMPLATE gcn_asm_trampoline_template
6282 #undef TARGET_ATTRIBUTE_TABLE
6283 #define TARGET_ATTRIBUTE_TABLE gcn_attribute_table
6284 #undef TARGET_BUILTIN_DECL
6285 #define TARGET_BUILTIN_DECL gcn_builtin_decl
6286 #undef TARGET_CAN_CHANGE_MODE_CLASS
6287 #define TARGET_CAN_CHANGE_MODE_CLASS gcn_can_change_mode_class
6288 #undef TARGET_CAN_ELIMINATE
6289 #define TARGET_CAN_ELIMINATE gcn_can_eliminate_p
6290 #undef TARGET_CANNOT_COPY_INSN_P
6291 #define TARGET_CANNOT_COPY_INSN_P gcn_cannot_copy_insn_p
6292 #undef TARGET_CLASS_LIKELY_SPILLED_P
6293 #define TARGET_CLASS_LIKELY_SPILLED_P gcn_class_likely_spilled_p
6294 #undef TARGET_CLASS_MAX_NREGS
6295 #define TARGET_CLASS_MAX_NREGS gcn_class_max_nregs
6296 #undef TARGET_CONDITIONAL_REGISTER_USAGE
6297 #define TARGET_CONDITIONAL_REGISTER_USAGE gcn_conditional_register_usage
6298 #undef TARGET_CONSTANT_ALIGNMENT
6299 #define TARGET_CONSTANT_ALIGNMENT gcn_constant_alignment
6300 #undef TARGET_DEBUG_UNWIND_INFO
6301 #define TARGET_DEBUG_UNWIND_INFO gcn_debug_unwind_info
6302 #undef TARGET_DWARF_REGISTER_SPAN
6303 #define TARGET_DWARF_REGISTER_SPAN gcn_dwarf_register_span
6304 #undef TARGET_EMUTLS_VAR_INIT
6305 #define TARGET_EMUTLS_VAR_INIT gcn_emutls_var_init
6306 #undef TARGET_EXPAND_BUILTIN
6307 #define TARGET_EXPAND_BUILTIN gcn_expand_builtin
6308 #undef TARGET_FUNCTION_ARG
6309 #undef TARGET_FUNCTION_ARG_ADVANCE
6310 #define TARGET_FUNCTION_ARG_ADVANCE gcn_function_arg_advance
6311 #define TARGET_FUNCTION_ARG gcn_function_arg
6312 #undef TARGET_FUNCTION_VALUE
6313 #define TARGET_FUNCTION_VALUE gcn_function_value
6314 #undef TARGET_FUNCTION_VALUE_REGNO_P
6315 #define TARGET_FUNCTION_VALUE_REGNO_P gcn_function_value_regno_p
6316 #undef TARGET_GIMPLIFY_VA_ARG_EXPR
6317 #define TARGET_GIMPLIFY_VA_ARG_EXPR gcn_gimplify_va_arg_expr
6318 #undef TARGET_OMP_DEVICE_KIND_ARCH_ISA
6319 #define TARGET_OMP_DEVICE_KIND_ARCH_ISA gcn_omp_device_kind_arch_isa
6320 #undef TARGET_GOACC_ADJUST_PROPAGATION_RECORD
6321 #define TARGET_GOACC_ADJUST_PROPAGATION_RECORD \
6322 gcn_goacc_adjust_propagation_record
6323 #undef TARGET_GOACC_ADJUST_GANGPRIVATE_DECL
6324 #define TARGET_GOACC_ADJUST_GANGPRIVATE_DECL gcn_goacc_adjust_gangprivate_decl
6325 #undef TARGET_GOACC_FORK_JOIN
6326 #define TARGET_GOACC_FORK_JOIN gcn_fork_join
6327 #undef TARGET_GOACC_REDUCTION
6328 #define TARGET_GOACC_REDUCTION gcn_goacc_reduction
6329 #undef TARGET_GOACC_VALIDATE_DIMS
6330 #define TARGET_GOACC_VALIDATE_DIMS gcn_goacc_validate_dims
6331 #undef TARGET_HARD_REGNO_MODE_OK
6332 #define TARGET_HARD_REGNO_MODE_OK gcn_hard_regno_mode_ok
6333 #undef TARGET_HARD_REGNO_NREGS
6334 #define TARGET_HARD_REGNO_NREGS gcn_hard_regno_nregs
6335 #undef TARGET_HAVE_SPECULATION_SAFE_VALUE
6336 #define TARGET_HAVE_SPECULATION_SAFE_VALUE speculation_safe_value_not_needed
6337 #undef TARGET_INIT_BUILTINS
6338 #define TARGET_INIT_BUILTINS gcn_init_builtins
6339 #undef TARGET_IRA_CHANGE_PSEUDO_ALLOCNO_CLASS
6340 #define TARGET_IRA_CHANGE_PSEUDO_ALLOCNO_CLASS \
6341 gcn_ira_change_pseudo_allocno_class
6342 #undef TARGET_LEGITIMATE_CONSTANT_P
6343 #define TARGET_LEGITIMATE_CONSTANT_P gcn_legitimate_constant_p
6345 #define TARGET_LRA_P hook_bool_void_true
6346 #undef TARGET_MACHINE_DEPENDENT_REORG
6347 #define TARGET_MACHINE_DEPENDENT_REORG gcn_md_reorg
6348 #undef TARGET_MEMORY_MOVE_COST
6349 #define TARGET_MEMORY_MOVE_COST gcn_memory_move_cost
6350 #undef TARGET_MODES_TIEABLE_P
6351 #define TARGET_MODES_TIEABLE_P gcn_modes_tieable_p
6352 #undef TARGET_OPTION_OVERRIDE
6353 #define TARGET_OPTION_OVERRIDE gcn_option_override
6354 #undef TARGET_PRETEND_OUTGOING_VARARGS_NAMED
6355 #define TARGET_PRETEND_OUTGOING_VARARGS_NAMED \
6356 gcn_pretend_outgoing_varargs_named
6357 #undef TARGET_PROMOTE_FUNCTION_MODE
6358 #define TARGET_PROMOTE_FUNCTION_MODE gcn_promote_function_mode
6359 #undef TARGET_REGISTER_MOVE_COST
6360 #define TARGET_REGISTER_MOVE_COST gcn_register_move_cost
6361 #undef TARGET_RETURN_IN_MEMORY
6362 #define TARGET_RETURN_IN_MEMORY gcn_return_in_memory
6363 #undef TARGET_RTX_COSTS
6364 #define TARGET_RTX_COSTS gcn_rtx_costs
6365 #undef TARGET_SECONDARY_RELOAD
6366 #define TARGET_SECONDARY_RELOAD gcn_secondary_reload
6367 #undef TARGET_SECTION_TYPE_FLAGS
6368 #define TARGET_SECTION_TYPE_FLAGS gcn_section_type_flags
6369 #undef TARGET_SCALAR_MODE_SUPPORTED_P
6370 #define TARGET_SCALAR_MODE_SUPPORTED_P gcn_scalar_mode_supported_p
6371 #undef TARGET_SMALL_REGISTER_CLASSES_FOR_MODE_P
6372 #define TARGET_SMALL_REGISTER_CLASSES_FOR_MODE_P \
6373 gcn_small_register_classes_for_mode_p
6374 #undef TARGET_SPILL_CLASS
6375 #define TARGET_SPILL_CLASS gcn_spill_class
6376 #undef TARGET_STRICT_ARGUMENT_NAMING
6377 #define TARGET_STRICT_ARGUMENT_NAMING gcn_strict_argument_naming
6378 #undef TARGET_TRAMPOLINE_INIT
6379 #define TARGET_TRAMPOLINE_INIT gcn_trampoline_init
6380 #undef TARGET_TRULY_NOOP_TRUNCATION
6381 #define TARGET_TRULY_NOOP_TRUNCATION gcn_truly_noop_truncation
6382 #undef TARGET_VECTORIZE_BUILTIN_VECTORIZATION_COST
6383 #define TARGET_VECTORIZE_BUILTIN_VECTORIZATION_COST gcn_vectorization_cost
6384 #undef TARGET_VECTORIZE_GET_MASK_MODE
6385 #define TARGET_VECTORIZE_GET_MASK_MODE gcn_vectorize_get_mask_mode
6386 #undef TARGET_VECTORIZE_PREFERRED_SIMD_MODE
6387 #define TARGET_VECTORIZE_PREFERRED_SIMD_MODE gcn_vectorize_preferred_simd_mode
6388 #undef TARGET_VECTORIZE_PREFERRED_VECTOR_ALIGNMENT
6389 #define TARGET_VECTORIZE_PREFERRED_VECTOR_ALIGNMENT \
6390 gcn_preferred_vector_alignment
6391 #undef TARGET_VECTORIZE_RELATED_MODE
6392 #define TARGET_VECTORIZE_RELATED_MODE gcn_related_vector_mode
6393 #undef TARGET_VECTORIZE_SUPPORT_VECTOR_MISALIGNMENT
6394 #define TARGET_VECTORIZE_SUPPORT_VECTOR_MISALIGNMENT \
6395 gcn_vectorize_support_vector_misalignment
6396 #undef TARGET_VECTORIZE_VEC_PERM_CONST
6397 #define TARGET_VECTORIZE_VEC_PERM_CONST gcn_vectorize_vec_perm_const
6398 #undef TARGET_VECTORIZE_VECTOR_ALIGNMENT_REACHABLE
6399 #define TARGET_VECTORIZE_VECTOR_ALIGNMENT_REACHABLE \
6400 gcn_vector_alignment_reachable
6401 #undef TARGET_VECTOR_MODE_SUPPORTED_P
6402 #define TARGET_VECTOR_MODE_SUPPORTED_P gcn_vector_mode_supported_p
6404 struct gcc_target targetm
= TARGET_INITIALIZER
;