34
34
//
35
35
// ******************************************************************************
36
36
#include < CL/sycl.hpp>
37
- #include < dpc_common.hpp>
38
37
#include < fstream>
39
38
#include < iomanip>
40
39
#include < iostream>
40
+ #include " dpc_common.hpp"
41
41
42
42
using namespace sycl ;
43
+ using namespace std ;
43
44
44
45
constexpr float dt = 0 .002f ;
45
46
constexpr float dx = 0 .01f ;
@@ -49,12 +50,12 @@ constexpr float temp = 100.0f; // Initial temperature.
49
50
// ************************************
50
51
// Function description: display input parameters used for this sample.
51
52
// ************************************
52
- void Usage (std:: string programName) {
53
- std:: cout << " Incorrect parameters \n " ;
54
- std:: cout << " Usage: " ;
55
- std:: cout << programName << " <n> <i>\n\n " ;
56
- std:: cout << " n : Number of points to simulate \n " ;
57
- std:: cout << " i : Number of timesteps \n " ;
53
+ void Usage (string programName) {
54
+ cout << " Incorrect parameters \n " ;
55
+ cout << " Usage: " ;
56
+ cout << programName << " <n> <i>\n\n " ;
57
+ cout << " n : Number of points to simulate \n " ;
58
+ cout << " i : Number of timesteps \n " ;
58
59
}
59
60
60
61
// ************************************
@@ -75,43 +76,40 @@ float* ComputeHeatDeviceParallel(float* arr, float* arr_next, float C,
75
76
try {
76
77
// Define the device queue
77
78
queue q = default_selector{};
78
- std:: cout << " Kernel runs on "
79
- << q. get_device (). get_info <info::device::name>() << " \n " ;
79
+ cout << " Kernel runs on " << q. get_device (). get_info <info::device::name>()
80
+ << " \n " ;
80
81
81
82
// Set boundary condition at one end.
82
83
arr[0 ] = arr_next[0 ] = temp;
83
84
84
85
float * current_data_ptr = arr;
85
86
float * next_data_ptr = arr_next;
86
- // current_data_ptr = arr;
87
- // next_data_ptr = arr_next;
88
87
89
88
// Buffer scope
90
89
{
91
- buffer< float , 1 > arr_buf (current_data_ptr, range< 1 >{ num_p + 2 } );
92
- buffer< float , 1 > arr_next_buf (next_data_ptr, range< 1 >{ num_p + 2 } );
90
+ buffer temperature_buf (current_data_ptr, range ( num_p + 2 ) );
91
+ buffer temperature_next_buf (next_data_ptr, range ( num_p + 2 ) );
93
92
94
93
// Iterate over timesteps
95
94
for (i = 1 ; i <= num_iter; i++) {
96
95
if (i % 2 != 0 ) {
97
- q.submit ([&](handler & h) {
96
+ q.submit ([&](auto & h) {
98
97
// The size of memory amount that will be given to the buffer.
99
98
range<1 > num_items{num_p + 2 };
100
99
101
- auto arr_acc = arr_buf.get_access <access::mode::read_write>(h);
102
- auto arr_next_acc =
103
- arr_next_buf.get_access <access::mode::read_write>(h);
100
+ accessor temperature (temperature_buf, h);
101
+ accessor temperature_next (temperature_next_buf, h);
104
102
105
103
h.parallel_for (num_items, [=](id<1 > k) {
106
104
size_t gid = k.get (0 );
107
105
108
106
if (gid == 0 ) {
109
107
} else if (gid == num_p + 1 ) {
110
- arr_next_acc [k] = arr_acc [k - 1 ];
108
+ temperature_next [k] = temperature [k - 1 ];
111
109
} else {
112
- arr_next_acc [k] =
113
- C * (arr_acc [k + 1 ] - 2 * arr_acc [k] + arr_acc [k - 1 ]) +
114
- arr_acc [k];
110
+ temperature_next [k] =
111
+ C * (temperature [k + 1 ] - 2 * temperature [k] + temperature [k - 1 ]) +
112
+ temperature [k];
115
113
}
116
114
}); // end parallel for loop in kernel1
117
115
}); // end device queue
@@ -121,20 +119,19 @@ float* ComputeHeatDeviceParallel(float* arr, float* arr_next, float C,
121
119
// The size of memory amount that will be given to the buffer.
122
120
range<1 > num_items{num_p + 2 };
123
121
124
- auto arr_acc = arr_buf.get_access <access::mode::read_write>(h);
125
- auto arr_next_acc =
126
- arr_next_buf.get_access <access::mode::read_write>(h);
122
+ accessor temperature (temperature_buf, h);
123
+ accessor temperature_next (temperature_next_buf, h);
127
124
128
125
h.parallel_for (num_items, [=](id<1 > k) {
129
126
size_t gid = k.get (0 );
130
127
131
128
if (gid == 0 ) {
132
129
} else if (gid == num_p + 1 ) {
133
- arr_acc [k] = arr_next_acc [k - 1 ];
130
+ temperature [k] = temperature_next [k - 1 ];
134
131
} else {
135
- arr_acc [k] = C * (arr_next_acc [k + 1 ] - 2 * arr_next_acc [k] +
136
- arr_next_acc [k - 1 ]) +
137
- arr_next_acc [k];
132
+ temperature [k] = C * (temperature_next [k + 1 ] - 2 * temperature_next [k] +
133
+ temperature_next [k - 1 ]) +
134
+ temperature_next [k];
138
135
}
139
136
}); // end parallel for loop in kernel2
140
137
}); // end device queue
@@ -144,8 +141,8 @@ float* ComputeHeatDeviceParallel(float* arr, float* arr_next, float C,
144
141
145
142
q.wait_and_throw ();
146
143
147
- } catch (cl:: sycl::exception e) {
148
- std:: cout << " SYCL exception caught: " << e.what () << " \n " ;
144
+ } catch (sycl::exception e) {
145
+ cout << " SYCL exception caught: " << e.what () << " \n " ;
149
146
}
150
147
151
148
if (i % 2 != 0 )
@@ -197,7 +194,7 @@ bool CompareResults(float* device_results, float* host_results,
197
194
double norm2 = 0 ;
198
195
bool err = false ;
199
196
200
- std:: ofstream err_file;
197
+ ofstream err_file;
201
198
err_file.open (" error_diff.txt" );
202
199
203
200
err_file << " \t idx\t heat[i]\t\t heat_CPU[i] \n " ;
@@ -225,16 +222,16 @@ int main(int argc, char* argv[]) {
225
222
226
223
// Read input parameters
227
224
try {
228
- n_point = std:: stoi (argv[1 ]);
229
- n_iteration = std:: stoi (argv[2 ]);
225
+ n_point = stoi (argv[1 ]);
226
+ n_iteration = stoi (argv[2 ]);
230
227
231
228
} catch (...) {
232
229
Usage (argv[0 ]);
233
230
return (-1 );
234
231
}
235
232
236
- std:: cout << " Number of points: " << n_point << " \n " ;
237
- std:: cout << " Number of iterations: " << n_iteration << " \n " ;
233
+ cout << " Number of points: " << n_point << " \n " ;
234
+ cout << " Number of iterations: " << n_iteration << " \n " ;
238
235
239
236
// Array heat and heat_next arrays store temperatures of the current and next
240
237
// iteration of n_point (calculated in kernel)
@@ -260,7 +257,7 @@ int main(int argc, char* argv[]) {
260
257
ComputeHeatDeviceParallel (heat, heat_next, C, n_point, n_iteration, temp);
261
258
262
259
// Display time used by device
263
- std:: cout << " Kernel time: " << t_par.Elapsed () << " sec\n " ;
260
+ cout << " Elapsed time: " << t_par.Elapsed () << " sec\n " ;
264
261
265
262
// Compute heat in CPU in (for comparision)
266
263
float * final_CPU = NULL ;
@@ -273,10 +270,10 @@ int main(int argc, char* argv[]) {
273
270
bool err = CompareResults (final_device, final_CPU, n_point, C);
274
271
275
272
if (err == true )
276
- std:: cout << " Please check the error_diff.txt file ...\n " ;
273
+ cout << " Please check the error_diff.txt file ...\n " ;
277
274
else
278
- std:: cout << " PASSED! There is no difference between the results computed "
279
- " in host and in kernel.\n " ;
275
+ cout << " PASSED! There is no difference between the results computed "
276
+ " in host and in kernel.\n " ;
280
277
281
278
// Cleanup
282
279
delete[] heat;
0 commit comments