ck.hpp 6.1 KB
Newer Older
Chao Liu's avatar
Chao Liu committed
1
2
3
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2022, Advanced Micro Devices, Inc. All rights reserved.

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

6
7
8
#ifdef CK_NOGPU
#define __host__
#define __device__
carlushuang's avatar
carlushuang committed
9
#include <stdint.h>
10
#else
11
12
#include "hip/hip_runtime.h"
#include "hip/hip_fp16.h"
Chao Liu's avatar
Chao Liu committed
13
#endif
14

Chao Liu's avatar
Chao Liu committed
15
16
#define CK_TIME_KERNEL 1

17
18
19
// 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
20
21

// launch bounds
zjing14's avatar
zjing14 committed
22
#define CK_USE_LAUNCH_BOUNDS 1
Chao Liu's avatar
Chao Liu committed
23
24
25

#ifdef CK_USE_LAUNCH_BOUNDS
#define CK_MAX_THREAD_PER_BLOCK 256
Chao Liu's avatar
Chao Liu committed
26
#define CK_MIN_BLOCK_PER_CU 2
Chao Liu's avatar
Chao Liu committed
27
28
#endif

29
30
31
32
33
34
35
36
// check GPU target
#ifdef __HIP_DEVICE_COMPILE__
#if !(defined(__gfx803__) || defined(__gfx900__) || defined(__gfx906__) || defined(__gfx908__) || \
      defined(__gfx90a__) || defined(__gfx1030__))
#error Not supported target
#endif
#endif

37
38
39
40
41
42
#if defined(__GNUC__) && !defined(__clang__) && !defined(__llvm__)
#if __GNUC__ < 9
#error "If use gcc, need make sure use at least gcc-9"
#endif
#endif

Chao Liu's avatar
Chao Liu committed
43
// buffer resource
44
45
46
47
#ifndef __HIP_DEVICE_COMPILE__ // for host code
#define CK_BUFFER_RESOURCE_3RD_DWORD -1
#elif defined(__gfx803__) || defined(__gfx900__) || defined(__gfx906__) || defined(__gfx908__) || \
    defined(__gfx90a__) // for GPU code
Chao Liu's avatar
Chao Liu committed
48
#define CK_BUFFER_RESOURCE_3RD_DWORD 0x00020000
49
#elif defined(__gfx1030__) // for GPU code
Chao Liu's avatar
Chao Liu committed
50
51
52
#define CK_BUFFER_RESOURCE_3RD_DWORD 0x31014000
#endif

53
// FMA instruction
54
55
#ifndef __HIP_DEVICE_COMPILE__                   // for host code, define nothing
#elif defined(__gfx803__) || defined(__gfx900__) // for GPU code
56
#define CK_USE_AMD_V_MAC_F32
57
58
#elif defined(__gfx906__) || defined(__gfx908__) || defined(__gfx90a__) || \
    defined(__gfx1030__) // for GPU code
59
60
61
62
63
#define CK_USE_AMD_V_FMAC_F32
#define CK_USE_AMD_V_DOT2_F32_F16
#define CK_USE_AMD_V_DOT4_I32_I8
#endif

64
65
66
67
68
// MFMA instruction
#ifndef __HIP_DEVICE_COMPILE__ // for host code
#define CK_USE_AMD_MFMA
#elif defined(__gfx908__) || defined(__gfx90a__) // for GPU code
#define CK_USE_AMD_MFMA
69
70
#endif

71
72
#if defined(__gfx90a__)
#define CK_USE_AMD_MFMA_BF16_1K_OP
Chao Liu's avatar
Chao Liu committed
73
74
#endif

75
// buffer load
Jianfeng Yan's avatar
Jianfeng Yan committed
76
#define CK_USE_AMD_BUFFER_LOAD 1
77

78
// buffer store
Jianfeng Yan's avatar
Jianfeng Yan committed
79
80
#define CK_USE_AMD_BUFFER_STORE 1

81
82
// buffer atomic add: integer
#define CK_USE_AMD_BUFFER_ATOMIC_ADD_INTEGER 1
83

84
85
86
87
88
89
90
// buffer atomic add: floating point
#ifndef __HIP_DEVICE_COMPILE__ // for host code
#define CK_USE_AMD_BUFFER_ATOMIC_ADD_FLOAT 1
#elif defined(__gfx908__) || defined(__gfx90a__) // for GPU code
#define CK_USE_AMD_BUFFER_ATOMIC_ADD_FLOAT 1
#else // for GPU code
#define CK_USE_AMD_BUFFER_ATOMIC_ADD_FLOAT 0
91
92
#endif

rocking5566's avatar
rocking5566 committed
93
94
95
96
97
98
#if defined(__gfx90a__) // for GPU code
#define CK_USE_AMD_BUFFER_ATOMIC_MAX_FLOAT64 1
#else
#define CK_USE_AMD_BUFFER_ATOMIC_MAX_FLOAT64 0
#endif

99
100
101
102
103
104
// inline asm
#define CK_USE_AMD_INLINE_ASM 1

// inner product (DLOP)
#define CK_USE_AMD_INNER_PRODUCT_INLINE_ASM 1

Chao Liu's avatar
Chao Liu committed
105
// block synchronization only s_wait lgkmcnt(0), not vmcnt(0)
106
#define CK_EXPERIMENTAL_BLOCK_SYNC_LDS_WITHOUT_SYNC_VMEM 1
Chao Liu's avatar
Chao Liu committed
107

108
109
// experimental feature: multi index implemented as array
#define CK_EXPERIMENTAL_USE_DYNAMICALLY_INDEXED_MULTI_INDEX 0
Chao Liu's avatar
Chao Liu committed
110

111
112
// experimental feature: static tensor descriptor
#define CK_EXPERIMENTAL_STATIC_TENSOR_DESCRIPTOR 0
Chao Liu's avatar
Chao Liu committed
113

rocking5566's avatar
rocking5566 committed
114
// experimental feature: buffer load/store/atomic-add/ OOB trick
115
// 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
116
117
// setting. Don't use this hack unless absolutely necessary!
// FIXME: make the behavior of buffer load a configurable (template) parameter for each usage
118
#ifndef CK_EXPERIMENTAL_USE_BUFFER_LOAD_OOB_CHECK_OFFSET_TRICK
119
#define CK_EXPERIMENTAL_USE_BUFFER_LOAD_OOB_CHECK_OFFSET_TRICK 0
120
#endif
121
#define CK_EXPERIMENTAL_USE_BUFFER_STORE_OOB_CHECK_OFFSET_TRICK 1
zjing14's avatar
zjing14 committed
122
#define CK_EXPERIMENTAL_USE_BUFFER_ATOMIC_ADD_OOB_CHECK_OFFSET_TRICK 1
rocking5566's avatar
rocking5566 committed
123
#define CK_EXPERIMENTAL_USE_BUFFER_ATOMIC_MAX_OOB_CHECK_OFFSET_TRICK 1
124

125
// experimental feature: in-regsiter sub-dword transpose
126
#define CK_EXPERIMENTAL_USE_IN_REGISTER_SUB_DWORD_TRANSPOSE 1
Chao Liu's avatar
Chao Liu committed
127

128
// experimental feature: merge transformation use magic number division
129
#define CK_EXPERIMENTAL_MERGE_USE_MAGIC_DIVISION 1
130

131
132
// experimental feature: use __builtin_memcpy instead of pointer cast to access a vector from
// pointer of scalar
133
134
#define CK_EXPERIMENTAL_USE_MEMCPY_FOR_VECTOR_ACCESS 0

135
// experimental feature: use __builtin_memcpy instead of union to do bit_cast
136
#define CK_EXPERIMENTAL_USE_MEMCPY_FOR_BIT_CAST 1
137

138
139
140
141
// experimental feature: optimize for inter-wave scheduling policy
#define CK_EXPERIMENTAL_INTER_WAVE_SCHEDULING 0
#define CK_EXPERIMENTAL_INTER_WAVE_SCHEDULING_MAC_CLUSTERS 1

Chao Liu's avatar
Chao Liu committed
142
143
144
145
// 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
146
#define CK_HACK_MERGE_CALCULATE_IDX_DIFF_LOW_CONST_USE_AMD_GCN_READ_FIRST_LANE 0
Chao Liu's avatar
Chao Liu committed
147

148
// workaround: compiler crash when compiling recursive lambda
Chao Liu's avatar
Chao Liu committed
149
#define CK_WORKAROUND_SWDEV_275126 1
150

151
// workaround: compiler crash when using buffer load/store for i8
152
153
#define CK_WORKAROUND_SWDEV_XXXXXX_INT8_BUFFER_LOAD_STORE_ISSUE 1

154
// workaround: compiler gnerating inefficient ds_write instructions
155
#define CK_WORKAROUND_SWDEV_XXXXXX_INT8_DS_WRITE_ISSUE 1
Chao Liu's avatar
Chao Liu committed
156

157
// workaround: verifaction failure, due to compiler regression, for conv bwd-data fp16 using some
158
159
160
// tuning parameter
#define CK_WORKAROUND_SWDEV_325164 1

carlushuang's avatar
carlushuang committed
161
162
163
164
#ifndef CK_USE_X86_INLINE_ASM
#define CK_USE_X86_INLINE_ASM 1
#endif

165
166
namespace ck {

167
enum struct InMemoryDataOperationEnum
Chao Liu's avatar
Chao Liu committed
168
{
Chao Liu's avatar
Chao Liu committed
169
    Set,
170
    AtomicAdd,
rocking5566's avatar
rocking5566 committed
171
    AtomicMax,
172
173
174
    Add
};

Chao Liu's avatar
Chao Liu committed
175
// FIXME: use regular Sequence and remove this
rocking5566's avatar
rocking5566 committed
176
177
178
179
180
181
182
183
184
185
186
187
188
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
189
// index type
Jianfeng Yan's avatar
Jianfeng Yan committed
190
191
using index_t      = int32_t;
using long_index_t = int64_t;
192
193

} // namespace ck