main.cpp

#include <CL/sycl.hpp>
#include <array>
#include <iostream>
#include <iomanip>
#include <cmath>

using namespace cl::sycl;

constexpr float KILL_RATE { 0.062f };
constexpr float FEED_RATE { 0.03f };
constexpr float DT { 1.0f };

constexpr float DIFFUSION_RATE_U { 0.1f };
constexpr float DIFFUSION_RATE_V { 0.05f };

void submit( queue & q, 
    float const * iu, float const * iv,
    float * ou, float * ov,
    float * iud, float * ivd,
    float * oud, float * ovd,
    std::size_t nb_rows, std::size_t nb_cols ) {
 
    // Submit command group for execution
    q.memcpy(iud,iu,(nb_rows+2)*(nb_cols+2)*sizeof(float)).wait();
    q.memcpy(ivd,iv,(nb_rows+2)*(nb_cols+2)*sizeof(float)).wait();
    q.submit([&](handler& h) {

        // 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 = iud[(row+1)*(nb_cols+2)+col+1];
			float v = ivd[(row+1)*(nb_cols+2)+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 += (iud[(row+k)*(nb_cols+2)+col+l] - u);
					full_v += (ivd[(row+k)*(nb_cols+2)+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;

			oud[(row+1)*(nb_cols+2)+col+1] = u + du*DT;
			ovd[(row+1)*(nb_cols+2)+col+1] = v + dv*DT;
        });

    }).wait();
    q.memcpy(ou,oud,(nb_rows+2)*(nb_cols+2)*sizeof(float)).wait();
    q.memcpy(ov,ovd,(nb_rows+2)*(nb_cols+2)*sizeof(float)).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 local arrays
        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 device arrays
        float * iud  = malloc_device<float>(size, q);
        float * ivd  = malloc_device<float>(size, q);
        float * oud  = malloc_device<float>(size, q);
        float * ovd  = malloc_device<float>(size, q);
        
        // iterations
        for ( std::size_t image = 0 ; image < nb_images ; ++image ) {
            for ( std::size_t iter = 0 ; iter < nb_iterations ; iter += 2 ) {
                submit( q, u1.data(), v1.data(), u2.data(), v2.data(), iud, ivd, oud, ovd, nb_rows, nb_cols );
                submit( q, u2.data(), v2.data(), u1.data(), v1.data(), iud, ivd, oud, ovd, 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 };
        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 << u1[i*padded_nb_cols+j] << " ";
            }
            std::cout << "\n";
        }
        std::cout << std::endl;
        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 << v1[i*padded_nb_cols+j] << " ";
            }
            std::cout << "\n";
        }
        std::cout << std::endl;

        // Release device arrays
        sycl::free(iud,q);
        sycl::free(ivd,q);
        sycl::free(oud,q);
        sycl::free(ovd,q);

    }
    catch (sycl::exception & e) {
      std::cout << e.what() << std::endl;
      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;
    }


    return 0;
}