libgomp/config/nvptx/oacc-parallel.c

   1 /* OpenACC constructs
   2
   3    Copyright (C) 2014-2018 Free Software Foundation, Inc.
   4
   5    Contributed by Mentor Embedded.
   6
   7    This file is part of the GNU Offloading and Multi Processing Library
   8    (libgomp).
   9
  10    Libgomp is free software; you can redistribute it and/or modify it
  11    under the terms of the GNU General Public License as published by
  12    the Free Software Foundation; either version 3, or (at your option)
  13    any later version.
  14
  15    Libgomp is distributed in the hope that it will be useful, but WITHOUT ANY
  16    WARRANTY; without even the implied warranty of MERCHANTABILITY or FITNESS
  17    FOR A PARTICULAR PURPOSE.  See the GNU General Public License for
  18    more details.
  19
  20    Under Section 7 of GPL version 3, you are granted additional
  21    permissions described in the GCC Runtime Library Exception, version
  22    3.1, as published by the Free Software Foundation.
  23
  24    You should have received a copy of the GNU General Public License and
  25    a copy of the GCC Runtime Library Exception along with this program;
  26    see the files COPYING3 and COPYING.RUNTIME respectively.  If not, see
  27    <http://www.gnu.org/licenses/>.  */
  28
  29 #include "libgomp_g.h"
  30
  31 __asm__ (".visible .func (.param .u32 %out_retval) GOACC_tid (.param .u32 %in_ar1);\n"
  32          ".visible .func (.param .u32 %out_retval) GOACC_ntid (.param .u32 %in_ar1);\n"
  33          ".visible .func (.param .u32 %out_retval) GOACC_ctaid (.param .u32 %in_ar1);\n"
  34          ".visible .func (.param .u32 %out_retval) GOACC_nctaid (.param .u32 %in_ar1);\n"
  35          "// BEGIN GLOBAL FUNCTION DECL: GOACC_get_num_threads\n"
  36          ".visible .func (.param .u32 %out_retval) GOACC_get_num_threads;\n"
  37          "// BEGIN GLOBAL FUNCTION DECL: GOACC_get_thread_num\n"
  38          ".visible .func (.param .u32 %out_retval) GOACC_get_thread_num;\n"
  39          "// BEGIN GLOBAL FUNCTION DECL: abort\n"
  40          ".extern .func abort;\n"
  41          ".visible .func (.param .u32 %out_retval) GOACC_tid (.param .u32 %in_ar1)\n"
  42          "{\n"
  43          ".reg .u32 %ar1;\n"
  44          ".reg .u32 %retval;\n"
  45          ".reg .u64 %hr10;\n"
  46          ".reg .u32 %r22;\n"
  47          ".reg .u32 %r23;\n"
  48          ".reg .u32 %r24;\n"
  49          ".reg .u32 %r25;\n"
  50          ".reg .u32 %r26;\n"
  51          ".reg .u32 %r27;\n"
  52          ".reg .u32 %r28;\n"
  53          ".reg .u32 %r29;\n"
  54          ".reg .pred %r30;\n"
  55          ".reg .u32 %r31;\n"
  56          ".reg .pred %r32;\n"
  57          ".reg .u32 %r33;\n"
  58          ".reg .pred %r34;\n"
  59          ".local .align 8 .b8 %frame[4];\n"
  60          "ld.param.u32 %ar1,[%in_ar1];\n"
  61          "mov.u32 %r27,%ar1;\n"
  62          "st.local.u32 [%frame],%r27;\n"
  63          "ld.local.u32 %r28,[%frame];\n"
  64          "mov.u32 %r29,1;\n"
  65          "setp.eq.u32 %r30,%r28,%r29;\n"
  66          "@%r30 bra $L4;\n"
  67          "mov.u32 %r31,2;\n"
  68          "setp.eq.u32 %r32,%r28,%r31;\n"
  69          "@%r32 bra $L5;\n"
  70          "mov.u32 %r33,0;\n"
  71          "setp.eq.u32 %r34,%r28,%r33;\n"
  72          "@!%r34 bra $L8;\n"
  73          "mov.u32 %r23,%tid.x;\n"
  74          "mov.u32 %r22,%r23;\n"
  75          "bra $L7;\n"
  76          "$L4:\n"
  77          "mov.u32 %r24,%tid.y;\n"
  78          "mov.u32 %r22,%r24;\n"
  79          "bra $L7;\n"
  80          "$L5:\n"
  81          "mov.u32 %r25,%tid.z;\n"
  82          "mov.u32 %r22,%r25;\n"
  83          "bra $L7;\n"
  84          "$L8:\n"
  85          "{\n"
  86          "{\n"
  87          "call abort;\n"
  88          "}\n"
  89          "}\n"
  90          "$L7:\n"
  91          "mov.u32 %r26,%r22;\n"
  92          "mov.u32 %retval,%r26;\n"
  93          "st.param.u32 [%out_retval],%retval;\n"
  94          "ret;\n"
  95          "}\n"
  96          ".visible .func (.param .u32 %out_retval) GOACC_ntid (.param .u32 %in_ar1)\n"
  97          "{\n"
  98          ".reg .u32 %ar1;\n"
  99          ".reg .u32 %retval;\n"
 100          ".reg .u64 %hr10;\n"
 101          ".reg .u32 %r22;\n"
 102          ".reg .u32 %r23;\n"
 103          ".reg .u32 %r24;\n"
 104          ".reg .u32 %r25;\n"
 105          ".reg .u32 %r26;\n"
 106          ".reg .u32 %r27;\n"
 107          ".reg .u32 %r28;\n"
 108          ".reg .u32 %r29;\n"
 109          ".reg .pred %r30;\n"
 110          ".reg .u32 %r31;\n"
 111          ".reg .pred %r32;\n"
 112          ".reg .u32 %r33;\n"
 113          ".reg .pred %r34;\n"
 114          ".local .align 8 .b8 %frame[4];\n"
 115          "ld.param.u32 %ar1,[%in_ar1];\n"
 116          "mov.u32 %r27,%ar1;\n"
 117          "st.local.u32 [%frame],%r27;\n"
 118          "ld.local.u32 %r28,[%frame];\n"
 119          "mov.u32 %r29,1;\n"
 120          "setp.eq.u32 %r30,%r28,%r29;\n"
 121          "@%r30 bra $L11;\n"
 122          "mov.u32 %r31,2;\n"
 123          "setp.eq.u32 %r32,%r28,%r31;\n"
 124          "@%r32 bra $L12;\n"
 125          "mov.u32 %r33,0;\n"
 126          "setp.eq.u32 %r34,%r28,%r33;\n"
 127          "@!%r34 bra $L15;\n"
 128          "mov.u32 %r23,%ntid.x;\n"
 129          "mov.u32 %r22,%r23;\n"
 130          "bra $L14;\n"
 131          "$L11:\n"
 132          "mov.u32 %r24,%ntid.y;\n"
 133          "mov.u32 %r22,%r24;\n"
 134          "bra $L14;\n"
 135          "$L12:\n"
 136          "mov.u32 %r25,%ntid.z;\n"
 137          "mov.u32 %r22,%r25;\n"
 138          "bra $L14;\n"
 139          "$L15:\n"
 140          "{\n"
 141          "{\n"
 142          "call abort;\n"
 143          "}\n"
 144          "}\n"
 145          "$L14:\n"
 146          "mov.u32 %r26,%r22;\n"
 147          "mov.u32 %retval,%r26;\n"
 148          "st.param.u32 [%out_retval],%retval;\n"
 149          "ret;\n"
 150          "}\n"
 151          ".visible .func (.param .u32 %out_retval) GOACC_ctaid (.param .u32 %in_ar1)\n"
 152          "{\n"
 153          ".reg .u32 %ar1;\n"
 154          ".reg .u32 %retval;\n"
 155          ".reg .u64 %hr10;\n"
 156          ".reg .u32 %r22;\n"
 157          ".reg .u32 %r23;\n"
 158          ".reg .u32 %r24;\n"
 159          ".reg .u32 %r25;\n"
 160          ".reg .u32 %r26;\n"
 161          ".reg .u32 %r27;\n"
 162          ".reg .u32 %r28;\n"
 163          ".reg .u32 %r29;\n"
 164          ".reg .pred %r30;\n"
 165          ".reg .u32 %r31;\n"
 166          ".reg .pred %r32;\n"
 167          ".reg .u32 %r33;\n"
 168          ".reg .pred %r34;\n"
 169          ".local .align 8 .b8 %frame[4];\n"
 170          "ld.param.u32 %ar1,[%in_ar1];\n"
 171          "mov.u32 %r27,%ar1;\n"
 172          "st.local.u32 [%frame],%r27;\n"
 173          "ld.local.u32 %r28,[%frame];\n"
 174          "mov.u32 %r29,1;\n"
 175          "setp.eq.u32 %r30,%r28,%r29;\n"
 176          "@%r30 bra $L18;\n"
 177          "mov.u32 %r31,2;\n"
 178          "setp.eq.u32 %r32,%r28,%r31;\n"
 179          "@%r32 bra $L19;\n"
 180          "mov.u32 %r33,0;\n"
 181          "setp.eq.u32 %r34,%r28,%r33;\n"
 182          "@!%r34 bra $L22;\n"
 183          "mov.u32 %r23,%ctaid.x;\n"
 184          "mov.u32 %r22,%r23;\n"
 185          "bra $L21;\n"
 186          "$L18:\n"
 187          "mov.u32 %r24,%ctaid.y;\n"
 188          "mov.u32 %r22,%r24;\n"
 189          "bra $L21;\n"
 190          "$L19:\n"
 191          "mov.u32 %r25,%ctaid.z;\n"
 192          "mov.u32 %r22,%r25;\n"
 193          "bra $L21;\n"
 194          "$L22:\n"
 195          "{\n"
 196          "{\n"
 197          "call abort;\n"
 198          "}\n"
 199          "}\n"
 200          "$L21:\n"
 201          "mov.u32 %r26,%r22;\n"
 202          "mov.u32 %retval,%r26;\n"
 203          "st.param.u32 [%out_retval],%retval;\n"
 204          "ret;\n"
 205          "}\n"
 206          ".visible .func (.param .u32 %out_retval) GOACC_nctaid (.param .u32 %in_ar1)\n"
 207          "{\n"
 208          ".reg .u32 %ar1;\n"
 209          ".reg .u32 %retval;\n"
 210          ".reg .u64 %hr10;\n"
 211          ".reg .u32 %r22;\n"
 212          ".reg .u32 %r23;\n"
 213          ".reg .u32 %r24;\n"
 214          ".reg .u32 %r25;\n"
 215          ".reg .u32 %r26;\n"
 216          ".reg .u32 %r27;\n"
 217          ".reg .u32 %r28;\n"
 218          ".reg .u32 %r29;\n"
 219          ".reg .pred %r30;\n"
 220          ".reg .u32 %r31;\n"
 221          ".reg .pred %r32;\n"
 222          ".reg .u32 %r33;\n"
 223          ".reg .pred %r34;\n"
 224          ".local .align 8 .b8 %frame[4];\n"
 225          "ld.param.u32 %ar1,[%in_ar1];\n"
 226          "mov.u32 %r27,%ar1;\n"
 227          "st.local.u32 [%frame],%r27;\n"
 228          "ld.local.u32 %r28,[%frame];\n"
 229          "mov.u32 %r29,1;\n"
 230          "setp.eq.u32 %r30,%r28,%r29;\n"
 231          "@%r30 bra $L25;\n"
 232          "mov.u32 %r31,2;\n"
 233          "setp.eq.u32 %r32,%r28,%r31;\n"
 234          "@%r32 bra $L26;\n"
 235          "mov.u32 %r33,0;\n"
 236          "setp.eq.u32 %r34,%r28,%r33;\n"
 237          "@!%r34 bra $L29;\n"
 238          "mov.u32 %r23,%nctaid.x;\n"
 239          "mov.u32 %r22,%r23;\n"
 240          "bra $L28;\n"
 241          "$L25:\n"
 242          "mov.u32 %r24,%nctaid.y;\n"
 243          "mov.u32 %r22,%r24;\n"
 244          "bra $L28;\n"
 245          "$L26:\n"
 246          "mov.u32 %r25,%nctaid.z;\n"
 247          "mov.u32 %r22,%r25;\n"
 248          "bra $L28;\n"
 249          "$L29:\n"
 250          "{\n"
 251          "{\n"
 252          "call abort;\n"
 253          "}\n"
 254          "}\n"
 255          "$L28:\n"
 256          "mov.u32 %r26,%r22;\n"
 257          "mov.u32 %retval,%r26;\n"
 258          "st.param.u32 [%out_retval],%retval;\n"
 259          "ret;\n"
 260          "}\n"
 261          "// BEGIN GLOBAL FUNCTION DEF: GOACC_get_num_threads\n"
 262          ".visible .func (.param .u32 %out_retval) GOACC_get_num_threads\n"
 263          "{\n"
 264          ".reg .u32 %retval;\n"
 265          ".reg .u64 %hr10;\n"
 266          ".reg .u32 %r22;\n"
 267          ".reg .u32 %r23;\n"
 268          ".reg .u32 %r24;\n"
 269          ".reg .u32 %r25;\n"
 270          ".reg .u32 %r26;\n"
 271          ".reg .u32 %r27;\n"
 272          ".reg .u32 %r28;\n"
 273          ".reg .u32 %r29;\n"
 274          "mov.u32 %r26,0;\n"
 275          "{\n"
 276          ".param .u32 %retval_in;\n"
 277          "{\n"
 278          ".param .u32 %out_arg0;\n"
 279          "st.param.u32 [%out_arg0],%r26;\n"
 280          "call (%retval_in),GOACC_ntid,(%out_arg0);\n"
 281          "}\n"
 282          "ld.param.u32 %r27,[%retval_in];\n"
 283          "}\n"
 284          "mov.u32 %r22,%r27;\n"
 285          "mov.u32 %r28,0;\n"
 286          "{\n"
 287          ".param .u32 %retval_in;\n"
 288          "{\n"
 289          ".param .u32 %out_arg0;\n"
 290          "st.param.u32 [%out_arg0],%r28;\n"
 291          "call (%retval_in),GOACC_nctaid,(%out_arg0);\n"
 292          "}\n"
 293          "ld.param.u32 %r29,[%retval_in];\n"
 294          "}\n"
 295          "mov.u32 %r23,%r29;\n"
 296          "mul.lo.u32 %r24,%r22,%r23;\n"
 297          "mov.u32 %r25,%r24;\n"
 298          "mov.u32 %retval,%r25;\n"
 299          "st.param.u32 [%out_retval],%retval;\n"
 300          "ret;\n"
 301          "}\n"
 302          "// BEGIN GLOBAL FUNCTION DEF: GOACC_get_thread_num\n"
 303          ".visible .func (.param .u32 %out_retval) GOACC_get_thread_num\n"
 304          "{\n"
 305          ".reg .u32 %retval;\n"
 306          ".reg .u64 %hr10;\n"
 307          ".reg .u32 %r22;\n"
 308          ".reg .u32 %r23;\n"
 309          ".reg .u32 %r24;\n"
 310          ".reg .u32 %r25;\n"
 311          ".reg .u32 %r26;\n"
 312          ".reg .u32 %r27;\n"
 313          ".reg .u32 %r28;\n"
 314          ".reg .u32 %r29;\n"
 315          ".reg .u32 %r30;\n"
 316          ".reg .u32 %r31;\n"
 317          ".reg .u32 %r32;\n"
 318          ".reg .u32 %r33;\n"
 319          "mov.u32 %r28,0;\n"
 320          "{\n"
 321          ".param .u32 %retval_in;\n"
 322          "{\n"
 323          ".param .u32 %out_arg0;\n"
 324          "st.param.u32 [%out_arg0],%r28;\n"
 325          "call (%retval_in),GOACC_ntid,(%out_arg0);\n"
 326          "}\n"
 327          "ld.param.u32 %r29,[%retval_in];\n"
 328          "}\n"
 329          "mov.u32 %r22,%r29;\n"
 330          "mov.u32 %r30,0;\n"
 331          "{\n"
 332          ".param .u32 %retval_in;\n"
 333          "{\n"
 334          ".param .u32 %out_arg0;\n"
 335          "st.param.u32 [%out_arg0],%r30;\n"
 336          "call (%retval_in),GOACC_ctaid,(%out_arg0);\n"
 337          "}\n"
 338          "ld.param.u32 %r31,[%retval_in];\n"
 339          "}\n"
 340          "mov.u32 %r23,%r31;\n"
 341          "mul.lo.u32 %r24,%r22,%r23;\n"
 342          "mov.u32 %r32,0;\n"
 343          "{\n"
 344          ".param .u32 %retval_in;\n"
 345          "{\n"
 346          ".param .u32 %out_arg0;\n"
 347          "st.param.u32 [%out_arg0],%r32;\n"
 348          "call (%retval_in),GOACC_tid,(%out_arg0);\n"
 349          "}\n"
 350          "ld.param.u32 %r33,[%retval_in];\n"
 351          "}\n"
 352          "mov.u32 %r25,%r33;\n"
 353          "add.u32 %r26,%r24,%r25;\n"
 354          "mov.u32 %r27,%r26;\n"
 355          "mov.u32 %retval,%r27;\n"
 356          "st.param.u32 [%out_retval],%retval;\n"
 357          "ret;\n"
 358          "}\n");