detectors/aare
25
0
Fork 0
mirror of https://github.com/slsdetectorgroup/aare.git synced 2026-09-02 23:20:43 +02:00
Code Issues Actions Packages Releases Activity

CUDA: expose in docs, drop dead headers
Build on RHEL8 / build (push) Successful in 4m12s
Details
Build on RHEL9 / build (push) Successful in 4m44s
Details
Run tests using data on local RHEL8 / build (push) Failing after 4m18s
Details

Browse Source 
Doxygen never read *.cuh and needs EXTENSION_MAPPING to parse it; add
both, plus C++/Python pages and AARE_CUDA build docs.

Delete ClusterFinderCUDA_old.hpp (declared the same class, polluting
the API page), utils/batch.hpp, and the deck. All unreferenced.

TODO:
- Frozen + CUDAGraph: bound in Python, headers not installed. Ship or drop?
- test_ClusterFinderCUDA.py guarded on _cuda_available()
- AARE_CUDA_ARCHITECTURES=native needs CMake 3.24, project min is 3.15
- make docs somewhere with sphinx+breathe; new .rst files never rendered
- fmt nvcc warnings: gone at -std=c++20 or with -isystem
This commit is contained in:
kferjaoui
2026-09-02 15:00:39 +02:00
parent 1d963a778b
commit 7870ccfc32
12 changed files with 681 additions and 1062 deletions
Download Patch File Download Diff File Expand all files Collapse all files
docs/ClusterFinderCUDA_optimizations.pptx
BIN
View File
Binary file not shown.
docs/Doxyfile.in
+3 -1
View File
@@ -288,7 +288,7 @@ OPTIMIZE_OUTPUT_VHDL = NO
# Note that for custom extensions you also need to set FILE_PATTERNS otherwise
# the files are not read by doxygen.
EXTENSION_MAPPING =
EXTENSION_MAPPING = cu=C++ cuh=C++
# If the MARKDOWN_SUPPORT tag is enabled then doxygen pre-processes all comments
# according to the Markdown format, which allows for more readable
@@ -831,6 +831,8 @@ FILE_PATTERNS = *.c \
*.hxx \
*.hpp \
*.h++ \
*.cu \
*.cuh \
*.cs \
*.d \
*.php \
docs/src/ClusterFinderCUDA.rst
+26
View File
@@ -0,0 +1,26 @@
ClusterFinderCUDA
=================
GPU cluster finder. Available only when aare is configured with
``-DAARE_CUDA=ON``; on a CPU-only build the class is not compiled and the
Python factory raises ``RuntimeError``.
Single frames go through :cpp:func:`aare::ClusterFinderCUDA::find_clusters`,
which mirrors the CPU :cpp:class:`aare::ClusterFinder` interface. Throughput
comes from ``find_clusters_batched``, which distributes a batch of frames
round-robin over ``n_streams`` CUDA streams (4 by default) and overlaps H2D,
kernel and D2H work. Results are read back either with ``collect``, which
returns one ``ClusterVector`` per frame, or with ``collect_view``, which hands
back a ``BatchView`` reading clusters in place from the pinned host buffer with
no per-frame allocation or copy.
.. doxygenclass:: aare::ClusterFinderCUDA
:members:
:undoc-members:
Device kernel
-------------
.. doxygennamespace:: aare::device
:members:
:undoc-members:
docs/src/Installation.rst
+33
View File
@@ -109,6 +109,39 @@ Build the Python bindings. Default option is off.
If you have a newer system Python compared to the one in your virtual environment,
you might have to pass -DPython_FIND_VIRTUALENV=ONLY to cmake.
**AARE_CUDA "Build CUDA cluster finder backend" OFF**
Build :doc:`ClusterFinderCUDA`, the GPU cluster finder. Default option is off,
in which case aare builds CPU-only and the class is not compiled at all.
Requires the CUDA toolkit (nvcc) and a compiler combination nvcc accepts; see
:doc:`Requirements`. With ``AARE_PYTHON_BINDINGS=ON`` the CUDA bindings are
built as a separate extension module, ``_aare_cuda``, which avoids nvcc/gcc ABI
conflicts in the main module. The classes are re-exported onto ``aare``, so
imports stay the same either way:
.. code-block:: bash
cmake ../aare -DAARE_CUDA=ON -DAARE_PYTHON_BINDINGS=ON
make -j4
.. code-block:: python
from aare import ClusterFinderCUDA, _cuda_available
On a CPU-only build ``_cuda_available()`` returns False and the
``ClusterFinderCUDA`` factory raises ``RuntimeError``.
**AARE_CUDA_ARCHITECTURES "CUDA architectures to compile for" native**
Only used when ``AARE_CUDA=ON``. Defaults to ``native``, which targets the GPU
in the building machine. Set it explicitly to build portable binaries or to
cross-compile for a GPU you are not building on:
.. code-block:: bash
cmake ../aare -DAARE_CUDA=ON -DAARE_CUDA_ARCHITECTURES="80;89"
**AARE_TESTS "Build tests" OFF**
Build unit tests. Default option is off.
docs/src/Requirements.rst
+11
View File
@@ -19,6 +19,17 @@ To simplify deployment we build and statically link a few libraries.
- pybind11
- ZeroMQ
**Extra dependencies for the CUDA backend (-DAARE_CUDA=ON)**
Only needed to build :doc:`ClusterFinderCUDA`; aare builds CPU-only without them.
- CUDA toolkit 11.0 or newer (nvcc), and a host compiler that nvcc supports.
The device code is compiled as CUDA C++17, which nvcc supports from CUDA 11.
- CMake 3.24+ if you keep the default ``AARE_CUDA_ARCHITECTURES=native``;
``native`` is not understood by older CMake. With an explicit architecture
list the project minimum of 3.15 is enough.
- An NVIDIA GPU at runtime.
**Extra dependencies for building documentation**
- Sphinx
docs/src/index.rst
+1
View File
@@ -49,6 +49,7 @@ AARE
Cluster
ClusterFinder
ClusterFinderMT
ClusterFinderCUDA
ClusterFile
ClusterVector
Interpolation
docs/src/python/cluster/index.rst
+1
View File
@@ -7,5 +7,6 @@ Cluster & Interpolation
pyCluster
pyClusterVector
pyClusterFinderCUDA
pyInterpolation
pyVarClusterFinder
docs/src/python/cluster/pyClusterFinderCUDA.rst
+11
View File
@@ -0,0 +1,11 @@
ClusterFinderCUDA
=================
.. py:currentmodule:: aare
Requires a build configured with ``-DAARE_CUDA=ON``. On a CPU-only build the
factory raises ``RuntimeError``.
.. autofunction:: ClusterFinderCUDA
.. autofunction:: find_cluster_views_batched_iter
include/aare/ClusterFinderCUDA_old.hpp
-435
View File
@@ -1,435 +0,0 @@
// SPDX-License-Identifier: MPL-2.0
#pragma once
#include "aare/ClusterFinder.hpp"
#include "aare/clusterfinder_kernel.cuh"
#include "aare/utils/cuda_check.cuh"
#include <algorithm>
#include <cmath>
#include <cstdint>
#include <cstdio>
#include <cstring>
#include <limits>
#include <stdexcept>
namespace aare {
// Per-stream device resources
template <typename ClusterType, typename FRAME_TYPE, typename PEDESTAL_TYPE>
struct StreamContext {
cudaStream_t stream = nullptr; // handle to the stream
FRAME_TYPE *d_frame = nullptr;
PEDESTAL_TYPE *d_pd_mean = nullptr;
PEDESTAL_TYPE *d_pd_sum = nullptr;
PEDESTAL_TYPE *d_pd_sum2 = nullptr;
ClusterType *d_clusters = nullptr;
uint32_t *d_cluster_count = nullptr;
// Pinned host staging buffers. These make cudaMemcpyAsync real async DMA
// transfers even when the caller's NDView points to pageable memory.
FRAME_TYPE *h_frame = nullptr;
uint32_t *h_cluster_count = nullptr;
ClusterType *h_clusters = nullptr;
cudaEvent_t kernel_start = nullptr;
cudaEvent_t kernel_stop = nullptr;
};
template <typename ClusterType = Cluster<int32_t, 3, 3>,
typename FRAME_TYPE = uint16_t, typename PEDESTAL_TYPE = double,
typename = std::enable_if_t<no_2x2_cluster<ClusterType>::value>>
class ClusterFinderCUDA {
using COMPUTE_TYPE =
device::COMPUTE_TYPE; // match the kernel's internal precision
static constexpr int BLOCK_X = 16;
static constexpr int BLOCK_Y = 16;
static constexpr int col_radius = ClusterType::cluster_size_x / 2;
static constexpr int row_radius = ClusterType::cluster_size_y / 2;
Shape<2> m_shape;
size_t nrows;
size_t ncols;
size_t m_image_size; // nrows * ncols
int n_streams;
size_t m_capacity;
size_t m_image_bytes;
size_t m_cluster_bytes;
COMPUTE_TYPE m_nSigma;
Pedestal<PEDESTAL_TYPE> m_pedestal;
ClusterVector<ClusterType> m_clusters;
bool m_pedestal_dirty = true;
using SC = StreamContext<ClusterType, FRAME_TYPE, PEDESTAL_TYPE>;
std::vector<SC> v_sc;
float m_total_kernel_ms = 0.0f;
size_t m_frames_processed = 0;
// Kernel parameters
dim3 grid;
dim3 block;
size_t shmem_bytes;
public:
/**
* @brief Construct a ClusterFinderCUDA
*
* @param m_image_size shape of the detector frame (rows, cols)
* @param nSigma threshold in units of per-pixel pedestal
* std
* @param capacity device-side cluster buffer size per stream
* @param n_streams number of CUDA streams for multi-frame
* overlap
*/
ClusterFinderCUDA(Shape<2> shape_, COMPUTE_TYPE nSigma = 5.0,
size_t capacity = 1000000, int n_streams_ = 1)
: m_shape(shape_), nrows(shape_[0]), ncols(shape_[1]),
m_image_size(nrows * ncols), n_streams(n_streams_),
m_capacity(capacity), m_nSigma(nSigma),
m_pedestal(shape_[0], shape_[1]), m_clusters(capacity) {
if (n_streams_ <= 0) {
throw std::invalid_argument(
"ClusterFinderCUDA: n_streams must be > 0");
}
if (capacity >
static_cast<size_t>(std::numeric_limits<uint32_t>::max())) {
throw std::invalid_argument(
"ClusterFinderCUDA: capacity must fit in uint32_t");
}
if (capacity == 0) {
throw std::invalid_argument(
"ClusterFinderCUDA: capacity must be > 0");
}
// Grid/Block dimensions
block = dim3(BLOCK_X, BLOCK_Y);
grid = dim3((static_cast<unsigned int>(ncols) + BLOCK_X - 1) / BLOCK_X,
(static_cast<unsigned int>(nrows) + BLOCK_Y - 1) / BLOCK_Y);
// Shared memory: one tile of (BLOCK_X + 2*col_radius) x (BLOCK_Y +
// 2*row_radius) elements
// Mixed precision used -> shmem takes COMPUTE_TYPE = floats (not
// PEDESTAL_TYPE)
shmem_bytes = (BLOCK_X + 2 * col_radius) * (BLOCK_Y + 2 * row_radius) *
sizeof(COMPUTE_TYPE);
m_image_bytes = m_image_size * sizeof(FRAME_TYPE);
m_cluster_bytes = m_capacity * sizeof(ClusterType);
v_sc.resize(n_streams);
for (int k = 0; k < n_streams; ++k) {
auto &sc = v_sc[k];
CUDA_CHECK(
cudaStreamCreateWithFlags(&sc.stream, cudaStreamNonBlocking));
CUDA_CHECK(cudaEventCreate(&sc.kernel_start));
CUDA_CHECK(cudaEventCreate(&sc.kernel_stop));
CUDA_CHECK(cudaMalloc(&sc.d_frame, m_image_bytes));
CUDA_CHECK(cudaMalloc(&sc.d_pd_mean,
m_image_size * sizeof(PEDESTAL_TYPE)));
CUDA_CHECK(
cudaMalloc(&sc.d_pd_sum, m_image_size * sizeof(PEDESTAL_TYPE)));
CUDA_CHECK(cudaMalloc(&sc.d_pd_sum2,
m_image_size * sizeof(PEDESTAL_TYPE)));
CUDA_CHECK(cudaMalloc(&sc.d_clusters, m_cluster_bytes));
CUDA_CHECK(cudaMalloc(&sc.d_cluster_count, sizeof(uint32_t)));
CUDA_CHECK(cudaMallocHost(reinterpret_cast<void **>(&sc.h_frame),
m_image_bytes));
CUDA_CHECK(
cudaMallocHost(reinterpret_cast<void **>(&sc.h_cluster_count),
sizeof(uint32_t)));
if (m_cluster_bytes > 0) {
CUDA_CHECK(
cudaMallocHost(reinterpret_cast<void **>(&sc.h_clusters),
m_cluster_bytes));
}
}
}
~ClusterFinderCUDA() {
for (auto &sc : v_sc) {
if (sc.stream)
cudaStreamSynchronize(sc.stream);
if (sc.d_frame)
cudaFree(sc.d_frame);
if (sc.d_pd_mean)
cudaFree(sc.d_pd_mean);
if (sc.d_pd_sum)
cudaFree(sc.d_pd_sum);
if (sc.d_pd_sum2)
cudaFree(sc.d_pd_sum2);
if (sc.d_clusters)
cudaFree(sc.d_clusters);
if (sc.d_cluster_count)
cudaFree(sc.d_cluster_count);
if (sc.h_frame)
cudaFreeHost(sc.h_frame);
if (sc.h_clusters)
cudaFreeHost(sc.h_clusters);
if (sc.h_cluster_count)
cudaFreeHost(sc.h_cluster_count);
if (sc.kernel_start)
cudaEventDestroy(sc.kernel_start);
if (sc.kernel_stop)
cudaEventDestroy(sc.kernel_stop);
if (sc.stream)
cudaStreamDestroy(sc.stream);
}
}
// Non-copyable, non-movable
ClusterFinderCUDA(const ClusterFinderCUDA &) = delete;
ClusterFinderCUDA &operator=(const ClusterFinderCUDA &) = delete;
ClusterFinderCUDA(ClusterFinderCUDA &&) = delete;
ClusterFinderCUDA &operator=(ClusterFinderCUDA &&) = delete;
void set_nSigma(COMPUTE_TYPE nSigma) { m_nSigma = nSigma; }
COMPUTE_TYPE get_nSigma() const { return m_nSigma; }
void push_pedestal_frame(NDView<FRAME_TYPE, 2> frame) {
m_pedestal.push(frame);
m_pedestal_dirty = true;
}
void clear_pedestal() {
m_pedestal.clear();
m_pedestal_dirty = true;
}
NDArray<PEDESTAL_TYPE, 2> pedestal() { return m_pedestal.mean(); }
NDArray<PEDESTAL_TYPE, 2> noise() { return m_pedestal.std(); }
/**
* @brief Move clusters out of the internal ClusterVector, optionally
* reallocating the internal one with the same capacity.
*/
ClusterVector<ClusterType>
steal_clusters(bool realloc_same_capacity = false) {
ClusterVector<ClusterType> tmp = std::move(m_clusters);
if (realloc_same_capacity)
m_clusters = ClusterVector<ClusterType>(tmp.capacity());
else
m_clusters = ClusterVector<ClusterType>{};
return tmp;
}
/**
* @brief Find clusters in a single frame, appending them to the internal
* ClusterVector.
*/
void find_clusters(NDView<FRAME_TYPE, 2> frame, uint64_t frame_number = 0) {
if (m_pedestal_dirty) { // need to update the pedestal on the gpu
sync_pedestal_to_device();
m_pedestal_dirty = false;
}
auto &sc = v_sc[0];
const uint32_t n_pd_samples =
static_cast<uint32_t>(m_pedestal.n_samples());
// First, CPU copies frame into a reusable pinned buffer
std::memcpy(sc.h_frame, frame.data(), m_image_bytes);
// Reset cluster counter
CUDA_CHECK(cudaMemsetAsync(sc.d_cluster_count, 0, sizeof(uint32_t),
sc.stream));
// Upload frame
CUDA_CHECK(cudaMemcpyAsync(sc.d_frame, sc.h_frame, m_image_bytes,
cudaMemcpyHostToDevice, sc.stream));
// Timed Kernel launch
CUDA_CHECK(cudaEventRecord(sc.kernel_start, sc.stream));
device::find_clusters_in_single_frame<ClusterType, FRAME_TYPE,
PEDESTAL_TYPE>
<<<grid, block, shmem_bytes, sc.stream>>>(
sc.d_frame, sc.d_pd_mean, sc.d_pd_sum, sc.d_pd_sum2,
n_pd_samples, m_nSigma, nrows, ncols, sc.d_clusters,
sc.d_cluster_count, static_cast<uint32_t>(m_capacity));
CUDA_CHECK(cudaEventRecord(sc.kernel_stop, sc.stream));
CUDA_CHECK(cudaGetLastError());
// Read back cluster count into pinned buffer
CUDA_CHECK(cudaMemcpyAsync(sc.h_cluster_count, sc.d_cluster_count,
sizeof(uint32_t), cudaMemcpyDeviceToHost,
sc.stream));
// Synchronize to ensure count is available before the CPU reads
// clusters
CUDA_CHECK(cudaStreamSynchronize(sc.stream));
record_kernel_time(sc);
// Clamp to max in case of overflow
uint32_t n_found = *sc.h_cluster_count;
n_found = std::min(n_found, static_cast<uint32_t>(m_capacity));
// Read back clusters
m_clusters.set_frame_number(frame_number);
if (n_found > 0) {
append_device_clusters_to(m_clusters, sc, n_found);
}
}
/**
* @brief Batched cluster finding across multiple frames, using n_streams
* CUDA streams to overlap H2D transfer, kernel, and D2H transfer.
*
* Returns one ClusterVector per input frame (with frame_number set to
* first_frame + i).
*/
std::vector<ClusterVector<ClusterType>>
find_clusters_batched(NDView<FRAME_TYPE, 3> frames,
uint64_t first_frame = 0) {
if (m_pedestal_dirty) {
sync_pedestal_to_device();
m_pedestal_dirty = false;
}
const size_t n_frames = frames.shape(0);
const uint32_t n_pd_samples =
static_cast<uint32_t>(m_pedestal.n_samples());
std::vector<ClusterVector<ClusterType>> results;
results.reserve(n_frames);
for (size_t i = 0; i < n_frames; ++i) {
results.emplace_back();
results.back().set_frame_number(first_frame + i);
}
const size_t n_rounds = (n_frames + n_streams - 1) / n_streams;
for (size_t round = 0; round < n_rounds; ++round) {
// Launch phase: fan out kernels on all streams for this round
for (int k = 0; k < n_streams; ++k) {
// OOB guard
const size_t frame_idx = round * n_streams + k;
if (frame_idx >= n_frames)
continue;
auto &sc_k = v_sc[k];
const FRAME_TYPE *h_src =
frames.data() + frame_idx * m_image_size;
std::memcpy(sc_k.h_frame, h_src, m_image_bytes);
CUDA_CHECK(cudaMemsetAsync(sc_k.d_cluster_count, 0,
sizeof(uint32_t), sc_k.stream));
CUDA_CHECK(
cudaMemcpyAsync(sc_k.d_frame, sc_k.h_frame, m_image_bytes,
cudaMemcpyHostToDevice, sc_k.stream));
CUDA_CHECK(cudaEventRecord(sc_k.kernel_start, sc_k.stream));
device::find_clusters_in_single_frame<ClusterType, FRAME_TYPE,
PEDESTAL_TYPE>
<<<grid, block, shmem_bytes, sc_k.stream>>>(
sc_k.d_frame, sc_k.d_pd_mean, sc_k.d_pd_sum,
sc_k.d_pd_sum2, n_pd_samples, m_nSigma, nrows, ncols,
sc_k.d_clusters, sc_k.d_cluster_count,
static_cast<uint32_t>(m_capacity));
CUDA_CHECK(cudaEventRecord(sc_k.kernel_stop, sc_k.stream));
CUDA_CHECK(cudaGetLastError());
// Queue count D2H immediately after the kernel
CUDA_CHECK(cudaMemcpyAsync(
sc_k.h_cluster_count, sc_k.d_cluster_count,
sizeof(uint32_t), cudaMemcpyDeviceToHost, sc_k.stream));
}
// Drain phase: fan in results from all streams
for (int k = 0; k < n_streams; ++k) {
const size_t frame_idx = round * n_streams + k;
if (frame_idx >= n_frames)
continue;
auto &sc_k = v_sc[k];
// Wait for memset -> H2D -> kernel -> count D2H
CUDA_CHECK(cudaStreamSynchronize(sc_k.stream));
record_kernel_time(sc_k);
uint32_t n_found = *sc_k.h_cluster_count;
n_found = std::min<uint32_t>(n_found,
static_cast<uint32_t>(m_capacity));
if (n_found > 0) {
append_device_clusters_to(results[frame_idx], sc_k,
n_found);
}
}
}
return results;
}
float avg_kernel_time_ms() const {
return m_frames_processed > 0 ? m_total_kernel_ms / m_frames_processed
: 0.0f;
}
void reset_timers() {
m_total_kernel_ms = 0.0f;
m_frames_processed = 0;
}
private:
/**
* Upload the current host pedestal (mean, sum, sum2) to every stream's
* device buffers. Called lazily before a find_clusters call when the
* host pedestal has been updated.
*/
void sync_pedestal_to_device() {
// These return-by-value NDArrays must stay alive until the async
// copies complete, so we synchronise at the end before they go out
// of scope.
NDArray<PEDESTAL_TYPE, 2> h_mean = m_pedestal.mean();
NDArray<PEDESTAL_TYPE, 2> h_sum = m_pedestal.get_sum();
NDArray<PEDESTAL_TYPE, 2> h_sum2 = m_pedestal.get_sum2();
const size_t bytes = m_image_size * sizeof(PEDESTAL_TYPE);
for (auto &sc : v_sc) {
CUDA_CHECK(cudaMemcpyAsync(sc.d_pd_mean, h_mean.data(), bytes,
cudaMemcpyHostToDevice, sc.stream));
CUDA_CHECK(cudaMemcpyAsync(sc.d_pd_sum, h_sum.data(), bytes,
cudaMemcpyHostToDevice, sc.stream));
CUDA_CHECK(cudaMemcpyAsync(sc.d_pd_sum2, h_sum2.data(), bytes,
cudaMemcpyHostToDevice, sc.stream));
}
for (auto &sc : v_sc)
CUDA_CHECK(cudaStreamSynchronize(sc.stream));
}
/**
* Copy n_found clusters from sc.d_clusters into the given ClusterVector
* and block on the transfer.
*/
void append_device_clusters_to(ClusterVector<ClusterType> &cv, SC &sc,
uint32_t n_found) {
CUDA_CHECK(cudaMemcpyAsync(sc.h_clusters, sc.d_clusters,
n_found * sizeof(ClusterType),
cudaMemcpyDeviceToHost, sc.stream));
CUDA_CHECK(cudaStreamSynchronize(sc.stream));
for (uint32_t i = 0; i < n_found; ++i)
cv.push_back(sc.h_clusters[i]);
}
void record_kernel_time(SC &sc) {
float ms = 0.0f;
CUDA_CHECK(cudaEventElapsedTime(&ms, sc.kernel_start, sc.kernel_stop));
m_total_kernel_ms += ms;
m_frames_processed++;
}
};
} // namespace aare
include/aare/clusterfinder_kernel.cuh
+5 -7
View File
@@ -6,14 +6,12 @@
namespace aare::device {
// Device arithmetic precision.
// COMPUTE_TYPE : per-frame stencil arithmetic (shared-memory tile, sums).
// DEVICE_PED_TYPE : device pedestal storage + running variance update.
// Shipped as double/double so the GPU result matches the double-precision CPU
// ClusterFinder to the floating-point floor. float/float is ~2x faster but
// reintroduces a small near-threshold mismatch; a build-time toggle for that
// trade-off is planned.
/// Stencil arithmetic precision. Set both aliases to double for an f64 build.
/// float is safe here because the device pedestal stores centered moments
/// against a host-computed baseline: f32 and f64 agree to 3e-7 in cluster
/// count, and f32 cuts the 9x9 kernel time by ~40%.
using COMPUTE_TYPE = float;
/// Device pedestal storage and running variance update. See COMPUTE_TYPE.
using DEVICE_PED_TYPE = float;
template <typename ClusterType = Cluster<int32_t, 3, 3>,
include/aare/utils/batch.hpp
-28
View File
@@ -1,28 +0,0 @@
// SPDX-License-Identifier: MPL-2.0
#pragma once
#include "aare/NDArray.hpp"
#include <cstring>
#include <vector>
template <typename FRAME_TYPE>
void pack_frame_batch(const std::vector<aare::NDArray<FRAME_TYPE, 2>> &frames,
size_t first_frame, size_t n_frames,
std::vector<FRAME_TYPE> &batch) {
if (n_frames == 0)
return;
const size_t rows = frames[first_frame].shape(0);
const size_t cols = frames[first_frame].shape(1);
const size_t image_size = rows * cols;
const size_t total_size = n_frames * image_size;
if (batch.size() != total_size) {
batch.resize(total_size);
}
for (size_t k = 0; k < n_frames; ++k) {
const FRAME_TYPE *src = frames[first_frame + k].data();
FRAME_TYPE *dst = batch.data() + k * image_size;
std::memcpy(dst, src, image_size * sizeof(FRAME_TYPE));
}
}
python/tests/ClusterFinderCUDA.ipynb
+590 -591
View File
File diff suppressed because one or more lines are too long