forked from yuhc/gpu-rodinia
-
Notifications
You must be signed in to change notification settings - Fork 0
/
Copy pathREADME
executable file
·214 lines (169 loc) · 9.31 KB
/
README
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
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167
168
169
170
171
172
173
174
175
176
177
178
179
180
181
182
183
184
185
186
187
188
189
190
191
192
193
194
195
196
197
198
199
200
201
202
203
204
205
206
207
208
209
210
211
212
213
214
Rodinia Benchmark Suite 3.1
===========================
0. Modifications
Please download the dataset from the [original package][data].
[data]: http://lava.cs.virginia.edu/Rodinia/download_links.htm
Make sure the Platform 0 is the one needed to be tested, and GPU is the
device 0 on this platform.
I. Overview
The University of Virginia Rodinia Benchmark Suite is a collection of parallel programs which targets
heterogeneous computing platforms with both multicore CPUs and GPUs.
II. Usage
1. Pakage Structure
rodinia_2.1/bin : binary executables
rodinia_2.1/common : common configuration file
rodinia_2.1/cuda : source code for the CUDA implementations
rodinia_2.1/data : input files
rodinia_2.1/openmp : source code for the OpenMP implementations
rodinia_2.1/opencl : source code for the OpenCL implementations
2. Build Rodinia
Install the CUDA/OCL drivers, SDK and toolkit on your machine.
Modify the rodinia_2.1/common/make.config file to change the settings of rodinia home directory and CUDA/OCL library paths.
To compile all the programs of the Rodinia benchmark suite, simply use the universal make file to compile all the programs, or go to each
benchmark directory and make individual programs.
3. Run Rodinia
There is a 'run' file specifying the sample command to run each program.
IV. Change Log
Dec. 12, 2015: Rodinia 3.1 is released
********************************************************
1. Bug fix
1). OpenCL version Hotspot (Thanks Shuai Che from AMD)
Delete this parameter "CL_MEM_ALLOC_HOST_PTR" for device-side buffer allocation.
2). OpenCL version Kmeans (Thanks Jeroen Ketema from Imperial College London, Tzu-Te from National Chiao Tung University, Shuai Che and Michael Boyer form AMD )
Fix data race problem for reduce kernel.
3). OpenCL version Leukocyte (Thanks Jeroen Ketema from Imperial College London)
Fix data race problem for find_ellipse kernel.
4). OpenCL version srad (Thanks Jeroen Ketema from Imperial College London)
Fix data race problem for reduce kernel
5). OpenCL version dwt2d (Thanks Tzu-Te from National Chiao Tung University)
Fix a bug for buffer size.
2. New benchmarks (Thanks Linh Nguyen from Hampden-Sydney College)
1). Hotspot3D(CUDA, OpenMP and OpenCL version)
2). Huffman (only CUDA version)
3. Performance improvement
1). Openmp version nn (Thanks Shuai Che from AMD)
2). OpenCL version nw (Thanks Shuai Che from AMD)
3). CUDA version cfd (Thanks Ke)
5. Several OpenMP benchmarks have been improved (Thanks Sergey Vinogradov and Julia Fedorova from Intel)
1). BFS
2). LUD
3). HotSpot
4). CFD
5). NW
Mar. 02, 2013: Rodinia 2.3 is released
***********************************************************************
A. General
Add -lOpenCL in the OPENCL_LIB definition in common/make.config
OPENCL_LIB = $(OPENCL_DIR)/OpenCL/common/lib -lOpenCL (gcc-4.6+ compatible)
B. OpenCL
1. Particlefilter OpenCL
a) Runtime work group size selection based on device limits
b) Several bugs of kernel fixed
c) Initialize all arrays on host side and device side
d) Fix objxy_GPU array across boundary access on device
objxy_GPU = clCreateBuffer(context, CL_MEM_READ_WRITE, 2*sizeof (int) *countOnes, NULL, &err);
and
err = clEnqueueWriteBuffer(cmd_queue, objxy_GPU, 1, 0, 2*sizeof (int) *countOnes, objxy, 0, 0, 0);
e) #define PI 3.1415926535897932 in ex_particle_OCL_naive_seq.cpp
f) put -lOpenCL just behind -L$(OPENCL_LIB) in Makefile.
g) delete an useless function tex1Dfetch() from particle_float.cl.
h) add single precision version!
2. B+Tree OpenCL
a) Replace CUDA function __syncthreads() with OpenCL barrier(CLK_LOCAL_MEM_FENCE) in kernel file
3. Heartwall OpenCL
a) Lower work item size from 512 to 256 (Better compatibility with AMD GPU)
b) Several bugs fixed on kernel codes
c) Several bugs fixed on host codes
4. BSF OpenCL
a). Replace all bool with char since bool is NOT a valid type for OpenCL arguments .
b). -lOpenCL just behind -L$(OPENCL_LIB) in Makefile. (gcc-4.6+ compatible)
c). remove NVIDIA-specific parameters and decrease thread block size for Better compatibility with AMD GPU
BFS/CLHelper.h:
//std::string options= "-cl-nv-verbose"; // doesn't work on AMD machines
resultCL = clBuildProgram(oclHandles.program, deviceListSize, oclHandles.devices, NULL, NULL,? NULL);
bfs.cpp:
#define MAX_THREADS_PER_BLOCK 256 // 512 is too big for my AMD Fusion GPU
d) Correct bad mallocs
BFS/CLHelper.h
oclHandles.devices = (cl_device_id *)malloc(deviceListSize * sizeof(cl_device_id));
d_mem = clCreateBuffer(oclHandles.context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, size, h_mem_ptr, &oclHandles.cl_status);
d_mem = clCreateBuffer(oclHandles.context, CL_MEM_WRITE_ONLY | CL_MEM_COPY_HOST_PTR, size, h_mem_ptr, &oclHandles.cl_status);
h_mem_pinned = (cl_float *)clEnqueueMapBuffer(oclHandles.queue, d_mem_pinned, CL_TRUE,? \
CL_MAP_WRITE, 0, size, 0, NULL,? \
bfs.cpp
d_graph_mask = _clMallocRW(no_of_nodes*sizeof(bool), h_graph_mask);
d_updating_graph_mask = _clMallocRW(no_of_nodes*sizeof(bool), h_updating_graph_mask);
d_graph_visited = _clMallocRW(no_of_nodes*sizeof(bool), h_graph_visited);
compare_results<int>(h_cost_ref, h_cost, no_of_nodes);
f) Add #include <cstdlib> in bfs.cpp
g) Conditional including time.h
5. CFD OpenCL
a) Comment out two useless clWaitForEvents commands in CLHelper.h. It will get 1.5X speedup on some GPUs.
b) -lOpenCL just behind -L$(OPENCL_LIB) in Makefile. (gcc-4.6+ compatible)
c) cfd/CLHelper.h
oclHandles.devices = (cl_device_id *)malloc(sizeof(cl_device_id) * deviceListSize);
6. Backprop OpenCL.
a) Opencl doesn’t support integer log2 and pow
backprop_kernel.cl 40 & 42 To:
for ( int i = 1 ; i <= HEIGHT ; i=i*2){
int power_two = i;
b) Change if( device_list ) delete device_list; to
if( device_list ) delete[] device_list;
7. gaussianElim OpenCL
a) Add codes to release device buffer at the end of ForwardSub() function (gaussianElim.cpp)
b) gaussian/gaussianElim.cpp
Add cl_cleanup(); after free(finalVec);
8. Lavamd OpenCL: In lavaMD/kernel/kernel_gpu_opencl_wrapper.c
add : #include <string.h>
9. pathfinder OpenCL
a) OpenCL.cpp: add #include <cstdlib>
b) Makefile: Changed the plase of -lOpenCL for better compatibility of gcc-4.6+.
10. streamcluster OpenCL: In CLHelper.h
oclHandles.devices = (cl_device_id *)malloc(sizeof(cl_device_id)*deviceListSize);
11. Hotspot OpenCL: In hotspot.c add clReleaseContext(context);
before main function return.
12. kmeans OpenCL: Add shutdown() in main function to release CL resource before quit.
C. CUDA
1. CFD CUDA: solve compatablity problem with CUDA 5.0.
2. Backprop CUDA: Correct include command in backprop_cuda.cu
3. BFS CUDA: Correct include command in backprop_cuda.cu
4. kmeans CUDA: Add “-lm” in link command.
5. nn CUDA: Fix makefile bugs
6. mummergpu CUDA
a) add #include <stdint.h> to
mummergpu_gold.cpp
mummergpu_main.cpp
suffix-tree.cpp
b) mummergpu.cu: correct void boardMemory function parameters types.
c) Rename getRef function to getRefGold in mummergpu_gold.cpp to avoid multiple definition
D. OpenMP
1. Kmeans OpenMP
Rename variable max_dist to min_dist in kmeans_clustering.c in kmeans_openmp/ and kmeans_serial/ folders to avoid misunderstanding.
***********************************************************************
For bug reports and fixes:
Thanks Alexey Kravets, Georgia Kouveli and Elena Stohr from CARP project. Thanks Maxim Perminov from Intel.Thanks Daniel Lustig from Princeton. Thanks John Andrew Stratton from UIUC. Thanks Mona Jalal from University of Wisconsin.
Oct. 09, 2012: Rodinia 2.2 is released
- BFS: Delete invalid flag CL_MEM_USE_HOST_PTR from _clMallocRW and _clMalloc functions in opencl verion. Thanks Alexey Kravets (CARP European research project).
- Hotspot: hotspot_kernel.cl:61 correct the index calculation as grid_cols *loadYidx + loadXidx. Correct the same problem in hotspot.cu:152. Thanks Alexey Kravets.
- Pathfinder: Added two __syncthreads in dynproc_kernel function of CUDA version to avoid data race. Thanks Ronny Krashinsky(Nvidia company) and Jiayuan Meng(Argonne National Laboratory). Alexey Kravets found and corrected the same problem in opencl version.
- SRAD: Replace CUDA function __syncthreads() in srad OpenCL kernel with OpenCL barrier(CLK_LOCAL_MEM_FENCE).
- NN: Fixed the bug of CUDA version on certain input sizes. The new version detects excess of x-dimension size limit of a CUDA block grid and executes a two-dimensional grid if needed.(Only cuda version has this problem)
- Promote B+Tree to main distribution (with output)
- Promote Myocyte to main distribution (with output)
June 27, 2012: Rodinia 2.1 is released
- Include fixes for SRAD, Heartwall, Particle Filter and Streamcluster
Nov 23, 2011: Rodinia 2.0.1 is released
- Include a CUDA version of NN comparable to the OCL version.
- Use a new version of clutils that is BSD, not GPL.
Nov 11, 2011: Rodinia 2.0 is released
- Include several applications into the main suite:
lavaMD, Gaussian Elimination, Pathfinder, k-Nearest Neighbor and Particle Filter.
Detailed application information can also be found at http://lava.cs.virginia.edu/wiki/rodinia
- Merge new OpenCL implementations into the main tarball.
Mar 01, 2010: Rodinia 1.0 is released
III. Contact
Ke Wang: [email protected]
Shuai Che: [email protected]
Kevin Skadron: [email protected]
Rodinia wiki:
http://lava.cs.virginia.edu/wiki/rodinia