get_id.hpp 480 Bytes
Newer Older
1
#pragma once
Chao Liu's avatar
Chao Liu committed
2
3
#include "config.hpp"

4
#ifndef CK_NOGPU
5
namespace ck {
Chao Liu's avatar
Chao Liu committed
6

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

Chao Liu's avatar
Chao Liu committed
13
__device__ index_t get_thread_local_1d_id() { return threadIdx.x; }
Chao Liu's avatar
Chao Liu committed
14

Chao Liu's avatar
Chao Liu committed
15
__device__ index_t get_warp_local_1d_id() { return threadIdx.x / get_warp_size(); }
Chao Liu's avatar
Chao Liu committed
16

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

Jianfeng Yan's avatar
Jianfeng Yan committed
19
20
__device__ index_t get_grid_size() { return gridDim.x; }

Chao Liu's avatar
Chao Liu committed
21
} // namespace ck
22
#endif