"vscode:/vscode.git/clone" did not exist on "f0e7fa62d920ee9a9875d12ce8825cb06a87bbff"
config_amd.hpp.in 1.1 KB
Newer Older
Chao Liu's avatar
Chao Liu committed
1
2
#ifndef CK_CONFIG_AMD_HPP
#define CK_CONFIG_AMD_HPP
Chao Liu's avatar
Chao Liu committed
3
4
5
6

#include "hip/hip_runtime.h"
#include "hip/hip_fp16.h"

Chao Liu's avatar
Chao Liu committed
7
8
#define CK_DEVICE_BACKEND_AMD 1
#define CK_USE_AMD_INLINE_ASM 1
Chao Liu's avatar
Chao Liu committed
9
10
#define CK_EXPERIMENTAL_USE_MORE_COMPILE_STATIC_BLOCKWISE_GENERIC_SLICE_COPY_V1 1
#define CK_EXPERIMENTAL_USE_MORE_COMPILE_STATIC_THREADWISE_GENERIC_TENSOR_SLICE_COPY_V1 0
11
#define CK_EXPERIMENTAL_USE_MORE_COMPILE_STATIC_THREADWISE_GENERIC_TENSOR_SLICE_COPY_V1R1 0
12
#define CK_EXPERIMENTAL_USE_MORE_COMPILE_STATIC_THREADWISE_GENERIC_TENSOR_SLICE_COPY_V1R2 0
13
#define CK_EXPERIMENTAL_USE_MORE_COMPILE_STATIC_THREADWISE_GENERIC_TENSOR_SLICE_COPY_V2 0
14
#define CK_EXPERIMENTAL_USE_MORE_COMPILE_STATIC_THREADWISE_GENERIC_TENSOR_SLICE_COPY_V2R1 0
Chao Liu's avatar
Chao Liu committed
15

Chao Liu's avatar
Chao Liu committed
16
17
18
19
20
21
22
23
24
namespace ck {

// For some reason, HIP compiler need this definition to generate optimal load and store
// instruction
typedef float float2_t __attribute__((ext_vector_type(2)));
typedef float float4_t __attribute__((ext_vector_type(4)));

using index_t = uint32_t;

Chao Liu's avatar
Chao Liu committed
25
26
template <class T>
__device__ void fused_multiply_accumulate(T& d, const T& s0, const T& s1)
Chao Liu's avatar
Chao Liu committed
27
28
29
30
31
32
33
{
    d += s0 * s1;
}

} // namespace ck

#endif