/[escript]/trunk/cusplibrary/performance/spmv/benchmark.h
ViewVC logotype

Contents of /trunk/cusplibrary/performance/spmv/benchmark.h

Parent Directory Parent Directory | Revision Log Revision Log


Revision 5148 - (show annotations)
Mon Sep 15 01:25:23 2014 UTC (6 years, 5 months ago) by caltinay
File MIME type: text/plain
File size: 10323 byte(s)
Merging ripley diagonal storage + CUDA support into trunk.
Options file version has been incremented due to new options
'cuda' and 'nvccflags'.

1 #pragma once
2
3 #include <string>
4 #include <iostream>
5 #include <stdio.h>
6
7 #include "bytes_per_spmv.h"
8
9 #include "../timer.h"
10 #include <cusp/detail/device/spmv/coo_flat_k.h>
11
12 const char * BENCHMARK_OUTPUT_FILE_NAME = "benchmark_output.log";
13
14 template <typename HostMatrix, typename TestMatrix, typename TestKernel>
15 float check_spmv(HostMatrix& host_matrix, TestMatrix& test_matrix, TestKernel test_kernel)
16 {
17 typedef typename TestMatrix::index_type IndexType; // ASSUME same as HostMatrix::index_type
18 typedef typename TestMatrix::value_type ValueType; // ASSUME same as HostMatrix::value_type
19 typedef typename TestMatrix::memory_space MemorySpace;
20
21 const IndexType M = host_matrix.num_rows;
22 const IndexType N = host_matrix.num_cols;
23
24 // create host input (x) and output (y) vectors
25 cusp::array1d<ValueType,cusp::host_memory> host_x(N);
26 cusp::array1d<ValueType,cusp::host_memory> host_y(M);
27 //for(IndexType i = 0; i < N; i++) host_x[i] = (rand() % 21) - 10;
28 for(IndexType i = 0; i < N; i++) host_x[i] = (int(i % 21) - 10);
29 for(IndexType i = 0; i < M; i++) host_y[i] = 0;
30
31 // create test input (x) and output (y) vectors
32 cusp::array1d<ValueType, MemorySpace> test_x(host_x.begin(), host_x.end());
33 cusp::array1d<ValueType, MemorySpace> test_y(host_y.begin(), host_y.end());
34
35 // compute SpMV on host and device
36 cusp::multiply(host_matrix, host_x, host_y);
37 test_kernel(test_matrix, thrust::raw_pointer_cast(&test_x[0]), thrust::raw_pointer_cast(&test_y[0]));
38
39 // compare results
40 cusp::array1d<ValueType,cusp::host_memory> test_y_copy(test_y.begin(), test_y.end());
41 double error = l2_error(M, thrust::raw_pointer_cast(&test_y_copy[0]), thrust::raw_pointer_cast(&host_y[0]));
42
43 // if (error > 0.0001)
44 // {
45 // for(int i = 0; i < std::min<int>(N,256); i++)
46 // printf("host_x[%5d] = %10.8f\n", i, (float) host_x[i]);
47 //
48 // int limit = 256;
49 // for(int i = 0; i < M; i++)
50 // {
51 // if(std::abs(host_y[i] - test_y[i]) > 0.0)
52 // {
53 // printf("host_y[%5d] = %10.8f test_y[%5d] = %10.8f\n", i, (float) host_y[i], i, (float) test_y[i]);
54 //
55 // if(--limit <= 0)
56 // break;
57 // }
58 // }
59 // }
60
61 return error;
62 }
63
64
65 template <typename TestMatrix, typename TestKernel>
66 float time_spmv(TestMatrix& test_matrix, TestKernel test_spmv, double seconds = 3.0, size_t min_iterations = 100, size_t max_iterations = 500)
67 {
68 typedef typename TestMatrix::index_type IndexType; // ASSUME same as HostMatrix::index_type
69 typedef typename TestMatrix::value_type ValueType; // ASSUME same as HostMatrix::value_type
70 typedef typename TestMatrix::memory_space MemorySpace;
71
72 const IndexType M = test_matrix.num_rows;
73 const IndexType N = test_matrix.num_cols;
74
75 // create test input (x) and output (y) vectors
76 cusp::array1d<ValueType, MemorySpace> test_x(N);
77 cusp::array1d<ValueType, MemorySpace> test_y(M);
78
79 // warmup
80 timer time_one_iteration;
81 test_spmv(test_matrix, thrust::raw_pointer_cast(&test_x[0]), thrust::raw_pointer_cast(&test_y[0]));
82 cudaThreadSynchronize();
83 double estimated_time = time_one_iteration.seconds_elapsed();
84
85 // determine # of seconds dynamically
86 size_t num_iterations;
87 if (estimated_time == 0)
88 num_iterations = max_iterations;
89 else
90 num_iterations = std::min(max_iterations, std::max(min_iterations, (size_t) (seconds / estimated_time)) );
91
92 // time several SpMV iterations
93 timer t;
94 for(size_t i = 0; i < num_iterations; i++)
95 test_spmv(test_matrix, thrust::raw_pointer_cast(&test_x[0]), thrust::raw_pointer_cast(&test_y[0]));
96 cudaThreadSynchronize();
97
98 float sec_per_iteration = t.seconds_elapsed() / num_iterations;
99
100 return sec_per_iteration;
101 }
102
103
104 template <typename HostMatrix, typename TestMatrixOnHost, typename TestMatrixOnDevice, typename TestKernel>
105 void test_spmv(std::string kernel_name,
106 HostMatrix& host_matrix,
107 TestMatrixOnHost& test_matrix_on_host,
108 TestMatrixOnDevice& test_matrix_on_device,
109 TestKernel test_spmv)
110 {
111 float error = check_spmv(host_matrix, test_matrix_on_device, test_spmv);
112 float time = time_spmv( test_matrix_on_device, test_spmv);
113 float gbyte = bytes_per_spmv(test_matrix_on_host);
114
115 float GFLOPs = (time == 0) ? 0 : (2 * host_matrix.num_entries / time) / 1e9;
116 float GBYTEs = (time == 0) ? 0 : (gbyte / time) / 1e9;
117
118 printf("\t%-20s: %8.4f ms ( %5.2f GFLOP/s %5.1f GB/s) [L2 error %f]\n", kernel_name.c_str(), 1e3 * time, GFLOPs, GBYTEs, error);
119
120 //record results to file
121 FILE * fid = fopen(BENCHMARK_OUTPUT_FILE_NAME, "a");
122 fprintf(fid, "kernel=%s gflops=%f gbytes=%f msec=%f\n", kernel_name.c_str(), GFLOPs, GBYTEs, 1e3 * time);
123 fclose(fid);
124 }
125
126
127 /////////////////////////////////////////////////////
128 // These methods test specific formats and kernels //
129 /////////////////////////////////////////////////////
130
131 template <typename HostMatrix>
132 void test_coo(HostMatrix& host_matrix)
133 {
134 typedef typename HostMatrix::index_type IndexType;
135 typedef typename HostMatrix::value_type ValueType;
136
137 // convert HostMatrix to TestMatrix on host
138 cusp::coo_matrix<IndexType, ValueType, cusp::host_memory> test_matrix_on_host(host_matrix);
139
140 // transfer TestMatrix to device
141 typedef typename cusp::coo_matrix<IndexType, ValueType, cusp::device_memory> DeviceMatrix;
142 DeviceMatrix test_matrix_on_device(test_matrix_on_host);
143
144 test_spmv("coo_flat", host_matrix, test_matrix_on_host, test_matrix_on_device, cusp::detail::device::spmv_coo_flat <DeviceMatrix,ValueType>);
145 test_spmv("coo_flat_tex", host_matrix, test_matrix_on_host, test_matrix_on_device, cusp::detail::device::spmv_coo_flat_tex<DeviceMatrix,ValueType>);
146
147 // test_spmv("coo_flat_k", host_matrix, test_matrix_on_host, test_matrix_on_device, cusp::detail::device::spmv_coo_flat_k <DeviceMatrix,ValueType>);
148 // test_spmv("coo_flat_k_tex", host_matrix, test_matrix_on_host, test_matrix_on_device, cusp::detail::device::spmv_coo_flat_k_tex<DeviceMatrix,ValueType>);
149 }
150
151 template <typename HostMatrix>
152 void test_csr(HostMatrix& host_matrix)
153 {
154 typedef typename HostMatrix::index_type IndexType;
155 typedef typename HostMatrix::value_type ValueType;
156
157 // convert HostMatrix to TestMatrix on host
158 cusp::csr_matrix<IndexType, ValueType, cusp::host_memory> test_matrix_on_host(host_matrix);
159
160 // transfer csr_matrix to device
161 typedef typename cusp::csr_matrix<IndexType, ValueType, cusp::device_memory> DeviceMatrix;
162 DeviceMatrix test_matrix_on_device(test_matrix_on_host);
163
164 test_spmv("csr_scalar", host_matrix, test_matrix_on_host, test_matrix_on_device, cusp::detail::device::spmv_csr_scalar <DeviceMatrix,ValueType>);
165 test_spmv("csr_scalar_tex", host_matrix, test_matrix_on_host, test_matrix_on_device, cusp::detail::device::spmv_csr_scalar_tex<DeviceMatrix,ValueType>);
166 test_spmv("csr_vector", host_matrix, test_matrix_on_host, test_matrix_on_device, cusp::detail::device::spmv_csr_vector <DeviceMatrix,ValueType>);
167 test_spmv("csr_vector_tex", host_matrix, test_matrix_on_host, test_matrix_on_device, cusp::detail::device::spmv_csr_vector_tex<DeviceMatrix,ValueType>);
168 }
169
170 template <typename HostMatrix>
171 void test_dia(HostMatrix& host_matrix)
172 {
173 typedef typename HostMatrix::index_type IndexType;
174 typedef typename HostMatrix::value_type ValueType;
175
176 // convert HostMatrix to TestMatrix on host
177 cusp::dia_matrix<IndexType, ValueType, cusp::host_memory> test_matrix_on_host;
178
179 try
180 {
181 test_matrix_on_host = host_matrix;
182 }
183 catch (cusp::format_conversion_exception)
184 {
185 std::cout << "\tRefusing to convert to DIA format" << std::endl;
186 return;
187 }
188
189 // transfer TestMatrix to device
190 typedef typename cusp::dia_matrix<IndexType, ValueType, cusp::device_memory> DeviceMatrix;
191 DeviceMatrix test_matrix_on_device(test_matrix_on_host);
192
193 test_spmv("dia", host_matrix, test_matrix_on_host, test_matrix_on_device, cusp::detail::device::spmv_dia <DeviceMatrix,ValueType>);
194 test_spmv("dia_tex", host_matrix, test_matrix_on_host, test_matrix_on_device, cusp::detail::device::spmv_dia_tex<DeviceMatrix,ValueType>);
195 }
196
197 template <typename HostMatrix>
198 void test_ell(HostMatrix& host_matrix)
199 {
200 typedef typename HostMatrix::index_type IndexType;
201 typedef typename HostMatrix::value_type ValueType;
202
203 // convert HostMatrix to TestMatrix on host
204 cusp::ell_matrix<IndexType, ValueType, cusp::host_memory> test_matrix_on_host;
205
206 try
207 {
208 test_matrix_on_host = host_matrix;
209 }
210 catch (cusp::format_conversion_exception)
211 {
212 std::cout << "\tRefusing to convert to ELL format" << std::endl;
213 return;
214 }
215
216 // transfer TestMatrix to device
217 typedef typename cusp::ell_matrix<IndexType, ValueType, cusp::device_memory> DeviceMatrix;
218 DeviceMatrix test_matrix_on_device(test_matrix_on_host);
219
220 test_spmv("ell", host_matrix, test_matrix_on_host, test_matrix_on_device, cusp::detail::device::spmv_ell <DeviceMatrix,ValueType>);
221 test_spmv("ell_tex", host_matrix, test_matrix_on_host, test_matrix_on_device, cusp::detail::device::spmv_ell_tex<DeviceMatrix,ValueType>);
222 }
223
224 template <typename HostMatrix>
225 void test_hyb(HostMatrix& host_matrix)
226 {
227 typedef typename HostMatrix::index_type IndexType;
228 typedef typename HostMatrix::value_type ValueType;
229
230 // convert HostMatrix to TestMatrix on host
231 cusp::hyb_matrix<IndexType, ValueType, cusp::host_memory> test_matrix_on_host(host_matrix);
232
233 // transfer TestMatrix to device
234 typedef typename cusp::hyb_matrix<IndexType, ValueType, cusp::device_memory> DeviceMatrix;
235 DeviceMatrix test_matrix_on_device(test_matrix_on_host);
236
237 test_spmv("hyb", host_matrix, test_matrix_on_host, test_matrix_on_device, cusp::detail::device::spmv_hyb <DeviceMatrix,ValueType>);
238 test_spmv("hyb_tex", host_matrix, test_matrix_on_host, test_matrix_on_device, cusp::detail::device::spmv_hyb_tex<DeviceMatrix,ValueType>);
239 }
240

  ViewVC Help
Powered by ViewVC 1.1.26