openmp/libomptarget/deviceRTLs/amdgcn/src/amdgcn_smid.hip

   1 //===-------- amdgcn_smid.hip - AMDGCN smid implementation -------- HIP -*-===//
   2 //
   3 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
   4 // See https://llvm.org/LICENSE.txt for license information.
   5 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
   6 //
   7 //===----------------------------------------------------------------------===//
   8 #pragma omp declare target
   9
  10 #include "target_impl.h"
  11
  12 // Partially derived fom hcc_detail/device_functions.h
  13
  14 // HW_ID Register bit structure
  15 // WAVE_ID     3:0     Wave buffer slot number. 0-9.
  16 // SIMD_ID     5:4     SIMD which the wave is assigned to within the CU.
  17 // PIPE_ID     7:6     Pipeline from which the wave was dispatched.
  18 // CU_ID       11:8    Compute Unit the wave is assigned to.
  19 // SH_ID       12      Shader Array (within an SE) the wave is assigned to.
  20 // SE_ID       14:13   Shader Engine the wave is assigned to.
  21 // TG_ID       19:16   Thread-group ID
  22 // VM_ID       23:20   Virtual Memory ID
  23 // QUEUE_ID    26:24   Queue from which this wave was dispatched.
  24 // STATE_ID    29:27   State ID (graphics only, not compute).
  25 // ME_ID       31:30   Micro-engine ID.
  26
  27 enum {
  28   HW_ID = 4, // specify that the hardware register to read is HW_ID
  29
  30   HW_ID_CU_ID_SIZE = 4,   // size of CU_ID field in bits
  31   HW_ID_CU_ID_OFFSET = 8, // offset of CU_ID from start of register
  32
  33   HW_ID_SE_ID_SIZE = 2,    // sizeof SE_ID field in bits
  34   HW_ID_SE_ID_OFFSET = 13, // offset of SE_ID from start of register
  35 };
  36
  37 // The s_getreg_b32 instruction, exposed as an intrinsic, takes a 16 bit
  38 // immediate and returns a 32 bit value.
  39 // The encoding of the immediate parameter is:
  40 // ID           5:0     Which register to read from
  41 // OFFSET       10:6    Range: 0..31
  42 // WIDTH        15:11   Range: 1..32
  43
  44 // The asm equivalent is s_getreg_b32 %0, hwreg(HW_REG_HW_ID, Offset, Width)
  45 // where hwreg forms a 16 bit immediate encoded by the assembler thus:
  46 // uint64_t encodeHwreg(uint64_t Id, uint64_t Offset, uint64_t Width) {
  47 //   return (Id << 0_) | (Offset << 6) | ((Width - 1) << 11);
  48 // }
  49 #define ENCODE_HWREG(WIDTH, OFF, REG) (REG | (OFF << 6) | ((WIDTH - 1) << 11))
  50
  51 // Note: The results can be changed by a context switch
  52 // Return value in [0 2^SE_ID_SIZE * 2^CU_ID_SIZE), which is an upper
  53 // bound on how many compute units are available. Some values in this
  54 // range may never be returned if there are fewer than 2^CU_ID_SIZE CUs.
  55
  56 EXTERN uint32_t __kmpc_impl_smid() {
  57   uint32_t cu_id = __builtin_amdgcn_s_getreg(
  58       ENCODE_HWREG(HW_ID_CU_ID_SIZE, HW_ID_CU_ID_OFFSET, HW_ID));
  59   uint32_t se_id = __builtin_amdgcn_s_getreg(
  60       ENCODE_HWREG(HW_ID_SE_ID_SIZE, HW_ID_SE_ID_OFFSET, HW_ID));
  61   return (se_id << HW_ID_CU_ID_SIZE) + cu_id;
  62 }
  63
  64 #pragma omp end declare target