get_id.hpp 937 Bytes
Newer Older
Chao Liu's avatar
Chao Liu committed
1
// SPDX-License-Identifier: MIT
Illia Silin's avatar
Illia Silin committed
2
// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved.
Chao Liu's avatar
Chao Liu committed
3

4
#pragma once
Chao Liu's avatar
Chao Liu committed
5
6

#include "ck/ck.hpp"
Chao Liu's avatar
Chao Liu committed
7

8
namespace ck {
Chao Liu's avatar
Chao Liu committed
9

Chao Liu's avatar
Chao Liu committed
10
11
12
13
14
__host__ __device__ constexpr index_t get_warp_size()
{
    // warpSize is defined by HIP
    return warpSize;
}
Chao Liu's avatar
Chao Liu committed
15

Chao Liu's avatar
Chao Liu committed
16
17
18
19
20
__device__ index_t get_grid_size() { return gridDim.x; }

__device__ index_t get_block_size() { return blockDim.x; }

// TODO: deprecate these
Chao Liu's avatar
Chao Liu committed
21
__device__ index_t get_thread_local_1d_id() { return threadIdx.x; }
Chao Liu's avatar
Chao Liu committed
22

rocking5566's avatar
rocking5566 committed
23
24
__device__ index_t get_thread_global_1d_id() { return blockIdx.x * blockDim.x + threadIdx.x; }

Chao Liu's avatar
Chao Liu committed
25
__device__ index_t get_block_1d_id() { return blockIdx.x; }
Chao Liu's avatar
Chao Liu committed
26

Chao Liu's avatar
Chao Liu committed
27
28
// Use these instead
__device__ index_t get_lane_id() { return __lane_id(); }
Jianfeng Yan's avatar
Jianfeng Yan committed
29

Chao Liu's avatar
Chao Liu committed
30
31
32
33
34
__device__ index_t get_warp_id() { return threadIdx.x / get_warp_size(); }

__device__ index_t get_thread_id() { return threadIdx.x; }

__device__ index_t get_block_id() { return blockIdx.x; }
rocking5566's avatar
rocking5566 committed
35

Chao Liu's avatar
Chao Liu committed
36
} // namespace ck