src/utils/thread_id.hpp

0001 /**
0002  * traccc library, part of the ACTS project (R&D line)
0003  *
0004  * (c) 2024-2025 CERN for the benefit of the ACTS project
0005  *
0006  * Mozilla Public License Version 2.0
0007  */
0008
0009 #pragma once
0010
0011 // Project include(s).
0012 #include "traccc/device/concepts/thread_id.hpp"
0013
0014 // HIP include(s).
0015 #include <hip/hip_runtime.h>
0016
0017 namespace traccc::hip::details {
0018
0019 /// A HIP thread identifier type
0020 struct thread_id1 {
0021   __device__ thread_id1() {}
0022
0023   inline unsigned int __device__ getLocalThreadId() const {
0024     return threadIdx.x;
0025   }
0026
0027   inline unsigned int __device__ getLocalThreadIdX() const {
0028     return threadIdx.x;
0029   }
0030
0031   inline unsigned int __device__ getGlobalThreadId() const {
0032     return threadIdx.x + blockIdx.x * blockDim.x;
0033   }
0034
0035   inline unsigned int __device__ getGlobalThreadIdX() const {
0036     return threadIdx.x + blockIdx.x * blockDim.x;
0037   }
0038
0039   inline unsigned int __device__ getBlockIdX() const { return blockIdx.x; }
0040
0041   inline unsigned int __device__ getBlockDimX() const { return blockDim.x; }
0042
0043   inline unsigned int __device__ getGridDimX() const { return gridDim.x; }
0044 };
0045
0046 static_assert(traccc::device::concepts::thread_id1<thread_id1>);
0047
0048 }  // namespace traccc::hip::details