/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck/utility/get_id.hpp Source File

/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck/utility/get_id.hpp Source File#

Composable Kernel: /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck/utility/get_id.hpp Source File
get_id.hpp
Go to the documentation of this file.
1 // SPDX-License-Identifier: MIT
2 // Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
3 
4 #pragma once
5 
6 #include "ck/ck.hpp"
7 
8 namespace ck {
9 
10 __device__ constexpr index_t get_warp_size()
11 {
12 #if defined(__HIP_DEVICE_COMPILE__)
13 #if defined(__GFX9__)
14  return 64;
15 #else
16  return 32;
17 #endif
18 #else
19  return 64;
20 #endif
21 }
22 
23 inline __host__ index_t get_warp_size()
24 {
25 #if !(defined(__HIPCC_RTC__) || defined(CK_CODE_GEN_RTC))
26  int device = 0;
27  int result = 0;
28  auto status = hipGetDevice(&device);
29  if(status == hipSuccess)
30  {
31  status = hipDeviceGetAttribute(&result, hipDeviceAttributeWarpSize, device);
32  if(status == hipSuccess)
33  {
34  return result;
35  }
36  }
37 #endif
38  return 64;
39 }
40 
41 __device__ index_t get_thread_local_1d_id() { return threadIdx.x; }
42 
43 __device__ index_t get_thread_global_1d_id() { return blockIdx.x * blockDim.x + threadIdx.x; }
44 
45 __device__ index_t get_warp_local_1d_id() { return threadIdx.x / get_warp_size(); }
46 
47 __device__ index_t get_block_1d_id() { return blockIdx.x; }
48 
49 __device__ index_t get_grid_size() { return gridDim.x; }
50 
51 __device__ index_t get_block_size() { return blockDim.x; }
52 
53 } // namespace ck
Definition: ck.hpp:268
__device__ index_t get_warp_local_1d_id()
Definition: get_id.hpp:45
__device__ index_t get_grid_size()
Definition: get_id.hpp:49
__device__ index_t get_block_size()
Definition: get_id.hpp:51
__device__ index_t get_block_1d_id()
Definition: get_id.hpp:47
__device__ index_t get_thread_global_1d_id()
Definition: get_id.hpp:43
constexpr __device__ index_t get_warp_size()
Definition: get_id.hpp:10
int32_t index_t
Definition: ck.hpp:299
__device__ index_t get_thread_local_1d_id()
Definition: get_id.hpp:41