Skip to content

GitLab

"vscode:/vscode.git/clone" did not exist on "e4a046e7d863cfe0365f2b334c7766fc0e5de203"
Switch branch/tag
Find file Normal viewHistoryPermalink
device_implicit_gemm_convolution.cuh 3.39 KB
Newer Older
Chao Liu's avatar
adding implicit gemm  
Chao Liu committed
1 
#pragma once
Chao Liu's avatar
adding implicit gemm  
Chao Liu committed
2
3 
#include "gridwise_implicit_gemm_convolution_nchw_kcsr.cuh"
#include "gridwise_implicit_gemm_convolution_nchw_srck.cuh"
Chao Liu's avatar
adding implicit gemm  
Chao Liu committed
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28 

template <class T, class InDesc, class WeiDesc, class OutDesc>
void device_implicit_gemm_convolution(
    InDesc, const Tensor<T>& in, WeiDesc, const Tensor<T>& wei, OutDesc, Tensor<T>& out)
{
    std::size_t data_sz = sizeof(T);
    DeviceMem in_device_buf(data_sz * in.mDesc.GetElementSpace());
    DeviceMem wei_device_buf(data_sz * wei.mDesc.GetElementSpace());
    DeviceMem out_device_buf(data_sz * out.mDesc.GetElementSpace());

    int num_thread = std::thread::hardware_concurrency();

    in_device_buf.ToDevice(in.mData.data());
    wei_device_buf.ToDevice(wei.mData.data());
    out_device_buf.ToDevice(out.mData.data());

    constexpr auto I0 = Number<0>{};
    constexpr auto I1 = Number<1>{};
    constexpr auto I2 = Number<2>{};
    constexpr auto I3 = Number<3>{};

    constexpr auto in_desc  = InDesc{};
    constexpr auto wei_desc = WeiDesc{};
    constexpr auto out_desc = OutDesc{};
Chao Liu's avatar
adding implicit gemm  
Chao Liu committed
29 
#if 0
Chao Liu's avatar
adding implicit gemm  
Chao Liu committed
30 
    constexpr unsigned NPerBlock  = 2;
Chao Liu's avatar
adding implicit gemm  
Chao Liu committed
31 
    constexpr unsigned KPerBlock  = 64;
Chao Liu's avatar
adding implicit gemm  
Chao Liu committed
32
33
34 
    constexpr unsigned CPerBlock  = 4;
    constexpr unsigned HoPerBlock = 2;
    constexpr unsigned WoPerBlock = 32;
Chao Liu's avatar
adding implicit gemm  
Chao Liu committed
35
Chao Liu's avatar
adding implicit gemm  
Chao Liu committed
36
37
38 
    constexpr unsigned NPerThread  = 2;
    constexpr unsigned KPerThread  = 8;
    constexpr unsigned CPerThread  = 2;
Chao Liu's avatar
adding implicit gemm  
Chao Liu committed
39
40 
    constexpr unsigned HoPerThread = 1;
    constexpr unsigned WoPerThread = 4;
Chao Liu's avatar
adding implicit gemm  
Chao Liu committed
41
42 

    constexpr unsigned BlockSize = 256;
Chao Liu's avatar
adding implicit gemm  
Chao Liu committed
43
44
45
46
47
48
49
50
51
52
53
54
55
56 
#elif 1
    constexpr unsigned NPerBlock  = 2;
    constexpr unsigned KPerBlock  = 32;
    constexpr unsigned CPerBlock  = 4;
    constexpr unsigned HoPerBlock = 2;
    constexpr unsigned WoPerBlock = 32;

    constexpr unsigned NPerThread  = 2;
    constexpr unsigned KPerThread  = 4;
    constexpr unsigned CPerThread  = 2;
    constexpr unsigned HoPerThread = 1;
    constexpr unsigned WoPerThread = 2;

    constexpr unsigned BlockSize = 128;
Chao Liu's avatar
adding implicit gemm  
Chao Liu committed
57
58 
#endif
Chao Liu's avatar
adding implicit gemm  
Chao Liu committed
59
60
61 
    constexpr unsigned GridSize =
        (out_desc.GetLength(I0) / NPerBlock) * (out_desc.GetLength(I1) / KPerBlock) *
        (out_desc.GetLength(I2) / HoPerBlock) * (out_desc.GetLength(I3) / WoPerBlock);
Chao Liu's avatar
adding implicit gemm  
Chao Liu committed
62
63
64
65
66
67
68
69
70
71
72
73 

    dim3 block_dim(BlockSize);
    dim3 grid_dim(GridSize);

    printf("%s: BlockSize %u, GridSize %u \n", __func__, BlockSize, GridSize);

    cudaEvent_t start, stop;
    float elapsedTime;

    cudaEventCreate(&start);
    cudaEventRecord(start, 0);
Chao Liu's avatar
adding implicit gemm  
Chao Liu committed
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98 
#if 0
    gridwise_implicit_gemm_convolution_nchw_kcsr
#elif 1
    gridwise_implicit_gemm_convolution_nchw_srck
#endif
    <GridSize,
     BlockSize,
     T,
     InDesc,
     WeiDesc,
     OutDesc,
     NPerBlock,
     KPerBlock,
     CPerBlock,
     HoPerBlock,
     WoPerBlock,
     KPerThread,
     CPerThread,
     HoPerThread,
     WoPerThread><<<grid_dim, block_dim>>>(InDesc{},
                                           static_cast<T*>(in_device_buf.GetDeviceBuffer()),
                                           WeiDesc{},
                                           static_cast<T*>(wei_device_buf.GetDeviceBuffer()),
                                           OutDesc{},
                                           static_cast<T*>(out_device_buf.GetDeviceBuffer()));
Chao Liu's avatar
adding implicit gemm  
Chao Liu committed
99
100
101
102
103
104
105
106
107
108
109 

    cudaEventCreate(&stop);
    cudaEventRecord(stop, 0);
    cudaEventSynchronize(stop);

    cudaEventElapsedTime(&elapsedTime, start, stop);
    printf("Elapsed time : %f ms\n", elapsedTime);

    checkCudaErrors(cudaGetLastError());
    out_device_buf.FromDevice(out.mData.data());
}