Newer
Older
#include <CL/sycl.hpp>
#include <array>
#include <iostream>
constexpr float KILL_RATE { 0.062f };
constexpr float FEED_RATE { 0.03f };
constexpr float DT { 1.0f };
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
constexpr float DIFFUSION_RATE_U { 0.1f };
constexpr float DIFFUSION_RATE_V { 0.05f };
//constexpr std::size_t nb_rows { 100ul }; //{ 1080ul };
//constexpr std::size_t nb_cols { 200ul }; //{ 1920ul };
//
//constexpr std::size_t nb_images { 5ul };
//constexpr std::size_t nb_iterations { 1000ul }; // 680000
void submit( queue & q,
buffer<float,2> & iub, buffer<float,2> & ivb,
buffer<float,2> & oub, buffer<float,2> & ovb,
std::size_t nb_rows, std::size_t nb_cols ) {
// Submit command group for execution
q.submit([&](handler& h) {
// Create accessors for input and output buffers
accessor iua(iub, h, read_only);
accessor iva(ivb, h, read_only);
accessor oua(oub, h, write_only);
accessor ova(ovb, h, write_only);
// Define the kernel
h.parallel_for(range<2>{nb_rows,nb_cols}, [=](item<2> it) {
id<2> xy = it.get_id();
std::size_t row = xy[0] ;
std::size_t col = xy[1] ;
float u = iua[id<2>{row+1,col+1}];
float v = iva[id<2>{row+1,col+1}];
float uvv = u*v*v;
float full_u = 0.0f;
float full_v = 0.0f;
for(long k = 0l; k < 3l; ++k){
for(long l = 0l; l < 3l; ++l){
full_u += (iua[id<2>{row+k,col+l}] - u);
full_v += (iva[id<2>{row+k,col+l}] - v);
}
}
float du = DIFFUSION_RATE_U*full_u - uvv + FEED_RATE*(1.0f - u);
float dv = DIFFUSION_RATE_V*full_v + uvv - (FEED_RATE + KILL_RATE)*v;
oua[id<2>{row+1,col+1}] = u + du*DT;
ova[id<2>{row+1,col+1}] = v + dv*DT;
});
});
// Wait for the command group to finish
q.wait();
}
int main( int argc, char * argv[] ) {
// runtime parameters
assert(argc=5) ;
std::size_t nb_rows {std::stoul(argv[1])} ;
std::size_t nb_cols {std::stoul(argv[2])} ;
std::size_t nb_images {std::stoul(argv[3])} ;
std::size_t nb_iterations {std::stoul(argv[4])} ;
assert(nb_iterations % 2 == 0); // nb_iterations must be even
try {
// Loop through available platforms and devices
for (auto const& this_platform : platform::get_platforms() ) {
std::cout << "Found platform: "
<< this_platform.get_info<info::platform::name>() << std::endl;
for (auto const& this_device : this_platform.get_devices() ) {
std::cout << " Device: "
<< this_device.get_info<info::device::name>() << std::endl;
}
// Create SYCL queue
queue q;
// Running platform and device
std::cout << "Running on platform: "
<< q.get_device().get_platform().get_info<info::platform::name>() << std::endl;
std::cout << " Device: "
<< q.get_device().get_info<info::device::name>() << std::endl;
std::cout << std::endl;
// Initialize input array
const std::size_t padded_nb_rows { nb_rows+2 };
const std::size_t padded_nb_cols { nb_cols+2 };
const std::size_t size { padded_nb_rows*padded_nb_cols };
std::vector<float> u1(size);
std::vector<float> v1(size);
std::vector<float> u2(size);
std::vector<float> v2(size);
for (int i = 0; i < padded_nb_rows; i++) {
for (int j = 0; j < padded_nb_cols; j++) {
u1[i*padded_nb_cols+j] = 1.f;
v1[i*padded_nb_cols+j] = 0.f;
u2[i*padded_nb_cols+j] = 1.f;
v2[i*padded_nb_cols+j] = 0.f;
}
const std::size_t v_row_begin { (7ul*padded_nb_rows+8ul)/16ul };
const std::size_t v_row_end { (9ul*padded_nb_rows+8ul)/16ul };
const std::size_t v_col_begin { (7ul*padded_nb_cols+8ul)/16ul };
const std::size_t v_col_end { (9ul*padded_nb_cols+8ul)/16ul };
std::cout << "v_row_begin: " << v_row_begin << std::endl;
std::cout << "v_row_end: " << v_row_end << std::endl;
std::cout << "v_col_begin: " << v_col_begin << std::endl;
std::cout << "v_col_end: " << v_col_end << std::endl;
std::cout << std::endl;
for (int i = v_row_begin; i < v_row_end; i++) {
for (int j = v_col_begin; j < v_col_end; j++) {
u1[i*padded_nb_cols+j] = 0.f;
v1[i*padded_nb_cols+j] = 1.f;
}
// Create buffers
buffer<float,2> u1b { u1.data(),range<2>{padded_nb_rows,padded_nb_cols} };
buffer<float,2> v1b { v1.data(),range<2>{padded_nb_rows,padded_nb_cols} };
buffer<float,2> u2b { u2.data(),range<2>{padded_nb_rows,padded_nb_cols} };
buffer<float,2> v2b { v2.data(),range<2>{padded_nb_rows,padded_nb_cols} };
// iterations
for ( std::size_t image = 0 ; image < nb_images ; ++image ) {
for ( std::size_t iter = 0 ; iter < nb_iterations ; iter += 2 ) {
submit( q, u1b, v1b, u2b, v2b, nb_rows, nb_cols );
submit( q, u2b, v2b, u1b, v1b, nb_rows, nb_cols );
}
// Print some result
const std::size_t row_center { padded_nb_rows/2ul };
const std::size_t col_center { padded_nb_cols/2ul };
host_accessor u1ha{u1b, read_only};
std::cout<<std::fixed<<std::setprecision(2) ;
for (std::size_t i = (row_center-5ul) ; i < (row_center+5ul); i++) {
for (std::size_t j = (col_center-5ul); j < (col_center+5ul); j++) {
std::cout << u1ha[id<2>{i,j}] << " ";
}
std::cout << "\n";
std::cout << std::endl;
host_accessor v1ha{v1b, read_only};
std::cout<<std::fixed<<std::setprecision(2) ;
for (std::size_t i = (row_center-5ul) ; i < (row_center+5ul); i++) {
for (std::size_t j = (col_center-5ul); j < (col_center+5ul); j++) {
std::cout << v1ha[id<2>{i,j}] << " ";
}
std::cout << "\n";
std::cout << std::endl;
}
catch (sycl::exception & e) {
std::cout << e.what() << std::endl;
CHAMONT David
committed
std::cout << e.code().message() << std::endl;
}
catch (std::exception & e) {
std::cout << e.what() << std::endl;
catch (const char * e) {
std::cout << e << std::endl;
}