-
Notifications
You must be signed in to change notification settings - Fork 1
Expand file tree
/
Copy pathdefault.cuh
More file actions
142 lines (118 loc) · 6.07 KB
/
Copy pathdefault.cuh
File metadata and controls
142 lines (118 loc) · 6.07 KB
1
2
3
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
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
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
71
72
73
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
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
#pragma once
#include "helpers.cuh"
#include "afl.cuh"
#include "aafl.cuh"
#include "pafl.cuh"
#include "delta_pafl.cuh"
#include "delta_aafl.cuh"
#include "afl_signed_experimental.cuh"
#include "delta_signed_experimental.cuh"
template < typename T, char CWARP_SIZE, typename CCONT>
__global__ void gpu_default_decompress_kernel (CCONT cdata, container_uncompressed<T> udata)
{
unsigned long data_id, cdata_id;
set_cmp_offset <T, CWARP_SIZE> (threadIdx.x, blockIdx.x * blockDim.x, cdata.bit_length, data_id, cdata_id);
fl_decompress_func <T, CWARP_SIZE> (cdata_id, data_id, cdata, udata);
}
template < typename T, char CWARP_SIZE , typename CCONT>
__global__ void gpu_default_compress_kernel (container_uncompressed<T> udata, CCONT cdata)
{
unsigned long data_id, cdata_id;
set_cmp_offset <T, CWARP_SIZE> (threadIdx.x, blockIdx.x * blockDim.x, cdata.bit_length, data_id, cdata_id);
fl_compress_func <T, CWARP_SIZE> (data_id, cdata_id, udata, cdata);
}
template < typename T, char CWARP_SIZE , typename CCONT>
struct gpu_fl_naive_launcher_compression {
__host__ static void compress (container_uncompressed<T> udata, CCONT cdata)
{
const unsigned int block_size = CWARP_SIZE * 8; // better occupancy
const unsigned long block_number = (udata.length + block_size * CWORD_SIZE(T) - 1) / (block_size * CWORD_SIZE(T));
gpu_default_compress_kernel <T, CWARP_SIZE> <<<block_number, block_size>>> (udata, cdata);
}
};
template < typename T, char CWARP_SIZE , typename CCONT>
struct gpu_fl_naive_launcher_decompression {
__host__ static void decompress (CCONT cdata, container_uncompressed<T> udata)
{
const unsigned int block_size = CWARP_SIZE * 8; // better occupancy
const unsigned long block_number = (udata.length + block_size * CWORD_SIZE(T) - 1) / (block_size * CWORD_SIZE(T));
gpu_default_decompress_kernel <T, CWARP_SIZE> <<<block_number, block_size>>> (cdata, udata);
}
};
//AAFL specialization
template < typename T, char CWARP_SIZE>
struct gpu_fl_naive_launcher_compression <T, CWARP_SIZE, container_aafl<T>>{
__host__ static void compress (container_uncompressed<T> udata, container_aafl<T> cdata)
{
const unsigned int block_size = CWARP_SIZE * 8; // better occupancy
const unsigned long block_number = (udata.length + block_size * CWORD_SIZE(T) - 1) / (block_size * CWORD_SIZE(T));
gpu_aafl_compress_kernel <T, CWARP_SIZE> <<<block_number, block_size>>> (udata, cdata);
}
};
template < typename T, char CWARP_SIZE>
struct gpu_fl_naive_launcher_decompression <T, CWARP_SIZE, container_aafl<T>>{
__host__ static void decompress (container_aafl<T> cdata, container_uncompressed<T> udata)
{
const unsigned int block_size = CWARP_SIZE * 8; // better occupancy
const unsigned long block_number = (udata.length + block_size * CWORD_SIZE(T) - 1) / (block_size * CWORD_SIZE(T));
gpu_aafl_decompress_kernel <T, CWARP_SIZE> <<<block_number, block_size>>> (cdata, udata);
}
};
//DELTA-AAFL specialization
template < typename T, char CWARP_SIZE>
struct gpu_fl_naive_launcher_compression <T, CWARP_SIZE, container_delta_aafl<T>>{
__host__ static void compress (container_uncompressed<T> udata, container_delta_aafl<T> cdata)
{
const unsigned int block_size = CWARP_SIZE * 8; // better occupancy
const unsigned long block_number = (udata.length + block_size * CWORD_SIZE(T) - 1) / (block_size * CWORD_SIZE(T));
gpu_delta_aafl_compress_kernel <T, CWARP_SIZE> <<<block_number, block_size>>> (udata, cdata);
}
};
template < typename T, char CWARP_SIZE>
struct gpu_fl_naive_launcher_decompression <T, CWARP_SIZE, container_delta_aafl<T>>{
__host__ static void decompress (container_delta_aafl<T> cdata, container_uncompressed<T> udata)
{
const unsigned int block_size = CWARP_SIZE * 8; // better occupancy
const unsigned long block_number = (udata.length + block_size * CWORD_SIZE(T) - 1) / (block_size * CWORD_SIZE(T));
gpu_delta_aafl_decompress_kernel <T, CWARP_SIZE> <<<block_number, block_size>>> (cdata, udata);
}
};
//PAFL specialization
template < typename T, char CWARP_SIZE>
struct gpu_fl_naive_launcher_decompression <T, CWARP_SIZE, container_pafl<T>>{
__host__ static void decompress (container_pafl<T> cdata, container_uncompressed<T> udata)
{
container_fl<T> cdata_fl = { cdata.bit_length, cdata.data, cdata.length};
gpu_fl_naive_launcher_decompression<T, CWARP_SIZE, container_fl<T>>::decompress(cdata_fl, udata);
cudaErrorCheck();
unsigned int block_size = CWARP_SIZE * 8; // better occupancy
unsigned long block_number = (cdata.length + block_size * CWARP_SIZE - 1) / (block_size * CWARP_SIZE);
patch_apply_kernel <T, CWARP_SIZE> <<<block_number * CWARP_SIZE, block_size>>> (udata, cdata);
cudaErrorCheck();
}
};
//PAFL specialization
template < typename T, char CWARP_SIZE>
struct gpu_fl_naive_launcher_decompression <T, CWARP_SIZE, container_delta_pafl<T>>{
__host__ static void decompress (container_delta_pafl<T> cdata, container_uncompressed<T> udata)
{
unsigned int block_size = CWARP_SIZE * 8; // better occupancy
unsigned long block_number = (cdata.length + block_size * CWARP_SIZE - 1) / (block_size * CWARP_SIZE);
container_pafl<T> cdata_pafl = {cdata.bit_length, cdata.data, cdata.length, cdata.patch_values, cdata.patch_index, cdata.patch_count};
patch_apply_kernel <T, CWARP_SIZE> <<<block_number * CWARP_SIZE, block_size>>> (udata, cdata_pafl);
cudaErrorCheck();
gpu_default_decompress_kernel <T, CWARP_SIZE> <<<block_number, block_size>>> (cdata, udata);
cudaErrorCheck();
}
};
// Launchers
template < typename T, char CWARP_SIZE, typename X>
__host__ void compress (container_uncompressed<T> udata, X cdata)
{
gpu_fl_naive_launcher_compression<T, CWARP_SIZE, X>::compress(udata, cdata);
}
template < typename T, char CWARP_SIZE, typename X>
__host__ void decompress (X cdata, container_uncompressed<T> udata)
{
gpu_fl_naive_launcher_decompression<T, CWARP_SIZE, X>::decompress(cdata, udata);
}