config_nvidia.hpp.in 2 KB
Newer Older
Chao Liu's avatar
Chao Liu committed
1
2
#ifndef CK_CONFIG_NVIDIA_HPP
#define CK_CONFIG_NVIDIA_HPP
3

Chao Liu's avatar
Chao Liu committed
4
#include "cuda_runtime.h"
Chao Liu's avatar
Chao Liu committed
5
#include "cuda_fp16.h"
Chao Liu's avatar
Chao Liu committed
6
7
#include "nvToolsExt.h"
#include "helper_cuda.h"
Chao Liu's avatar
Chao Liu committed
8

Chao Liu's avatar
Chao Liu committed
9
#define CK_UNSIGNED_INDEX_TYPE 0
Chao Liu's avatar
Chao Liu committed
10
#define CK_DEVICE_BACKEND_NVIDIA 1
Chao Liu's avatar
Chao Liu committed
11
#define CK_USE_AMD_INTRINSIC 0
Chao Liu's avatar
Chao Liu committed
12
#define CK_USE_AMD_INLINE_ASM 0
13
#define CK_USE_AMD_INTRINSIC_BUFFER_LOAD_STORE 0
Chao Liu's avatar
Chao Liu committed
14
#define CK_EXPERIMENTAL_USE_MORE_COMPILE_STATIC_BLOCKWISE_GENERIC_SLICE_COPY_V1 0
15
#define CK_EXPERIMENTAL_USE_MORE_COMPILE_STATIC_THREADWISE_GENERIC_TENSOR_SLICE_COPY_V1R1 0
16
#define CK_EXPERIMENTAL_USE_MORE_COMPILE_STATIC_THREADWISE_GENERIC_TENSOR_SLICE_COPY_V1R2 0
17
#define CK_EXPERIMENTAL_USE_MORE_COMPILE_STATIC_THREADWISE_GENERIC_TENSOR_SLICE_COPY_V2R1 0
Chao Liu's avatar
Chao Liu committed
18

19
20
namespace ck {

21
22
enum address_space_t
{
Chao Liu's avatar
Chao Liu committed
23
24
25
26
    generic = 0,
    vgpr    = generic,
    lds     = generic,
    global  = generic
27
28
};

Chao Liu's avatar
Chao Liu committed
29
#if CK_UNSIGNED_INDEX_TYPE
30
using index_t = uint32_t;
Chao Liu's avatar
Chao Liu committed
31
#else
32
using index_t = int32_t;
Chao Liu's avatar
Chao Liu committed
33
34
#endif

Chao Liu's avatar
Chao Liu committed
35
36
37
38
39
40
// For some reason, CUDA need this definition, otherwise
//   compiler won't generate optimal load and store instruction, and
//   kernel would produce wrong result, indicating the compiler fail to generate correct
//   instruction,
using float2_t = float2;
using float4_t = float4;
Chao Liu's avatar
Chao Liu committed
41

Chao Liu's avatar
Chao Liu committed
42
43
template <class T>
__device__ void fused_multiply_accumulate(T& d, const T& s0, const T& s1)
Chao Liu's avatar
Chao Liu committed
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
{
    d += s0 * s1;
}

#if 0
__device__ void fused_multiply_accumulate(half& d, const half& s0, const half& s1) { d += s0 * s1; }

__device__ void fused_multiply_accumulate(half& d, const half2& s0, const half2& s1)
{
    d += s0.x * s1.x;
    d += s0.y * s1.y;
}

__device__ void fused_multiply_accumulate(float& d, const half2& s0, const half2& s1)
{
    d += s0.x * s1.x + s0.y * s1.y;
}

__device__ void fused_multiply_accumulate(char& d, const char& s0, const char& s1) { d += s0 * s1; }

// TODO:: this interface is misleading, s0, s1 are actually int8x4
//  need to make a better interface
__device__ void fused_multiply_accumulate(int32_t& d, const int32_t& s0, const int32_t& s1)
{
    d = __dp4a(s0, s1, d);
}
#endif
71
72
73
74

} // namespace ck

#endif