summaryrefslogtreecommitdiff
path: root/gcc/brig/brigfrontend/phsa.h
blob: 02b2bd73c80d0ec5c6132aca0f2faf01c9660c95 (plain)
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
/* phsa.h -- interfacing between the gcc BRIG FE and the phsa runtime
   Copyright (C) 2016-2020 Free Software Foundation, Inc.
   Contributed by Pekka Jaaskelainen <pekka.jaaskelainen@parmance.com>
   for General Processor Tech.

This file is part of GCC.

GCC is free software; you can redistribute it and/or modify it under
the terms of the GNU General Public License as published by the Free
Software Foundation; either version 3, or (at your option) any later
version.

GCC is distributed in the hope that it will be useful, but WITHOUT ANY
WARRANTY; without even the implied warranty of MERCHANTABILITY or
FITNESS FOR A PARTICULAR PURPOSE.  See the GNU General Public License
for more details.

You should have received a copy of the GNU General Public License
along with GCC; see the file COPYING3.  If not see
<http://www.gnu.org/licenses/>.  */

#ifndef PHSA_H
#define PHSA_H

#include <stdint.h>

/* This struct is used to pass information from the BRIG FE to the
   runtime of the finalizer kernel, its control directives etc.
   The data is passed raw in a special ELF section named
   phsa.kerneldesc.kernel_function_name.  */

typedef struct __attribute__((__packed__))
{
  /* Set to 1 in case the function is a kernel.  */
  uint8_t is_kernel;
  /* The size of the group segment used by the kernel.  */
  uint32_t group_segment_size;
  /* Size of the private segment used by a single work-item.  */
  uint32_t private_segment_size;
  /* Total size of the kernel arguments.  */
  uint32_t kernarg_segment_size;
  /* Maximum alignment of a kernel argument variable.  */
  uint16_t kernarg_max_align;
  /* Maximum size (in bytes) of dynamic group memory.  */
  uint32_t max_dynamic_group_size;
  /* Max number of work-items used to launch the kernel.  */
  uint64_t max_flat_grid_size;
  /* Max number of work-items in a work-group used to launch the kernel.  */
  uint32_t max_flat_workgroup_size;
  /* The grid size required by the kernel.  */
  uint64_t required_grid_size[3];
  /* The work group size required by the kernel.  */
  uint32_t required_workgroup_size[3];
  /* The number of dimensions required by the kernel.  */
  uint8_t required_dim;

} phsa_descriptor;

/* The prefix to use in the ELF section containing descriptor for
   a function.  */

#define PHSA_DESC_SECTION_PREFIX "phsa.desc."
#define PHSA_HOST_DEF_PTR_PREFIX "__phsa.host_def."

/* The frontend error messages are parsed by the host runtime.  Known
   prefix strings are used to separate the different runtime error
   codes.  */

#define PHSA_ERROR_PREFIX_INCOMPATIBLE_MODULE "Incompatible module: "
#define PHSA_ERROR_PREFIX_CORRUPTED_MODULE "Corrupted module: "

/* Offsets of attributes in the PHSA context structs.
   Used by -fphsa-wi-context-opt.  */
#define PHSA_CONTEXT_OFFS_WI_IDS 0
#define PHSA_CONTEXT_OFFS_WG_IDS (PHSA_CONTEXT_OFFS_WI_IDS + 3 * 4)
#define PHSA_CONTEXT_WG_SIZES (PHSA_CONTEXT_OFFS_WG_IDS + 3 * 4)
#define PHSA_CONTEXT_CURRENT_WG_SIZES (PHSA_CONTEXT_WG_SIZES + 3 * 4)

#endif