Commit 6ec4cc03 authored by Alejandro Homs Puron's avatar Alejandro Homs Puron Committed by Generic Bliss account for Control Software
Browse files

[WIP][SMX] Add average mode

parent a12d405c
Pipeline #77183 failed with stages
in 7 seconds
// Copyright (C) 2020 Alejandro Homs Puron, ESRF.
// Use, modification and distribution is subject to the Boost Software
// License, Version 1.0. (See accompanying file LICENSE_1_0.txt or copy at
// http://www.boost.org/LICENSE_1_0.txt)
#pragma once
#include <lima/logging.hpp>
#include <lima/processing/fai/opencl.hpp>
#include <lima/processing/fai/enums.hpp>
#include <lima/processing/fai/typedefs.hpp>
namespace lima
{
namespace processing::fai
{
using ave_ret_t = std::size_t;
inline auto make_average(bcl::context& context, //
std::size_t width, std::size_t height, //
std::size_t start_frame, // First frame to average
std::size_t nb_frames, // Number of frames to average
std::filesystem::path const& cl_source_path, // Path to the CL sources
bcl::command_queue& queue) //
{
const std::size_t nb_pixels = width * height;
// Preconditions
if (nb_frames == 0)
LIMA_THROW_EXCEPTION(lima::invalid_argument("Invalid null average nb_frames"));
auto build_program = [&context, &cl_source_path](auto&& file, const std::string& options = std::string()) {
return build_with_source_file(context, cl_source_path, file, options);
};
auto program = build_program("average.cl");
bcl::vector<unsigned long> acc_d(nb_pixels, 0, queue); //!< Accumulation buffer
bcl::kernel kernel(program, "average");
std::uint32_t curr_acc_frames = 0;
return [=, //
acc = std::move(acc_d),
to_skip = start_frame](bcl::command_queue& queue, //
vector<std::uint16_t> const& raw, // Input
vector<float>& ave, // Output average
bool force_ave) mutable -> ave_ret_t {
// Skip the first & last frames
if (to_skip > 0) {
--to_skip;
return 0;
} else if (curr_acc_frames == nb_frames)
return 0;
kernel.set_args(raw, acc, nb_pixels);
queue.enqueue_1d_range_kernel(kernel, 0, nb_pixels, 0 /*max = 1024 for Tesla*/);
++curr_acc_frames;
if ((curr_acc_frames < nb_frames) && !force_ave)
return 0;
BOOST_COMPUTE_CLOSURE(float, calc_ave, (unsigned long acc), (curr_acc_frames),
{ return acc / curr_acc_frames; });
bcl::transform(acc.begin(), acc.begin() + nb_pixels, ave.begin(), calc_ave, queue);
return curr_acc_frames;
};
};
using average = std::function<ave_ret_t(bcl::command_queue&, //
vector<std::uint16_t> const&, // Input
vector<float>&, // Output average
bool // Force output average
)>;
} // namespace processing::fai
} // namespace lima
// Copyright (C) 2020 Samuel Debionne, ESRF.
// Use, modification and distribution is subject to the Boost Software
// License, Version 1.0. (See accompanying file LICENSE_1_0.txt or copy at
// http://www.boost.org/LICENSE_1_0.txt)
#pragma once
#include <filesystem>
#include <boost/compute.hpp>
#include <lima/exceptions.hpp>
namespace lima
{
namespace processing::fai
{
namespace bcl = boost::compute;
inline auto build_with_source_file(bcl::context& context, const std::filesystem::path& cl_source_path,
const std::string& file, const std::string& options = std::string())
{
std::string final_options = options;
auto file_path = file;
if (!cl_source_path.string().empty()) {
std::string sep = final_options.empty() ? "" : " ";
final_options += sep + "-I" + cl_source_path.string();
file_path = cl_source_path / file;
}
bcl::program res = bcl::program::create_with_source_file(file_path, context);
try {
// attempt to compile to program
res.build(final_options);
} catch (boost::compute::opencl_error& e) {
// program failed to compile, print out the build log
std::cerr << res.build_log();
LIMA_THROW_EXCEPTION(std::runtime_error(res.build_log()));
}
return res;
}
} // namespace processing::fai
} // namespace lima
......@@ -10,13 +10,10 @@
#include <functional>
#include <memory>
#include <filesystem>
#include <boost/compute.hpp>
#include <lima/exceptions.hpp>
#include <lima/logging.hpp>
#include <lima/processing/fai/opencl.hpp>
#include <lima/processing/fai/enums.hpp>
#include <lima/processing/fai/typedefs.hpp>
......@@ -24,30 +21,7 @@ namespace lima
{
namespace processing::fai
{
namespace bcl = boost::compute;
inline auto build_with_source_file(bcl::context& context, const std::filesystem::path& cl_source_path,
const std::string& file, const std::string& options = std::string())
{
std::string final_options = options;
auto file_path = file;
if (!cl_source_path.string().empty()) {
std::string sep = final_options.empty() ? "" : " ";
final_options += sep + "-I" + cl_source_path.string();
file_path = cl_source_path / file;
}
bcl::program res = bcl::program::create_with_source_file(file_path, context);
try {
// attempt to compile to program
res.build(final_options);
} catch (boost::compute::opencl_error& e) {
// program failed to compile, print out the build log
std::cerr << res.build_log();
LIMA_THROW_EXCEPTION(std::runtime_error(res.build_log()));
}
return res;
}
using pf_ret_t = std::tuple<std::size_t, std::size_t>;
inline auto make_peak_finder(bcl::context& context, //
std::size_t width, std::size_t height, //
......@@ -81,24 +55,23 @@ namespace processing::fai
// Preconditions
auto build_program = [&cl_source_path](auto&& context, auto&& file,
const std::string& options = std::string()) {
auto build_program = [&context, &cl_source_path](auto&& file, const std::string& options = std::string()) {
return build_with_source_file(context, cl_source_path, file, options);
};
// Stage 0 - Initialization
auto stage0_program = build_program(context, "init.cl");
auto stage0_program = build_program("init.cl");
bcl::kernel stage0(stage0_program, "init");
// Stage 1 - Pedestal gain correction
auto stage1_program = build_program(context, "pedestal_gain_correction.cl");
auto stage1_program = build_program("pedestal_gain_correction.cl");
bcl::vector<float> gain_d(gain, queue); //!< Gain device buffer
bcl::vector<float> pedestal_d(pedestal, queue); //!< Pedestal device buffer
bcl::vector<float> corrected_d(nb_pixels, context); //!< Gain / pedestal corrected data
bcl::kernel stage1(stage1_program, "pedestal_gain_correction");
// Stage 2 - Preprocessing
auto stage2_program = build_program(context, "preprocess.cl");
auto stage2_program = build_program("preprocess.cl");
bcl::vector<float> variance_d(variance, queue);
bcl::vector<float> dark_d(dark, queue);
bcl::vector<float> dark_variance_d(dark_variance, queue);
......@@ -122,7 +95,7 @@ namespace processing::fai
bcl::kernel stage2(stage2_program, "corrections4");
// Stage 3 - Sigma clipping
auto stage3_program = build_program(context, "csr_sigma_clip.cl", "-D WORKGROUP_SIZE=32");
auto stage3_program = build_program("csr_sigma_clip.cl", "-D WORKGROUP_SIZE=32");
bcl::vector<float> csr_data_d(csr_data, queue);
bcl::vector<int> csr_indices_d(csr_indices, queue);
bcl::vector<int> csr_indptr_d(csr_indptr, queue);
......@@ -134,7 +107,7 @@ namespace processing::fai
// Stage 4 - Peak finding
constexpr int peak_finder_wg_size = 1024;
std::string wg_option = "-D WORKGROUP_SIZE=" + std::to_string(peak_finder_wg_size);
auto stage4_program = build_program(context, "peak_finder.cl", wg_option.c_str());
auto stage4_program = build_program("peak_finder.cl", wg_option.c_str());
bcl::vector<float> radius2d_d(radius2d, queue);
bcl::vector<float> radius1d_d(radius1d, queue);
......@@ -265,17 +238,16 @@ namespace processing::fai
};
//using peak_finder = decltype(make_peak_finder(std::declval<bcl::context>(), std::declval<bcl::command_queue>()));
using ret_t = std::tuple<std::size_t, std::size_t>;
using peak_finder = std::function<ret_t(bcl::command_queue&, //
vector<std::uint16_t> const&, // Input
vector<float>& corrected_dbg, // Corrected
vector<bcl::float4_>& preprocessed_dbg, // Preprocessed
vector<bcl::float8_>& cliped_dbg, // Summed (after sigma clip)
vector<bcl::float4_>& found_dbg, // Finded (after peak find)
vector<int>&, vector<float>&, // Outputs Peaks
vector<float>&, vector<float>&, // Background
vector<float>&, vector<float>&, bool // Accumulation
)>;
using peak_finder = std::function<pf_ret_t(bcl::command_queue&, //
vector<std::uint16_t> const&, // Input
vector<float>& corrected_dbg, // Corrected
vector<bcl::float4_>& preprocessed_dbg, // Preprocessed
vector<bcl::float8_>& cliped_dbg, // Summed (after sigma clip)
vector<bcl::float4_>& found_dbg, // Finded (after peak find)
vector<int>&, vector<float>&, // Outputs Peaks
vector<float>&, vector<float>&, // Background
vector<float>&, vector<float>&, bool // Accumulation
)>;
} // namespace processing::fai
} // namespace lima
__kernel void average(__global const unsigned short* in, __global unsigned long* acc, unsigned int length)
{
size_t i = get_global_id(0);
// bound check
if (i < length)
acc[i] += in[i];
}
\ No newline at end of file
// Copyright (C) 2020 Alejandro Homs Puron, ESRF.
// Use, modification and distribution is subject to the Boost Software
// License, Version 1.0. (See accompanying file LICENSE_1_0.txt or copy at
// http://www.boost.org/LICENSE_1_0.txt)
#pragma once
namespace lima
{
namespace processing::pipelines
{
namespace smx
{
/// Frame (reconstruction) assembly mode
enum class proc_mode_enum : int
{
fai,
average
};
} // namespace smx
} // namespace processing::pipelines
} // namespace lima
blissadm@lid29p9jfrau1.83821:1648048898
\ No newline at end of file
// Copyright (C) 2018 Alejandro Homs, ESRF.
// Use, modification and distribution is subject to the Boost Software
// License, Version 1.0. (See accompanying file LICENSE_1_0.txt or copy at
// http://www.boost.org/LICENSE_1_0.txt)
#pragma once
#include <memory>
#include <tbb/flow_graph.h>
#include <lima/exceptions.hpp>
#include <lima/logging.hpp>
#include <lima/core/enums.hpp>
#include <lima/core/image/view.hpp>
#include <lima/processing/fai/average.hpp>
namespace lima
{
namespace processing::pipelines
{
namespace smx
{
namespace bcl = boost::compute;
struct average_node : public tbb::flow::multifunction_node<lima::frame, std::tuple<lima::frame>>
{
using parent_t = tbb::flow::multifunction_node<lima::frame, std::tuple<lima::frame>>;
struct average_body
{
average_body(std::size_t width, std::size_t height, std::size_t start_frame, std::size_t nb_frames,
std::filesystem::path const& cl_source_path, bcl::context& context,
bcl::command_queue& queue) :
context(context),
queue(queue),
nb_pixels(width * height),
raw_d(nb_pixels, context), // Input
ave_d(nb_pixels, context) // Output Average
{
LIMA_LOG(trace, proc) << "nb_pixels:" << nb_pixels;
// Create the average calculator
ave_calc = fai::make_average(context, //
width, height, //
start_frame, nb_frames, cl_source_path, queue);
LIMA_LOG(trace, proc) << "Peak finder created";
}
void operator()(lima::frame const& in, typename parent_t::output_ports_type& ports) const
{
assert(nb_pixels == in.size());
// Write input to the device (non-blocking)
auto v_in = lima::const_view<gray16c_view_t>(in);
bcl::copy(v_in.begin(), v_in.end(), raw_d.begin(), queue);
// Run the kernel
auto force_ave = in.metadata.is_final;
auto nb_ave_frames = ave_calc(queue, raw_d, ave_d, force_ave);
LIMA_LOG(trace, proc) << nb_ave_frames << " frames averaged";
if (nb_ave_frames == 0)
return;
LIMA_LOG(trace, proc) << "Average available";
auto dims = in.dimensions();
lima::frame ave(dims, pixel_enum::gray32f);
// Copy metadata - TODO: include nb_ave_frames
ave.metadata = in.metadata;
ave.metadata.idx = ave_frame_idx;
ave.attributes = in.attributes;
// Copy data
auto view = lima::view<lima::gray32f_view_t>(ave);
bcl::copy(ave_d.begin(), ave_d.end(), &boost::gil::at_c<0>(view(0, 0)), queue);
++ave_frame_idx;
std::get<0>(ports).try_put(ave);
}
bcl::context context;
mutable bcl::command_queue queue;
std::size_t nb_pixels;
fai::average ave_calc;
mutable fai::vector<std::uint16_t> raw_d; // Input
mutable fai::vector<float> ave_d; // Output average
mutable std::size_t ave_frame_idx{0}; //
};
average_node(tbb::flow::graph& g, std::size_t width, std::size_t height, pixel_enum pixel,
std::size_t start_frame, std::size_t nb_frames, std::filesystem::path const& cl_source_path,
bcl::context& context, bcl::command_queue& queue) :
parent_t(g, tbb::flow::serial /*serial fo now*/,
average_body(width, height, start_frame, nb_frames, cl_source_path, context, queue))
{
if (pixel != pixel_enum::gray16)
LIMA_THROW_EXCEPTION(lima::invalid_argument("average_node only supports gray16 pixel type"));
}
};
} // namespace smx
} // namespace processing::pipelines
} // namespace lima
// Copyright (C) 2020 Alejandro Homs Puron, ESRF.
// Use, modification and distribution is subject to the Boost Software
// License, Version 1.0. (See accompanying file LICENSE_1_0.txt or copy at
// http://www.boost.org/LICENSE_1_0.txt)
#pragma once
#include <boost/describe.hpp>
#include <boost/describe/io_enums.hpp>
#include <lima/processing/pipelines/smx/enums.hpp>
namespace lima
{
namespace processing::pipelines
{
namespace smx
{
BOOST_DESCRIBE_ENUM(proc_mode_enum, fai, average)
using boost::describe::operator<<;
using boost::describe::operator>>;
} // namespace smx
} // namespace processing::pipelines
} // namespace lima
// Copyright (C) 2020 Alejandro Homs Puron, ESRF.
// Use, modification and distribution is subject to the Boost Software
// License, Version 1.0. (See accompanying file LICENSE_1_0.txt or copy at
// http://www.boost.org/LICENSE_1_0.txt)
#pragma once
namespace lima
{
namespace processing::pipelines
{
namespace smx
{
/// Frame (reconstruction) assembly mode
enum class proc_mode_enum : int
{
fai,
average
};
} // namespace smx
} // namespace processing::pipelines
} // namespace lima
......@@ -10,6 +10,7 @@
#include <lima/processing/fai/enums.describe.hpp>
#include <lima/processing/pipelines/smx/params.hpp>
#include <lima/processing/pipelines/smx/enums.describe.hpp>
// This file is part of the user interface of the library and should not include any private dependencies
// or other thing that expose implementation details
......@@ -64,6 +65,18 @@ namespace processing::pipelines
(doc, "The path to the OpenCL kernels"))
// clang-format on
BOOST_DESCRIBE_STRUCT(average_params, (), (start_frame, nb_frames))
// clang-format off
BOOST_ANNOTATE_MEMBER(average_params, start_frame,
(desc, "start frame"),
(doc, "The first frame to start average (nb of frames to skip)"))
BOOST_ANNOTATE_MEMBER(average_params, nb_frames,
(desc, "nb of frames"),
(doc, "The number of frames to average"))
// clang-format on
BOOST_DESCRIBE_STRUCT(gain_pedestal_params, (), (gain_path, pedestal_path, photon_energy))
// clang-format off
......@@ -177,10 +190,15 @@ namespace processing::pipelines
// clang-format on
BOOST_DESCRIBE_STRUCT(proc_params, (),
(fifo, buffers, saving_dense, saving_sparse, saving_accumulation_corrected,
saving_accumulation_peak, gpu, jfrau, fai, counters))
(proc_mode, fifo, buffers, saving_dense, saving_average, saving_sparse,
saving_accumulation_corrected, saving_accumulation_peak, gpu, average, jfrau, fai,
counters))
// clang-format off
BOOST_ANNOTATE_MEMBER(proc_params, proc_mode,
(desc, "processing mode"),
(doc, "The processing mode: fai, average"))
BOOST_ANNOTATE_MEMBER(proc_params, fifo,
(desc, "fifo parameters"),
(doc, "The processing FIFO parameters"))
......@@ -193,6 +211,10 @@ namespace processing::pipelines
(desc, "saving parameters raw"),
(doc, "The HDF5 saving parameters for the raw frame"))
BOOST_ANNOTATE_MEMBER(proc_params, saving_average,
(desc, "saving parameters averaged"),
(doc, "The HDF5 saving parameters for the averaged frame"))
BOOST_ANNOTATE_MEMBER(proc_params, saving_sparse,
(desc, "saving parameters assembled"),
(doc, "The HDF5 saving parameters for the assembled frame"))
......@@ -209,6 +231,10 @@ namespace processing::pipelines
(desc, "gpu params"),
(doc, "The GPU (OpenCL) parameters"))
BOOST_ANNOTATE_MEMBER(proc_params, average,
(desc, "average params"),
(doc, "The averaging parameters"))
BOOST_ANNOTATE_MEMBER(proc_params, jfrau,
(desc, "jungfrau gain pedestal params"),
(doc, "The Gain / Pesdetal correction parameters"))
......
......@@ -12,6 +12,8 @@
#include <lima/processing/fai/enums.hpp>
#include <lima/processing/pipelines/smx/enums.hpp>
// This file is part of the user interface of the library and should not include any private dependencies
// or other thing that expose implementation details
......@@ -45,6 +47,12 @@ namespace processing::pipelines
std::filesystem::path cl_source_path = "detectors/processing/psi/src"; // CL source file path
};
struct average_params
{
std::size_t start_frame = 0;
std::size_t nb_frames = 0;
};
struct gain_pedestal_params
{
std::filesystem::path gain_path; // JungFrau Gain/Pedestal
......@@ -84,13 +92,16 @@ namespace processing::pipelines
struct proc_params
{
proc_mode_enum proc_mode = proc_mode_enum::fai;
fifo_params fifo;
buffer_params buffers;
saving_params saving_dense;
saving_params saving_average;
saving_params saving_sparse;
saving_params saving_accumulation_corrected;
saving_params saving_accumulation_peak;
gpu_params gpu;
average_params average;
gain_pedestal_params jfrau;
fai_params fai;
roi_counters_params counters;
......
......@@ -20,7 +20,7 @@ namespace processing::pipelines
BOOST_DESCRIBE_STRUCT(counters, (),
(nb_frames_source, nb_frames_reconstructed, nb_frames_counters, nb_frames_sparsified,
nb_frames_sparse_saved, nb_frames_dense_saved, nb_frames_acc_corrected_saved,
nb_frames_acc_peak_saved))
nb_frames_acc_peak_saved, nb_frames_average_generated, nb_frames_average_saved))
// clang-format off
BOOST_ANNOTATE_MEMBER(counters, nb_frames_source,
......@@ -54,6 +54,14 @@ namespace processing::pipelines
BOOST_ANNOTATE_MEMBER(counters, nb_frames_acc_peak_saved,
(desc, "nb frames accumulation peak saved"),
(doc, "The number of peak accumulation frame saved to file"))
BOOST_ANNOTATE_MEMBER(counters, nb_frames_average_generated,
(desc, "nb frames average generated"),
(doc, "The number of average frames generated"))
BOOST_ANNOTATE_MEMBER(counters, nb_frames_average_saved,
(desc, "nb frames average saved"),
(doc, "The number of average frames saved to file"))
// clang-format on