device.hpp 2.54 KB
Newer Older
Chao Liu's avatar
Chao Liu committed
1
2
#ifndef DEVICE_HPP
#define DEVICE_HPP
3

Chao Liu's avatar
Chao Liu committed
4
#include <memory>
Chao Liu's avatar
Chao Liu committed
5
#include "config.hpp"
6

Chao Liu's avatar
Chao Liu committed
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
struct DeviceMem
{
    DeviceMem() = delete;
    DeviceMem(std::size_t mem_size);
    void* GetDeviceBuffer();
    void ToDevice(const void* p);
    void FromDevice(void* p);
    ~DeviceMem();

    void* mpDeviceBuf;
    std::size_t mMemSize;
};

struct KernelTimerImpl;

struct KernelTimer
{
    KernelTimer();
    ~KernelTimer();
    void Start();
    void End();
    float GetElapsedTime() const;

    std::unique_ptr<KernelTimerImpl> impl;
};

33
34
35
#if CK_DEVICE_BACKEND_AMD
using device_stream_t = hipStream_t;

Chao Liu's avatar
Chao Liu committed
36
template <typename... Args, typename F>
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
void launch_kernel(F kernel,
                   dim3 grid_dim,
                   dim3 block_dim,
                   std::size_t lds_byte,
                   hipStream_t stream_id,
                   Args... args)
{
    hipLaunchKernelGGL(kernel, grid_dim, block_dim, lds_byte, stream_id, args...);
}

template <typename... Args, typename F>
float launch_and_time_kernel(F kernel,
                             dim3 grid_dim,
                             dim3 block_dim,
                             std::size_t lds_byte,
                             hipStream_t stream_id,
                             Args... args)
Chao Liu's avatar
Chao Liu committed
54
55
56
57
58
{
    KernelTimer timer;

    timer.Start();

59
    hipLaunchKernelGGL(kernel, grid_dim, block_dim, lds_byte, stream_id, args...);
Chao Liu's avatar
Chao Liu committed
60
61
62

    timer.End();

Chao Liu's avatar
Chao Liu committed
63
    hipGetLastError();
64
65
66
67

    return timer.GetElapsedTime();
}

Chao Liu's avatar
Chao Liu committed
68
#elif CK_DEVICE_BACKEND_NVIDIA
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
using device_stream_t = cudaStream_t;

template <typename... Args, typename F>
void launch_kernel(F kernel,
                   dim3 grid_dim,
                   dim3 block_dim,
                   std::size_t lds_byte,
                   cudaStream_t stream_id,
                   Args... args)
{
    const void* f  = reinterpret_cast<const void*>(kernel);
    void* p_args[] = {&args...};

    cudaError_t error = cudaLaunchKernel(f, grid_dim, block_dim, p_args, lds_byte, stream_id);
}

template <typename... Args, typename F>
float launch_and_time_kernel(F kernel,
                             dim3 grid_dim,
                             dim3 block_dim,
                             std::size_t lds_byte,
                             cudaStream_t stream_id,
                             Args... args)
{
    KernelTimer timer;

95
96
    const void* f  = reinterpret_cast<const void*>(kernel);
    void* p_args[] = {&args...};
Chao Liu's avatar
Chao Liu committed
97
98
99

    timer.Start();

100
    cudaError_t error = cudaLaunchKernel(f, grid_dim, block_dim, p_args, lds_byte, stream_id);
Chao Liu's avatar
Chao Liu committed
101
102
103
104
105

    timer.End();

    return timer.GetElapsedTime();
}
106
#endif
107
108

#endif