ck.hpp 8.55 KB
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

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

6
7
#include "ck/config.h"

JD's avatar
JD committed
8
#ifndef CK_DONT_USE_HIP_RUNTIME_HEADERS
9
10
#include "hip/hip_runtime.h"
#include "hip/hip_fp16.h"
Chao Liu's avatar
Chao Liu committed
11
#endif
12

Chao Liu's avatar
Chao Liu committed
13
14
#define CK_TIME_KERNEL 1

15
16
17
// constant address space for kernel parameter
// https://llvm.org/docs/AMDGPUUsage.html#address-spaces
#define CK_CONSTANT_ADDRESS_SPACE __attribute__((address_space(4)))
Chao Liu's avatar
Chao Liu committed
18
19

// launch bounds
zjing14's avatar
zjing14 committed
20
#define CK_USE_LAUNCH_BOUNDS 1
Chao Liu's avatar
Chao Liu committed
21
22

#ifdef CK_USE_LAUNCH_BOUNDS
23
// for most kernels
Chao Liu's avatar
Chao Liu committed
24
#define CK_MAX_THREAD_PER_BLOCK 256
Chao Liu's avatar
Chao Liu committed
25
#define CK_MIN_BLOCK_PER_CU 2
26
27
28
29

// for wavelet GEMM kernel
#define CK_WAVELET_MAX_THREAD_PER_BLOCK 512
#define CK_WAVELET_MIN_BLOCK_PER_CU 2
Chao Liu's avatar
Chao Liu committed
30
31
#endif

32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
// kernel attribute: amdgpu_waves_per_eu()
#ifdef CK_USE_WAVES_PER_EU
// for 1-wave kernels, control arguments of amdgpu_waves_per_eu() attribute
#ifndef CK_MIN_WAVES_PER_EU
#define CK_MIN_WAVES_PER_EU 0
#endif

#ifndef CK_MAX_WAVES_PER_EU
#define CK_MAX_WAVES_PER_EU 0
#endif

#else
#define CK_USE_WAVES_PER_EU 0
#endif

47
// define general macros for various architectures
48
49
50
51
#if defined(__gfx908__) || defined(__gfx90a__) || defined(__gfx940__) || defined(__gfx941__) || \
    defined(__gfx942__)
#define __gfx9__
#endif
52
53
54
55
56
57
58
59
60
61
62
63
64
65
#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__)
#define __gfx94__
#endif
#if defined(__gfx1010__) || defined(__gfx1011__) || defined(__gfx1012__)
#define __gfx101__
#endif
#if defined(__gfx1030__) || defined(__gfx1031__) || defined(__gfx1032__) || \
    defined(__gfx1034__) || defined(__gfx1035__) || defined(__gfx1036__)
#define __gfx103__
#endif
#if defined(__gfx1100__) || defined(__gfx1101__) || defined(__gfx1102__) || defined(__gfx1103__)
#define __gfx11__
#endif

Chao Liu's avatar
Chao Liu committed
66
// buffer resource
67
68
#ifndef __HIP_DEVICE_COMPILE__ // for host code
#define CK_BUFFER_RESOURCE_3RD_DWORD -1
69
#elif defined(__gfx803__) || defined(__gfx900__) || defined(__gfx906__) || defined(__gfx9__)
Chao Liu's avatar
Chao Liu committed
70
#define CK_BUFFER_RESOURCE_3RD_DWORD 0x00020000
71
#elif defined(__gfx103__)
Chao Liu's avatar
Chao Liu committed
72
#define CK_BUFFER_RESOURCE_3RD_DWORD 0x31014000
73
#elif defined(__gfx11__)
74
#define CK_BUFFER_RESOURCE_3RD_DWORD 0x31004000
Chao Liu's avatar
Chao Liu committed
75
76
#endif

77
// FMA instruction
78
79
#ifndef __HIP_DEVICE_COMPILE__                   // for host code, define nothing
#elif defined(__gfx803__) || defined(__gfx900__) // for GPU code
80
#define CK_USE_AMD_V_MAC_F32
81
#elif defined(__gfx906__) || defined(__gfx9__) || defined(__gfx103__) // for GPU code
82
83
84
#define CK_USE_AMD_V_FMAC_F32
#define CK_USE_AMD_V_DOT2_F32_F16
#define CK_USE_AMD_V_DOT4_I32_I8
85
#elif defined(__gfx11__)
86
87
88
#define CK_USE_AMD_V_FMAC_F32
#define CK_USE_AMD_V_DOT2_F32_F16
#define CK_USE_AMD_V_DOT4_I32_I8_GFX11
89
90
#endif

91
92
93
// MFMA instruction
#ifndef __HIP_DEVICE_COMPILE__ // for host code
#define CK_USE_AMD_MFMA
94
#elif defined(__gfx9__) // for GPU code
95
#define CK_USE_AMD_MFMA
96
97
#endif

98
#if(defined(__gfx90a__) || defined(__gfx94__))
99
#define CK_USE_AMD_MFMA_BF16_1K_OP
Chao Liu's avatar
Chao Liu committed
100
101
#endif

102
#if defined(__gfx94__)
103
104
105
#define CK_USE_AMD_MFMA_GFX940
#endif

106
107
108
// WMMA instruction
#ifndef __HIP_DEVICE_COMPILE__ // for host code
#define CK_USE_AMD_WMMA
109
#elif defined(__gfx11__) // for GPU code
110
111
112
#define CK_USE_AMD_WMMA
#endif

113
// buffer load
Jianfeng Yan's avatar
Jianfeng Yan committed
114
#define CK_USE_AMD_BUFFER_LOAD 1
115

116
// buffer store
Jianfeng Yan's avatar
Jianfeng Yan committed
117
118
#define CK_USE_AMD_BUFFER_STORE 1

119
120
// buffer atomic add: integer
#define CK_USE_AMD_BUFFER_ATOMIC_ADD_INTEGER 1
121

122
123
124
// buffer atomic add: floating point
#ifndef __HIP_DEVICE_COMPILE__ // for host code
#define CK_USE_AMD_BUFFER_ATOMIC_ADD_FLOAT 1
125
#elif defined(__gfx9__) // for GPU code
126
127
128
#define CK_USE_AMD_BUFFER_ATOMIC_ADD_FLOAT 1
#else // for GPU code
#define CK_USE_AMD_BUFFER_ATOMIC_ADD_FLOAT 0
129
130
#endif

131
#if(defined(__gfx90a__) || defined(__gfx94__)) // for GPU code
rocking5566's avatar
rocking5566 committed
132
133
134
135
136
#define CK_USE_AMD_BUFFER_ATOMIC_MAX_FLOAT64 1
#else
#define CK_USE_AMD_BUFFER_ATOMIC_MAX_FLOAT64 0
#endif

137
138
139
// inline asm
#define CK_USE_AMD_INLINE_ASM 1

140
141
142
143
144
145
// inner product (V_MAC/V_FMAC)
#define CK_USE_AMD_V_MAC_INLINE_ASM 1

// V_DOT inline instructions, less efficient since they require adding
// `s_nop`s to avoid hazard
#define CK_USE_AMD_V_DOT_INLINE_ASM 0
146

147
148
149
// inner product using V_DOT with DPP8 modifiers
#define CK_USE_AMD_V_DOT_DPP8_INLINE_ASM 1

150
151
152
// LDS direct loads using inline assembly
#define CK_USE_AMD_LDS_DIRECT_LOAD_INLINE_ASM 1

153
154
155
// set stochastic rounding as default for f8 conversions
#define CK_USE_SR_F8_CONVERSION 1

Chao Liu's avatar
Chao Liu committed
156
// block synchronization only s_wait lgkmcnt(0), not vmcnt(0)
157
#define CK_EXPERIMENTAL_BLOCK_SYNC_LDS_WITHOUT_SYNC_VMEM 1
Chao Liu's avatar
Chao Liu committed
158

159
160
// experimental feature: multi index implemented as array
#define CK_EXPERIMENTAL_USE_DYNAMICALLY_INDEXED_MULTI_INDEX 0
Chao Liu's avatar
Chao Liu committed
161

162
163
// experimental feature: static tensor descriptor
#define CK_EXPERIMENTAL_STATIC_TENSOR_DESCRIPTOR 0
Chao Liu's avatar
Chao Liu committed
164

rocking5566's avatar
rocking5566 committed
165
// experimental feature: buffer load/store/atomic-add/ OOB trick
166
// This (ifndef) is a hack to use customized behavior for buffer load rather than using default
Po Yen Chen's avatar
Po Yen Chen committed
167
168
// setting. Don't use this hack unless absolutely necessary!
// FIXME: make the behavior of buffer load a configurable (template) parameter for each usage
169
#ifndef CK_EXPERIMENTAL_USE_BUFFER_LOAD_OOB_CHECK_OFFSET_TRICK
170
#define CK_EXPERIMENTAL_USE_BUFFER_LOAD_OOB_CHECK_OFFSET_TRICK 0
171
#endif
172
#define CK_EXPERIMENTAL_USE_BUFFER_STORE_OOB_CHECK_OFFSET_TRICK 1
zjing14's avatar
zjing14 committed
173
#define CK_EXPERIMENTAL_USE_BUFFER_ATOMIC_ADD_OOB_CHECK_OFFSET_TRICK 1
rocking5566's avatar
rocking5566 committed
174
#define CK_EXPERIMENTAL_USE_BUFFER_ATOMIC_MAX_OOB_CHECK_OFFSET_TRICK 1
175

176
// experimental feature: in-regsiter sub-dword transpose
177
#define CK_EXPERIMENTAL_USE_IN_REGISTER_SUB_DWORD_TRANSPOSE 1
Chao Liu's avatar
Chao Liu committed
178

179
// experimental feature: merge transformation use magic number division
180
#define CK_EXPERIMENTAL_MERGE_USE_MAGIC_DIVISION 1
181

182
183
// experimental feature: use __builtin_memcpy instead of pointer cast to access a vector from
// pointer of scalar
184
185
#define CK_EXPERIMENTAL_USE_MEMCPY_FOR_VECTOR_ACCESS 0

186
// experimental feature: use __builtin_memcpy instead of union to do bit_cast
187
#define CK_EXPERIMENTAL_USE_MEMCPY_FOR_BIT_CAST 1
188

189
// experimental feature: optimize for inter-wave scheduling policy
190
#define CK_EXPERIMENTAL_INTER_WAVE_SCHEDULING 1
191
#define CK_EXPERIMENTAL_INTER_WAVE_SCHEDULING_MAC_CLUSTERS 1
192
193
194
195
196
197
// this will let make_default_loop_scheduler() return interwave scheduling flag by default
#define CK_EXPERIMENTAL_DEFAULT_TO_INTER_WAVE_SCHEDULING 0
// experimental feature: add instances using interwave scheduling
#define CK_EXPERIMENTAL_INTER_WAVE_INSTANCES 1
// experimental feature: add instances using pipeline v2
#define CK_EXPERIMENTAL_PIPELINE_V2_INSTANCES 1
198
199
200
201
// experimental feature: optimize pipeline v2 by IGLP strategy (value=ID of strategy)
#ifndef CK_EXPERIMENTAL_PIPELINE_V2_IGLP_OPT
#define CK_EXPERIMENTAL_PIPELINE_V2_IGLP_OPT 0
#endif
202

Chao Liu's avatar
Chao Liu committed
203
204
205
206
// hack: have underlying assumption that need to be satsified, otherwise it's a bug
// hack for forcing register to keep idx_diff_low_const in SGPR. idx_diff_low_const must be
// thread-invariant, otherwise it's a bug
// TODO: separate index calculation into "compile-time", "global", "block", "wave", "thread"
Chao Liu's avatar
rename  
Chao Liu committed
207
#define CK_HACK_MERGE_CALCULATE_IDX_DIFF_LOW_CONST_USE_AMD_GCN_READ_FIRST_LANE 0
Chao Liu's avatar
Chao Liu committed
208

209
// workaround: compiler crash when compiling recursive lambda
Chao Liu's avatar
Chao Liu committed
210
#define CK_WORKAROUND_SWDEV_275126 1
211

212
// workaround: compiler crash when using buffer load/store for i8
213
214
#define CK_WORKAROUND_SWDEV_XXXXXX_INT8_BUFFER_LOAD_STORE_ISSUE 1

215
// workaround: compiler gnerating inefficient ds_write instructions
216
#define CK_WORKAROUND_SWDEV_XXXXXX_INT8_DS_WRITE_ISSUE 1
Chao Liu's avatar
Chao Liu committed
217

218
// workaround: verifaction failure, due to compiler regression, for conv bwd-data fp16 using some
219
// tuning parameter
220
#define CK_WORKAROUND_SWDEV_325164 0
221

222
223
224
// workaround: compiler not emiting reciprocal instruction frm __frcp_rn()
#define CK_WORKAROUND_SWDEV_383542 1

ltqin's avatar
ltqin committed
225
226
// workaround: compiler issue on gfx908
#define CK_WORKAROUND_SWDEV_388832 1
227

228
229
230
// flag to enable (1) or disable (0) the debugging output in some kernels
#define DEBUG_LOG 0

231
232
233
// denorm test fix, required to work around dissue
#ifndef CK_WORKAROUND_DENORM_FIX
#define CK_WORKAROUND_DENORM_FIX 0
Illia Silin's avatar
Illia Silin committed
234
#else
235
236
237
// enable only on MI200
#define CK_WORKAROUND_DENORM_FIX = CK_WORKAROUND_DENORM_FIX && defined(__gfx90a__)
#endif // CK_WORKAROUND_DENORM_FIX
238

239
240
namespace ck {

241
enum struct InMemoryDataOperationEnum
Chao Liu's avatar
Chao Liu committed
242
{
Chao Liu's avatar
Chao Liu committed
243
    Set,
244
    AtomicAdd,
rocking5566's avatar
rocking5566 committed
245
    AtomicMax,
246
247
248
    Add
};

Chao Liu's avatar
Chao Liu committed
249
// FIXME: use regular Sequence and remove this
rocking5566's avatar
rocking5566 committed
250
251
252
253
254
255
256
257
258
259
260
261
262
template <InMemoryDataOperationEnum... Is>
struct InMemoryDataOperationEnumSequence
{
    static constexpr int mSize = sizeof...(Is);

    __host__ __device__ static constexpr InMemoryDataOperationEnum At(int I)
    {
        // the last dummy element is to prevent compiler complain about empty array, when mSize = 0
        const InMemoryDataOperationEnum mData[mSize + 1] = {Is..., InMemoryDataOperationEnum::Set};
        return mData[I];
    }
};

Chao Liu's avatar
Chao Liu committed
263
// index type
Jianfeng Yan's avatar
Jianfeng Yan committed
264
265
using index_t      = int32_t;
using long_index_t = int64_t;
266
267

} // namespace ck