refactoring

Author: Aympab

benchmark/conv1d_bench.cpp

@@ -7,11 +7,6 @@
 static constexpr auto __WG_SIZE = 1024;
 static constexpr real_t __INIT_VALUE = 7.3;
 
-auto
-sycl_alloc(size_t size, sycl::queue &q) {
-    return sycl::malloc_device<real_t>(size, q);
-}
-
 // ==========================================
 struct BenchmarkConv1dParams {
     int batch_size;

src/CMakeLists.txt

@@ -1,13 +1,11 @@
 add_subdirectory(config)
-# add_subdirectory(core)
 
 function(add_bkma_executable name)
   add_executable(${name} ${name}.cpp)
 
   target_link_libraries(${name}
   PUBLIC
   bkma::config)
-  # bkma::core)
 
   target_include_directories(${name}
     PUBLIC
@@ -27,5 +25,5 @@ function(add_bkma_executable name)
 endfunction()
 
 # Add executables
-# add_bkma_executable(advection)
+add_bkma_executable(advection)
 add_bkma_executable(conv1d)

src/advection.cpp

@@ -1,6 +1,5 @@
 #include <AdvectionParams.hpp>
 #include <AdvectionSolver.hpp>
-#include <advectors.hpp>
 #include <iostream>
 #include <sycl/sycl.hpp>
 #include <unique_ref.hpp>
@@ -15,16 +14,18 @@
 // ==========================================
 // returns duration for maxIter-1 iterations
 std::chrono::duration<double>
-advection(sycl::queue &Q, real_t* fidst_dev,
-          sref::unique_ref<IAdvectorX> &advector, const AdvectionSolver &solver) {
+advection(sycl::queue &Q, span3d_t data, const AdvectionSolver &solver,
+          BkmaOptimParams &optim_params) {
 
     auto static const maxIter = solver.params.maxIter;
 
     auto start = std::chrono::high_resolution_clock::now();
     // Time loop
     for (size_t t = 0; t < maxIter; ++t) {
-        advector(Q, fidst_dev, solver);
+        bkma_run<AdvectionSolver, BkmaImpl::AdaptiveWg>(Q, data, solver,
+                                                        optim_params);
         Q.wait();
+
     }   // end for t < T
     auto end = std::chrono::high_resolution_clock::now();
 
@@ -42,24 +43,11 @@ main(int argc, char **argv) {
     ADVParamsNonCopyable strParams;// = ADVParamsNonCopyable();
     strParams.setup(configMap);
 
-    const auto run_on_gpu = strParams.gpu;
-
-    sycl::device d;
-    if (run_on_gpu)
-        try {
-            d = sycl::device{sycl::gpu_selector_v};
-        } catch (const sycl::exception e) {
-            std::cout
-                << "GPU was requested but none is available, running kernels "
-                   "on the CPU\n"
-                << std::endl;
-            d = sycl::device{sycl::cpu_selector_v};
-            strParams.gpu = false;
-        }
-    else
-        d = sycl::device{sycl::cpu_selector_v};
-
-    sycl::queue Q{d};
+    const bool run_on_gpu = strParams.gpu;
+    auto device = pick_device(run_on_gpu);
+    strParams.gpu = device.is_gpu() ? true : false;
+
+    sycl::queue Q{device};
 
     /* Display infos on current device */
     std::cout << "Using device: "
@@ -76,24 +64,21 @@ main(int argc, char **argv) {
 
     /* Buffer for the distribution function containing the probabilities of
     having a particle at a particular speed and position, plus a fictive dim */
-    real_t* fdist = sycl::malloc_device<real_t>(n0*n1*n2, Q);
+    span3d_t data(sycl_alloc(n0*n1*n2, Q), n0, n1, n2);
     Q.wait();
-    fill_buffer(Q, fdist, params);
+    fill_buffer_adv(Q, data, params);
 
     AdvectionSolver solver(params);
-    auto advector = kernel_impl_factory(Q, strParams, solver);
+    auto optim_params = create_optim_params<ADVParams>(Q, params);
 
-    auto elapsed_seconds = advection(Q, fdist, advector, solver);
+    auto elapsed_seconds = advection(Q, data, solver, optim_params);
 
-    std::cout << "\nRESULTS_VALIDATION:" << std::endl;
-    validate_result(Q, fdist, params);
+    validate_result_adv(Q, data, params);
 
-    std::cout << "PERF_DIAGS:" << std::endl;
-    std::cout << "elapsed_time: " << elapsed_seconds.count() << " s\n";
+    auto const n_cells = n0 * n1 * n2 * (maxIter);
+    print_perf(elapsed_seconds.count(), n_cells);
 
-    auto gcells = ((n0*n1*n2*(maxIter)) / elapsed_seconds.count()) / 1e9;
-    std::cout << "upd_cells_per_sec: " << gcells << " Gcell/sec\n";
-    std::cout << "estimated_throughput: " << gcells * sizeof(real_t) * 2
-              << " GB/s" << std::endl;
+    sycl::free(data.data_handle(), Q);
+    Q.wait();
     return 0;
 }

src/config/CMakeLists.txt

@@ -11,6 +11,7 @@ target_sources(config
 target_include_directories(config
   PUBLIC
   ${CMAKE_SOURCE_DIR}/src/config
+  ${CMAKE_SOURCE_DIR}/src/tools
   ${CMAKE_SOURCE_DIR}/src)
 
 target_link_libraries(config PUBLIC mdspan)

src/conv1d.cpp

@@ -8,11 +8,6 @@
 #include <init.hpp>
 #include <validation.hpp>
 
-auto
-sycl_alloc(size_t size, sycl::queue &q) {
-    return sycl::malloc_shared<real_t>(size, q);
-}
-
 // ==========================================
 // ==========================================
 int
@@ -110,5 +105,7 @@ main(int argc, char **argv) {
     sycl::free(weight.data_handle(), Q);
     sycl::free(bias.data_handle(), Q);
     sycl::free(data.data_handle(), Q);
+    Q.wait();
+
     return 0;
 }

src/tools/init.hpp

@@ -2,7 +2,10 @@
 #include <AdvectionParams.hpp>
 #include <cmath>
 #include <sycl/sycl.hpp>
+#include <bkma.hpp>
 
+// ==========================================
+// ==========================================
 inline sycl::device
 pick_device(bool run_on_gpu) {
     sycl::device d;
@@ -21,27 +24,26 @@ pick_device(bool run_on_gpu) {
         d = sycl::device{sycl::cpu_selector_v};
 
     return d;
-}
+} //end pick_device
 
 // ==========================================
 // ==========================================
 void
-fill_buffer_adv(sycl::queue &q, real_t *fdist_dev, const ADVParams &params) {
+fill_buffer_adv(sycl::queue &q, span3d_t &data, const ADVParams &params) {
     const auto n0 = params.n0, n1 = params.n1, n2 = params.n2;
 
     sycl::range r3d(n0, n1, n2);
     q.submit([&](sycl::handler &cgh) {
          cgh.parallel_for(r3d, [=](auto i) {
-             span3d_t fdist(fdist_dev, n0, n1, n2);
              const size_t i0 = i[0];
              const size_t i1 = i[1];
              const size_t i2 = i[2];
 
              real_t x = params.minRealX + i1 * params.dx;
-             fdist(i0, i1, i2) = sycl::sin(4 * x * M_PI);
+             data(i0, i1, i2) = sycl::sin(4 * x * M_PI);
          });      // end parallel_for
      }).wait();   // end q.submit
-}
+} // end fill_buffer_adv
 
 // ==========================================
 // ==========================================
@@ -67,4 +69,33 @@ fill_buffer_conv1d(sycl::queue &q, span3d_t &data, span3d_t &warmup_data,
                    [=](unsigned itm) { bias(itm) = 1.0; });
 
     q.wait();
-}
+} // end fill_buffer_conv1d
+
+// ==========================================
+// ==========================================
+template <typename Params>
+BkmaOptimParams create_optim_params(sycl::queue &q, const Params &params) {
+    const auto n0 = params.n0;
+    const auto n1 = params.n1;
+    const auto n2 = params.n2;
+
+    WorkItemDispatch wi_dispatch;
+    wi_dispatch.set_ideal_sizes(params.pref_wg_size, n0, n1, n2);
+    auto max_elem_local_mem =
+        q.get_device().get_info<sycl::info::device::local_mem_size>() /
+        sizeof(real_t);
+    wi_dispatch.adjust_sizes_mem_limit(max_elem_local_mem, n1);
+
+    WorkGroupDispatch wg_dispatch;
+    wg_dispatch.set_num_work_groups(n0, n2, params.seq_size0, params.seq_size2,
+                                    wi_dispatch.w0_, wi_dispatch.w2_);
+
+    /* TODO : here compute the number of batchs */
+    return BkmaOptimParams{{1, n0, n0},        // BatchConfig1D dispatch_d0
+                           {1, n2, n2},        // BatchConfig1D dispatch_d2
+                           wi_dispatch.w0_,    // size_t w0
+                           wi_dispatch.w1_,    // size_t w1
+                           wi_dispatch.w2_,    // size_t w2
+                           wg_dispatch,        // WorkGroupDispatch wg_disp
+                           MemorySpace::Local};
+} //end create_optim_params

src/tools/validation.hpp

@@ -6,8 +6,10 @@
 // ==========================================
 // ==========================================
 real_t
-validate_result_adv(sycl::queue &Q, real_t *fdist_dev, const ADVParams &params,
+validate_result_adv(sycl::queue &Q, span3d_t &data, const ADVParams &params,
                     bool do_print = true) {
+    std::cout << "\nRESULTS_VALIDATION:" << std::endl;
+
     auto const dx = params.dx;
     auto const dvx = params.dvx;
     auto const dt = params.dt;
@@ -41,11 +43,9 @@ validate_result_adv(sycl::queue &Q, real_t *fdist_dev, const ADVParams &params,
 
                  cgh.parallel_for(
                      r2d, errorl1_reduc, [=](auto itm, auto &errorl1_reduc) {
-                         span3d_t fdist(fdist_dev, params.n0, params.n1,
-                                        params.n2);
                          auto i1 = itm[1];
                          auto i0 = itm[0];
-                         auto f = fdist(i0, i1, i2);
+                         auto f = data(i0, i1, i2);
 
                          real_t const x = minRealX + i1 * dx;
                          real_t const v = minRealVx + i0 * dvx;
@@ -147,3 +147,16 @@ validate_conv1d(sycl::queue &Q, span3d_t &data, size_t nw) {
     else
         std::cout << "All values data[:,i1,:] are equal." << std::endl;
 }   // end validate_conv1d
+
+// ==========================================
+// ==========================================
+void print_perf(const double elapsed_seconds, const size_t n_cells){
+
+    std::cout << "PERF_DIAGS:" << std::endl;
+    std::cout << "elapsed_time: " << elapsed_seconds << " s\n";
+
+    auto gcells = (n_cells / elapsed_seconds) / 1e9;
+    std::cout << "upd_cells_per_sec: " << gcells << " Gcell/sec\n";
+    std::cout << "estimated_throughput: " << gcells * sizeof(real_t) * 2
+              << " GB/s" << std::endl;
+}

src/types.hpp

@@ -4,6 +4,11 @@
 
 using real_t = double;
 
+[[nodiscard]] inline auto
+sycl_alloc(size_t size, sycl::queue &q) {
+    return sycl::malloc_shared<real_t>(size, q);
+}
+
 using span0d_t =
     std::experimental::mdspan<real_t, std::experimental::dextents<size_t, 0>,
                               std::experimental::layout_right>;