From 922e76b427d2f3c572c357acaab676bd119e1218 Mon Sep 17 00:00:00 2001 From: Piotr Jarosik Date: Mon, 1 Nov 2021 07:11:53 +0100 Subject: [PATCH 1/8] From now on FCM gives information about the TARGET US4OEM, channel and frame. --- api/python/arrus/utils/imaging.py | 27 ++-- api/python/arrus/utils/us4r_remap_gpu.py | 50 ++++-- arrus/core/api/devices/DeviceId.h | 4 +- .../api/devices/us4r/FrameChannelMapping.h | 12 +- arrus/core/devices/probe/ProbeImpl.cpp | 100 ++++++------ arrus/core/devices/probe/ProbeImplTest.cpp | 50 +++--- .../devices/us4r/FrameChannelMappingImpl.cpp | 39 +++-- .../devices/us4r/FrameChannelMappingImpl.h | 21 +-- .../us4r/probeadapter/ProbeAdapterImpl.cpp | 49 +++--- .../probeadapter/ProbeAdapterImplTest.cpp | 150 +++++++++--------- arrus/core/devices/us4r/us4oem/Us4OEMImpl.cpp | 13 +- .../devices/us4r/us4oem/Us4OEMImplTest.cpp | 23 ++- 12 files changed, 286 insertions(+), 252 deletions(-) diff --git a/api/python/arrus/utils/imaging.py b/api/python/arrus/utils/imaging.py index 33a006ac5..5cde98fa1 100644 --- a/api/python/arrus/utils/imaging.py +++ b/api/python/arrus/utils/imaging.py @@ -1713,38 +1713,33 @@ def prepare(self, const_metadata: arrus.metadata.ConstMetadata): n_frames, n_channels = fcm.frames.shape n_samples_set = {op.rx.get_n_samples() for op in const_metadata.context.raw_sequence.ops} + + # get (unique) number of samples in a frame if len(n_samples_set) > 1: raise arrus.exceptions.IllegalArgumentError( f"Each tx/rx in the sequence should acquire the same number of " f"samples (actual: {n_samples_set})") n_samples = next(iter(n_samples_set)) - self.output_shape = (n_frames, n_samples, n_channels) + batch_size = fcm.batch_size + self.output_shape = (batch_size, n_frames, n_samples, n_channels) self._output_buffer = xp.zeros(shape=self.output_shape, dtype=xp.int16) - - n_samples_raw, n_channels_raw = const_metadata.input_shape - self._input_shape = (n_samples_raw//n_samples, n_samples, - n_channels_raw) - self.batch_size = fcm.batch_size - if xp == np: # CPU - self._transfers = __group_transfers(fcm) - def cpu_remap_fn(data): - __remap(self._output_buffer, - data.reshape(self._input_shape), - transfers=self._transfers) - self._remap_fn = cpu_remap_fn + raise ValueError(f"'{type(self).__name__}' is not implemented for CPU") else: # GPU import cupy as cp from arrus.utils.us4r_remap_gpu import get_default_grid_block_size, run_remap self._fcm_frames = cp.asarray(fcm.frames) self._fcm_channels = cp.asarray(fcm.channels) - self.grid_size, self.block_size = get_default_grid_block_size(self._fcm_frames, n_samples) + self._fcm_us4oems = cp.asarray(fcm.us4oems) + self.grid_size, self.block_size = get_default_grid_block_size( + self._fcm_frames, n_samples, + batch_size + ) def gpu_remap_fn(data): - run_remap( - self.grid_size, self.block_size, + run_remap(self.grid_size, self.block_size, [self._output_buffer, data, self._fcm_frames, self._fcm_channels, n_frames, n_samples, n_channels]) diff --git a/api/python/arrus/utils/us4r_remap_gpu.py b/api/python/arrus/utils/us4r_remap_gpu.py index 434e2f0e6..450ad14ab 100644 --- a/api/python/arrus/utils/us4r_remap_gpu.py +++ b/api/python/arrus/utils/us4r_remap_gpu.py @@ -1,29 +1,53 @@ import cupy as cp +# TODO strategy for case batch size == 1 + _arrus_remap_str = r''' // Naive implementation of data remapping (physical -> logical order). extern "C" __global__ void arrus_remap(short* out, short* in, const short* fcmFrames, const char* fcmChannels, - const unsigned nFrames, const unsigned nSamples, const unsigned nChannels) + const unsigned char *fcmUs4oems, + const int frameOffsets, + // Output shape + const unsigned nSequences, const unsigned nFrames, const unsigned nSamples, const unsigned nChannels) { - int x = blockIdx.x * 32 + threadIdx.x; // logical channel - int y = blockIdx.y * 32 + threadIdx.y; // logical sample - int z = blockIdx.z; // logical frame - if(x >= nChannels || y >= nSamples || z >= nFrames) { + int channel = blockIdx.x * 32 + threadIdx.x; // logical channel + int sample = blockIdx.y * 32 + threadIdx.y; // logical sample + int frame = blockIdx.z; // logical frame, global in the whole batch of sequences + // Determine sequence number (in batch) and frame number (within sequence) + int sequence = frame / batchSize; + int localFrame = frame % batchSize; + if(channel >= nChannels || sample >= nSamples || localFrame >= nFrames || sequence >= nSequences) { // outside the range return; } - int indexOut = x + y*nChannels + z*nChannels*nSamples; - int physicalChannel = fcmChannels[x + nChannels*z]; + // FCM describes here a single sequence + int physicalChannel = fcmChannels[channel + nChannels*localFrame]; if(physicalChannel < 0) { // channel is turned off return; } - int physicalFrame = fcmFrames[x + nChannels*z]; + + // [sequence, frame, sample, channel] + int indexOut = sequence*nFrames*nSamples*nChannels + + frame*nSamples*nChannels + + sample*nChannels + + channel; + + int physicalFrame = fcmFrames[channel + nChannels*localFrame]; // 32 - number of channels in the physical mapping - int indexIn = physicalChannel + y*32 + physicalFrame*32*nSamples; + // [us4oem, sequence, physicalFrame, sample, physicalChannel] + + int us4oem = fcmUs4oems[channel + nChannels*localFrame]; + int us4oemOffset = frameOffsets[us4oem]; + + int indexIn = us4oemOffset // nbytes + // physicalFrame should be calculated relative to the us4oem module begin (first acquired frame should be 0) + + physicalFrame*nSamples*32 + + sample*32 + + physicalChannel; out[indexOut] = in[indexIn]; }''' @@ -31,11 +55,15 @@ remap_kernel = cp.RawKernel(_arrus_remap_str, "arrus_remap") -def get_default_grid_block_size(fcm_frames, n_samples): +def get_default_grid_block_size(fcm_frames, n_samples, batch_size): # Note the kernel implementation block_size = (32, 32) n_frames, n_channels = fcm_frames.shape - grid_size = (int((n_channels - 1) // block_size[0] + 1), int((n_samples - 1) // block_size[1] + 1), n_frames) + grid_size = ( + (n_channels - 1) // block_size[0] + 1, + (n_samples - 1) // block_size[1] + 1, + n_frames*batch_size + ) return (grid_size, block_size) diff --git a/arrus/core/api/devices/DeviceId.h b/arrus/core/api/devices/DeviceId.h index 4cf1d3dce..25a852aa1 100644 --- a/arrus/core/api/devices/DeviceId.h +++ b/arrus/core/api/devices/DeviceId.h @@ -49,9 +49,7 @@ using Ordinal = unsigned short; */ class DeviceId { public: - DeviceId(const DeviceType dt, - const Ordinal ordinal) - : deviceType(dt), ordinal(ordinal) {} + DeviceId(const DeviceType dt, const Ordinal ordinal) : deviceType(dt), ordinal(ordinal) {} DeviceType getDeviceType() const { return deviceType; diff --git a/arrus/core/api/devices/us4r/FrameChannelMapping.h b/arrus/core/api/devices/us4r/FrameChannelMapping.h index 1ec6423a3..dfbc8d24c 100644 --- a/arrus/core/api/devices/us4r/FrameChannelMapping.h +++ b/arrus/core/api/devices/us4r/FrameChannelMapping.h @@ -14,19 +14,21 @@ class FrameChannelMapping { public: using Handle = std::unique_ptr; using SharedHandle = std::shared_ptr; + // Frame Channel Mapping supports up to 256 Us4OEMs. + using Us4OEMNumber = uint8; using FrameNumber = uint16; constexpr static int8 UNAVAILABLE = -1; + /** - * Returns physical frame number and channel number for a given, - * logical, frame number and a **rx aperture** channel. + * Returns us4oem module number, physical frame number and channel number for a given, + * logical, frame number and an **rx aperture** channel. * * @param frame logical frame number * @param channel logical channel number - * @return actual frame number and channel number + * @return a tuple: us4oem module number, frame number (within a single sequence), channel number */ - // TODO use FrameNumber typedef (current implementation is simplified for swig) - virtual std::pair getLogical(FrameNumber frame, ChannelIdx channel) = 0; + virtual std::tuple getLogical(FrameNumber frame, ChannelIdx channel) = 0; virtual FrameNumber getNumberOfLogicalFrames() = 0; virtual ChannelIdx getNumberOfLogicalChannels() = 0; diff --git a/arrus/core/devices/probe/ProbeImpl.cpp b/arrus/core/devices/probe/ProbeImpl.cpp index 51bda1e61..3291db95d 100644 --- a/arrus/core/devices/probe/ProbeImpl.cpp +++ b/arrus/core/devices/probe/ProbeImpl.cpp @@ -50,15 +50,10 @@ class ProbeTxRxValidator : public Validator { }; std::tuple -ProbeImpl::setTxRxSequence(const std::vector &seq, - const ops::us4r::TGCCurve &tgcSamples, - uint16 rxBufferSize, - uint16 rxBatchSize, std::optional sri, - bool triggerSync) { +ProbeImpl::setTxRxSequence(const std::vector &seq, const ops::us4r::TGCCurve &tgcSamples, + uint16 rxBufferSize, uint16 rxBatchSize, std::optional sri, bool triggerSync) { // Validate input sequence - ProbeTxRxValidator validator( - ::arrus::format("tx rx sequence for {}", getDeviceId().toString()), - model); + ProbeTxRxValidator validator(format("tx rx sequence for {}", getDeviceId().toString()), model); validator.validate(seq); validator.throwOnErrors(); @@ -69,7 +64,7 @@ ProbeImpl::setTxRxSequence(const std::vector &seq, // Each vector contains mapping: // probe's rx aperture element number -> adapter rx aperture channel number - // Where each element and channel is the active bit element/channel number. + // Where each element is the active bit element/channel number. std::vector> rxApertureChannelMappings; // TODO the below list is used only in the remapFcm function, consider simplifying it @@ -77,8 +72,7 @@ ProbeImpl::setTxRxSequence(const std::vector &seq, std::vector rxPaddingRight; for (const auto &op: seq) { - logger->log(LogSeverity::TRACE, arrus::format( - "Setting tx/rx {}", ::arrus::toString(op))); + logger->log(LogSeverity::TRACE, format("Setting tx/rx {}", ::arrus::toString(op))); std::vector rxApertureChannelMapping; BitMask txAperture(adapter->getNumberOfChannels()); @@ -87,11 +81,10 @@ ProbeImpl::setTxRxSequence(const std::vector &seq, ARRUS_REQUIRES_TRUE( op.getTxAperture().size() == op.getRxAperture().size() - && op.getTxAperture().size() == op.getTxDelays().size() - && op.getTxAperture().size() == probeNumberOfElements, - arrus::format("Probe's tx, rx apertures and tx delays " - "array should have the same size: {}", - model.getNumberOfElements().product())); + && op.getTxAperture().size() == op.getTxDelays().size() + && op.getTxAperture().size() == probeNumberOfElements, + format("Probe's tx, rx apertures and tx delays array should have the same size: {}", + model.getNumberOfElements().product())); for (size_t pch = 0; pch < op.getTxAperture().size(); ++pch) { auto ach = channelMapping[pch]; @@ -103,21 +96,16 @@ ProbeImpl::setTxRxSequence(const std::vector &seq, rxApertureChannelMapping.push_back(ach); } } - adapterSeq.emplace_back(txAperture, txDelays, op.getTxPulse(), - rxAperture, op.getRxSampleRange(), - op.getRxDecimationFactor(), op.getPri(), - op.getRxPadding()); + adapterSeq.emplace_back(txAperture, txDelays, op.getTxPulse(), rxAperture, op.getRxSampleRange(), + op.getRxDecimationFactor(), op.getPri(), op.getRxPadding()); rxApertureChannelMappings.push_back(rxApertureChannelMapping); rxPaddingLeft.push_back(op.getRxPadding()[0]); rxPaddingRight.push_back(op.getRxPadding()[1]); } - auto[buffer, fcm] = adapter->setTxRxSequence(adapterSeq, tgcSamples, - rxBufferSize, rxBatchSize, - sri, triggerSync); - FrameChannelMapping::Handle actualFcm = remapFcm( - fcm, rxApertureChannelMappings, rxPaddingLeft, rxPaddingRight); + auto[buffer, fcm] = adapter->setTxRxSequence(adapterSeq, tgcSamples, rxBufferSize, rxBatchSize, sri, triggerSync); + FrameChannelMapping::Handle actualFcm = remapFcm(fcm, rxApertureChannelMappings, rxPaddingLeft, rxPaddingRight); return std::make_tuple(std::move(buffer), std::move(actualFcm)); } @@ -137,23 +125,32 @@ void ProbeImpl::syncTrigger() { adapter->syncTrigger(); } -void ProbeImpl::registerOutputBuffer(Us4ROutputBuffer *buffer, - const Us4RBuffer::Handle &us4rBuffer, +void ProbeImpl::registerOutputBuffer(Us4ROutputBuffer *buffer, const Us4RBuffer::Handle &us4rBuffer, ::arrus::ops::us4r::Scheme::WorkMode workMode) { adapter->registerOutputBuffer(buffer, us4rBuffer, workMode); } // Remaps FCM according to given rx aperture active channels mappings. -FrameChannelMapping::Handle ProbeImpl::remapFcm( - const FrameChannelMapping::Handle &adapterFcm, - const std::vector> &adapterActiveChannels, - const std::vector &rxPaddingLeft, - const std::vector &rxPaddingRight) -{ +/** + * This function reorders channels in FCM produced by ProbeAdapterImpl, so the order of channel + * is correct even in case of some permutation between probe and adapter channels (e.g. like + * for ALS probes - esaote adapters). + * + * Basically, this function reads adapter FCM and sets the order channels according to the mapping + * probe2AdpaterMap, which is + * probe's aperture channel number -> adapter's aperture channel number + * + * e.g. in the case of probe-adapter mapping: 1-3, 2-1, 3-2, 3-element aperture, + * the output FCM data (internal arrays) will be reordered 3, 1, 2 (i.e. probe's channel 1 will point to + * adapter channel 3, and so on). + */ +FrameChannelMapping::Handle ProbeImpl::remapFcm(const FrameChannelMapping::Handle &adapterFcm, + const std::vector> &adapterActiveChannels, + const std::vector &rxPaddingLeft, + const std::vector &rxPaddingRight) { auto nOps = adapterActiveChannels.size(); if (adapterFcm->getNumberOfLogicalFrames() != nOps) { - throw std::runtime_error( - "Inconsistent mapping and op number of probe's Rx apertures"); + throw std::runtime_error("Inconsistent mapping and op number of probe's Rx apertures"); } FrameChannelMappingBuilder builder(adapterFcm->getNumberOfLogicalFrames(), adapterFcm->getNumberOfLogicalChannels()); @@ -161,39 +158,42 @@ FrameChannelMapping::Handle ProbeImpl::remapFcm( unsigned short frameNumber = 0; for (const auto &mapping : adapterActiveChannels) { // mapping[i] = dst adapter channel number - // (e.g. from 0 to 256 (number of channels system have)) + // (e.g. from 0 to 256 (number of channels the system have)) // where i is the probe rx active element + // EXAMPLE: mapping = {3, 1, 10} auto paddingLeft = rxPaddingLeft[frameNumber]; auto paddingRight = rxPaddingRight[frameNumber]; - // pairs: channel position, adapter channel + // pairs: probe's APERTURE channel, adapter channel std::vector> posChannel; - // adapterRxChannel[i] = dst adapter aperture channel number - // (e.g. from 0 to 64 (aperture size)). auto nRxChannels = mapping.size(); - std::vector adapterRxChannel(nRxChannels, 0); + // probe2AdapterMap[i] = dst adapter aperture channel number (e.g. from 0 to 64 (aperture size)). + std::vector probe2AdapterMap(nRxChannels, 0); - std::transform(std::begin(mapping), std::end(mapping), - std::back_insert_iterator(posChannel), + std::transform(std::begin(mapping), std::end(mapping), std::back_insert_iterator(posChannel), [i = 0](ChannelIdx channel) mutable { return std::make_pair(static_cast(i++), channel); }); + // EXAMPLE: posChannel = {{0, 3}, {1, 1}, {2, 10}} std::sort(std::begin(posChannel), std::end(posChannel), [](const auto &a, const auto &b) { return a.second < b.second; }); + // Now the position in the vector `posChannel` is equal to the adapter APERTURE channel. + // EXAMPLE: posChannel = {{1, 1}, {0, 3}, {2, 10}} ChannelIdx i = 0; - for (const auto& pos_ch: posChannel) { - adapterRxChannel[std::get<0>(pos_ch)] = i++; + + // probe aperture channel -> adapter aperture channel + // EXAMPLE: probe2AdapterMap = {1, 0, 2} + for (const auto& posCh: posChannel) { + probe2AdapterMap[std::get<0>(posCh)] = i++; } - // probe aperture rx number -> adapter aperture rx number -> - // physical channel + // probe aperture rx number -> adapter aperture rx number -> physical channel auto nChannels = adapterFcm->getNumberOfLogicalChannels(); for (ChannelIdx pch = 0; pch < nChannels; ++pch) { if(pch >= paddingLeft && pch < (nChannels-paddingRight)) { - auto[physicalFrame, physicalChannel] = - adapterFcm->getLogical(frameNumber, adapterRxChannel[pch-paddingLeft]+paddingLeft); + auto [us4oem, physicalFrame, physicalChannel] = + adapterFcm->getLogical(frameNumber, probe2AdapterMap[pch-paddingLeft]+paddingLeft); - builder.setChannelMapping(frameNumber, pch, - physicalFrame, physicalChannel); + builder.setChannelMapping(frameNumber, pch, us4oem, physicalFrame, physicalChannel); } } diff --git a/arrus/core/devices/probe/ProbeImplTest.cpp b/arrus/core/devices/probe/ProbeImplTest.cpp index 6c9a0cc44..97bff84c5 100644 --- a/arrus/core/devices/probe/ProbeImplTest.cpp +++ b/arrus/core/devices/probe/ProbeImplTest.cpp @@ -13,15 +13,15 @@ class ProbeImplFcmRemapTest : public ::testing::Test { void SetUp() override { FrameChannelMappingBuilder fcmBuilder(N_FRAMES, N_CHANNELS); - fcmBuilder.setChannelMapping(0, 0, 0, 0); - fcmBuilder.setChannelMapping(0, 1, 0, 1); - fcmBuilder.setChannelMapping(0, 2, 1, 0); - fcmBuilder.setChannelMapping(0, 3, 1, 1); - - fcmBuilder.setChannelMapping(1, 0, 2, 0); - fcmBuilder.setChannelMapping(1, 1, 2, 1); - fcmBuilder.setChannelMapping(1, 2, 3, 0); - fcmBuilder.setChannelMapping(1, 3, 3, 1); + fcmBuilder.setChannelMapping(0, 0, 0, 0, 0); + fcmBuilder.setChannelMapping(0, 1, 0, 0, 1); + fcmBuilder.setChannelMapping(0, 2, 0, 1, 0); + fcmBuilder.setChannelMapping(0, 3, 0, 1, 1); + + fcmBuilder.setChannelMapping(1, 0, 0, 2, 0); + fcmBuilder.setChannelMapping(1, 1, 0, 2, 1); + fcmBuilder.setChannelMapping(1, 2, 0, 3, 0); + fcmBuilder.setChannelMapping(1, 3, 0, 3, 1); fcm = fcmBuilder.build(); rxPaddingLeft = std::vector(fcm->getNumberOfLogicalFrames(), 0); rxPaddingRight = std::vector(fcm->getNumberOfLogicalFrames(), 0); @@ -54,15 +54,15 @@ TEST_F(ProbeImplFcmRemapTest, OneToOne) { EXPECT_EQ(actualNFrames, N_FRAMES); EXPECT_EQ(actualNChannels, N_CHANNELS); - EXPECT_EQ(actualFcm->getLogical(0, 0), (std::pair(0, 0))); - EXPECT_EQ(actualFcm->getLogical(0, 1), (std::pair(0, 1))); - EXPECT_EQ(actualFcm->getLogical(0, 2), (std::pair(1, 0))); - EXPECT_EQ(actualFcm->getLogical(0, 3), (std::pair(1, 1))); + EXPECT_EQ(actualFcm->getLogical(0, 0), (std::tuple(0, 0, 0))); + EXPECT_EQ(actualFcm->getLogical(0, 1), (std::tuple(0, 0, 1))); + EXPECT_EQ(actualFcm->getLogical(0, 2), (std::tuple(0, 1, 0))); + EXPECT_EQ(actualFcm->getLogical(0, 3), (std::tuple(0, 1, 1))); - EXPECT_EQ(actualFcm->getLogical(1, 0), (std::pair(2, 0))); - EXPECT_EQ(actualFcm->getLogical(1, 1), (std::pair(2, 1))); - EXPECT_EQ(actualFcm->getLogical(1, 2), (std::pair(3, 0))); - EXPECT_EQ(actualFcm->getLogical(1, 3), (std::pair(3, 1))); + EXPECT_EQ(actualFcm->getLogical(1, 0), (std::tuple(0, 2, 0))); + EXPECT_EQ(actualFcm->getLogical(1, 1), (std::tuple(0, 2, 1))); + EXPECT_EQ(actualFcm->getLogical(1, 2), (std::tuple(0, 3, 0))); + EXPECT_EQ(actualFcm->getLogical(1, 3), (std::tuple(0, 3, 1))); } TEST_F(ProbeImplFcmRemapTest, NonStandard) { @@ -81,16 +81,16 @@ TEST_F(ProbeImplFcmRemapTest, NonStandard) { EXPECT_EQ(actualNFrames, N_FRAMES); EXPECT_EQ(actualNChannels, N_CHANNELS); - EXPECT_EQ(actualFcm->getLogical(0, 0), (std::pair(0, 0))); - EXPECT_EQ(actualFcm->getLogical(0, 1), (std::pair(0, 1))); - EXPECT_EQ(actualFcm->getLogical(0, 3), (std::pair(1, 0))); - EXPECT_EQ(actualFcm->getLogical(0, 2), (std::pair(1, 1))); + EXPECT_EQ(actualFcm->getLogical(0, 0), (std::tuple(0, 0, 0))); + EXPECT_EQ(actualFcm->getLogical(0, 1), (std::tuple(0, 0, 1))); + EXPECT_EQ(actualFcm->getLogical(0, 3), (std::tuple(0, 1, 0))); + EXPECT_EQ(actualFcm->getLogical(0, 2), (std::tuple(0, 1, 1))); // Change - EXPECT_EQ(actualFcm->getLogical(1, 3), (std::pair(2, 0))); - EXPECT_EQ(actualFcm->getLogical(1, 1), (std::pair(2, 1))); - EXPECT_EQ(actualFcm->getLogical(1, 2), (std::pair(3, 0))); - EXPECT_EQ(actualFcm->getLogical(1, 0), (std::pair(3, 1))); + EXPECT_EQ(actualFcm->getLogical(1, 3), (std::tuple(0, 2, 0))); + EXPECT_EQ(actualFcm->getLogical(1, 1), (std::tuple(0, 2, 1))); + EXPECT_EQ(actualFcm->getLogical(1, 2), (std::tuple(0, 3, 0))); + EXPECT_EQ(actualFcm->getLogical(1, 0), (std::tuple(0, 3, 1))); } } diff --git a/arrus/core/devices/us4r/FrameChannelMappingImpl.cpp b/arrus/core/devices/us4r/FrameChannelMappingImpl.cpp index 645f374c9..5d8a55362 100644 --- a/arrus/core/devices/us4r/FrameChannelMappingImpl.cpp +++ b/arrus/core/devices/us4r/FrameChannelMappingImpl.cpp @@ -7,52 +7,57 @@ namespace arrus::devices { -FrameChannelMappingImpl::FrameChannelMappingImpl(FrameMapping &frameMapping, - ChannelMapping &channelMapping) - : frameMapping(std::move(frameMapping)), channelMapping(std::move(channelMapping)) { +FrameChannelMappingImpl::FrameChannelMappingImpl( + Us4OEMMapping &us4oemMapping, FrameMapping &frameMapping, ChannelMapping &channelMapping) + : us4oemMapping(std::move(us4oemMapping)), frameMapping(std::move(frameMapping)), + channelMapping(std::move(channelMapping)) { ARRUS_REQUIRES_TRUE_E(frameMapping.rows() == channelMapping.rows() - && frameMapping.cols() == channelMapping.cols(), - ArrusException("Frame and channel mapping arrays should have the " - "same shape")); + && frameMapping.cols() == channelMapping.cols() + && frameMapping.rows() == us4oemMapping.rows() + && frameMapping.cols() == us4oemMapping.cols(), + ArrusException("All channel mapping structures should have the same shape")); } -std::pair +std::tuple FrameChannelMappingImpl::getLogical(FrameNumber frame, ChannelIdx channel) { + auto us4oem = us4oemMapping(frame, channel); auto physicalFrame = frameMapping(frame, channel); auto physicalChannel = channelMapping(frame, channel); - return {physicalFrame, physicalChannel}; + return {us4oem, physicalFrame, physicalChannel}; } FrameChannelMapping::FrameNumber FrameChannelMappingImpl::getNumberOfLogicalFrames() { - assert(frameMapping.rows() >= 0 - && frameMapping.rows() <= std::numeric_limits::max()); + ARRUS_REQUIRES_TRUE(frameMapping.rows() >= 0 && frameMapping.rows() <= std::numeric_limits::max(), + "FCM number of logical frames exceeds the maximum number of frames (uint16::max)."); return static_cast(frameMapping.rows()); } ChannelIdx FrameChannelMappingImpl::getNumberOfLogicalChannels() { - assert(frameMapping.cols() >= 0 - && frameMapping.cols() <= std::numeric_limits::max()); + ARRUS_REQUIRES_TRUE(frameMapping.cols() >= 0 && frameMapping.cols() <= std::numeric_limits::max(), + "FCM number of logical channels exceeds the maximum number of channels (uint16::max)."); return static_cast(frameMapping.cols()); } FrameChannelMappingImpl::~FrameChannelMappingImpl() = default; -void -FrameChannelMappingBuilder::setChannelMapping(FrameNumber logicalFrame, ChannelIdx logicalChannel, - FrameNumber physicalFrame, int8 physicalChannel) { +void FrameChannelMappingBuilder::setChannelMapping(FrameNumber logicalFrame, ChannelIdx logicalChannel, + uint8 us4oem, FrameNumber physicalFrame, int8 physicalChannel) { + us4oemMapping(logicalFrame, logicalChannel) = us4oem; frameMapping(logicalFrame, logicalChannel) = physicalFrame; channelMapping(logicalFrame, logicalChannel) = physicalChannel; } FrameChannelMappingImpl::Handle FrameChannelMappingBuilder::build() { - return std::make_unique(this->frameMapping, this->channelMapping); + return std::make_unique(this->us4oemMapping, this->frameMapping, this->channelMapping); } FrameChannelMappingBuilder::FrameChannelMappingBuilder(FrameNumber nFrames, ChannelIdx nChannels) - : frameMapping(FrameChannelMappingImpl::FrameMapping(nFrames, nChannels)), + : us4oemMapping(FrameChannelMappingImpl::Us4OEMMapping(nFrames, nChannels)), + frameMapping(FrameChannelMappingImpl::FrameMapping(nFrames, nChannels)), channelMapping(FrameChannelMappingImpl::ChannelMapping(nFrames, nChannels)) { // Creates empty frame mapping. + us4oemMapping.fill(0); frameMapping.fill(0); channelMapping.fill(FrameChannelMapping::UNAVAILABLE); } diff --git a/arrus/core/devices/us4r/FrameChannelMappingImpl.h b/arrus/core/devices/us4r/FrameChannelMappingImpl.h index 3a1466d6b..2d1bd26f7 100644 --- a/arrus/core/devices/us4r/FrameChannelMappingImpl.h +++ b/arrus/core/devices/us4r/FrameChannelMappingImpl.h @@ -5,6 +5,8 @@ #include #include + +#include "arrus/core/api/devices/DeviceId.h" #include "arrus/core/api/devices/us4r/FrameChannelMapping.h" @@ -13,20 +15,16 @@ namespace arrus::devices { class FrameChannelMappingImpl : public FrameChannelMapping { public: using Handle = std::unique_ptr; + using Us4OEMMapping = Eigen::Matrix; using FrameMapping = Eigen::Matrix; using ChannelMapping = Eigen::Matrix; /** * Takes ownership for the provided frames. */ - FrameChannelMappingImpl(FrameMapping &frameMapping, ChannelMapping &channelMapping); + FrameChannelMappingImpl(Us4OEMMapping &us4oemMapping, FrameMapping &frameMapping, ChannelMapping &channelMapping); - /** - * @param frame logical frame to acquire - * @param channel channel in the logical frame to acquire - * @return frame and channel number of the physical signal data (the one returned by us4r device) - */ - std::pair getLogical(FrameNumber frame, ChannelIdx channel) override; + std::tuple getLogical(FrameNumber frame, ChannelIdx channel) override; FrameNumber getNumberOfLogicalFrames() override; @@ -35,7 +33,8 @@ class FrameChannelMappingImpl : public FrameChannelMapping { ~FrameChannelMappingImpl() override; private: - // logical (frame, number) -> physical (frame, number) + // logical (frame, number) -> physical (us4oem, frame, number) + Us4OEMMapping us4oemMapping; FrameMapping frameMapping; ChannelMapping channelMapping; }; @@ -43,16 +42,18 @@ class FrameChannelMappingImpl : public FrameChannelMapping { class FrameChannelMappingBuilder { public: using FrameNumber = FrameChannelMapping::FrameNumber; + using Us4OEMNumber = FrameChannelMapping::Us4OEMNumber; FrameChannelMappingBuilder(FrameNumber nFrames, ChannelIdx nChannels); - void setChannelMapping(FrameNumber logicalFrame, ChannelIdx logicalChannel, - FrameNumber physicalFrame, int8 physicalChannel); + void setChannelMapping(FrameNumber logicalFrame, ChannelIdx logicalChannel, // -> + Us4OEMNumber us4oem, FrameNumber physicalFrame, int8 physicalChannel); FrameChannelMappingImpl::Handle build(); private: // logical (frame, number) -> physical (frame, number) + FrameChannelMappingImpl::Us4OEMMapping us4oemMapping; FrameChannelMappingImpl::FrameMapping frameMapping; FrameChannelMappingImpl::ChannelMapping channelMapping; }; diff --git a/arrus/core/devices/us4r/probeadapter/ProbeAdapterImpl.cpp b/arrus/core/devices/us4r/probeadapter/ProbeAdapterImpl.cpp index fab31b571..90ea1e66f 100644 --- a/arrus/core/devices/us4r/probeadapter/ProbeAdapterImpl.cpp +++ b/arrus/core/devices/us4r/probeadapter/ProbeAdapterImpl.cpp @@ -69,9 +69,8 @@ ProbeAdapterImpl::setTxRxSequence(const std::vector &seq, uint16 batchSize, std::optional sri, bool triggerSync) { // Validate input sequence - ProbeAdapterTxRxValidator validator( - ::arrus::format("{} tx rx sequence", getDeviceId().toString()), - numberOfChannels); + ProbeAdapterTxRxValidator validator(::arrus::format("{} tx rx sequence", getDeviceId().toString()), + numberOfChannels); validator.validate(seq); validator.throwOnErrors(); @@ -106,8 +105,7 @@ ProbeAdapterImpl::setTxRxSequence(const std::vector &seq, txDelaysList.emplace(ordinal, std::vector>(seq.size())); } - // Split Tx, Rx apertures and tx delays into sub-apertures specific for - // each us4oem module. + // Split Tx, Rx apertures and tx delays into sub-apertures specific for each us4oem module. uint32 opNumber = 0; uint32 frameNumber = 0; @@ -118,9 +116,8 @@ ProbeAdapterImpl::setTxRxSequence(const std::vector &seq, const auto &rxAperture = op.getRxAperture(); const auto &txDelays = op.getTxDelays(); - // TODO change the below to an 'assert' ARRUS_REQUIRES_TRUE(txAperture.size() == rxAperture.size() && txAperture.size() == numberOfChannels, - format("Tx and Rx apertures should have a size: {}", numberOfChannels)); + format("TX and RX apertures should have a size: {}", numberOfChannels)); for(Ordinal ordinal = 0; ordinal < us4oems.size(); ++ordinal) { txApertures[ordinal][opNumber].resize(Us4OEMImpl::N_ADDR_CHANNELS); @@ -145,9 +142,8 @@ ProbeAdapterImpl::setTxRxSequence(const std::vector &seq, // FC Mapping stuff if(op.getRxAperture()[ach]) { isRxNop = false; - frameModule(frameNumber, activeAdapterCh + op.getRxPadding()[0]) = dstModule; - frameChannel(frameNumber, activeAdapterCh + op.getRxPadding()[0]) = - static_cast(activeUs4oemCh[dstModule]); + frameModule(frameNumber, activeAdapterCh+op.getRxPadding()[0]) = dstModule; + frameChannel(frameNumber, activeAdapterCh+op.getRxPadding()[0]) = (int32)(activeUs4oemCh[dstModule]); ++activeAdapterCh; ++activeUs4oemCh[dstModule]; } @@ -185,24 +181,18 @@ ProbeAdapterImpl::setTxRxSequence(const std::vector &seq, for(auto &us4oem: us4oems) { us4oemL2PChannelMappings.push_back(us4oem->getChannelMapping()); } - auto[splittedOps, opDstSplittedOp, opDestSplittedCh] = splitRxAperturesIfNecessary( - seqs, us4oemL2PChannelMappings); + auto[splittedOps, opDstSplittedOp, opDestSplittedCh] = splitRxAperturesIfNecessary(seqs, us4oemL2PChannelMappings); // set sequence on each us4oem std::vector fcMappings; - FrameChannelMapping::FrameNumber totalNumberOfFrames = 0; - std::vector frameOffsets(seqs.size(), 0); - // section -> us4oem -> transfer std::vector> outputTransfers; Us4RBufferBuilder us4RBufferBuilder; for(Ordinal us4oemOrdinal = 0; us4oemOrdinal < us4oems.size(); ++us4oemOrdinal) { auto &us4oem = us4oems[us4oemOrdinal]; - auto[buffer, fcMapping] = us4oem->setTxRxSequence( - splittedOps[us4oemOrdinal], tgcSamples, rxBufferSize, batchSize, sri, triggerSync); - frameOffsets[us4oemOrdinal] = totalNumberOfFrames; - totalNumberOfFrames += fcMapping->getNumberOfLogicalFrames(); + auto[buffer, fcMapping] = us4oem->setTxRxSequence(splittedOps[us4oemOrdinal], tgcSamples, rxBufferSize, + batchSize, sri, triggerSync); fcMappings.push_back(std::move(fcMapping)); // fcMapping is not valid anymore here us4RBufferBuilder.pushBack(buffer); @@ -227,24 +217,29 @@ ProbeAdapterImpl::setTxRxSequence(const std::vector &seq, // and has no assigned value. ARRUS_REQUIRES_DATA_TYPE_E(dstModuleChannel, int8,ArrusException("Invalid dstModuleChannel data type")); if(FrameChannelMapping::isChannelUnavailable((int8) dstModuleChannel)) { - outFcBuilder.setChannelMapping(frameIdx, activeRxChIdx + op.getRxPadding()[0], 0, - FrameChannelMapping::UNAVAILABLE); + outFcBuilder.setChannelMapping(frameIdx, activeRxChIdx + op.getRxPadding()[0], + 0, 0, FrameChannelMapping::UNAVAILABLE); } else { // Otherwise, we have an actual channel. ARRUS_REQUIRES_TRUE_E(dstModule >= 0 && dstModuleChannel >= 0, ArrusException("Dst module and dst channel should be non-negative")); + // dstOp, dstChannel - frame and channel after considering that the aperture ops are + // into multiple smaller ops for each us4oem separately. + // dstOp, dstChannel - frame and channel of a given module auto dstOp = opDstSplittedOp(dstModule, frameIdx, dstModuleChannel); auto dstChannel = opDestSplittedCh(dstModule, frameIdx, dstModuleChannel); - FrameChannelMapping::FrameNumber destFrame = 0; - int8 destFrameChannel = -1; + FrameChannelMapping::Us4OEMNumber us4oem = 0; + FrameChannelMapping::FrameNumber dstFrame = 0; + int8 dstFrameChannel = -1; if(!FrameChannelMapping::isChannelUnavailable(dstChannel)) { auto res = fcMappings[dstModule]->getLogical(dstOp, dstChannel); - destFrame = res.first; - destFrameChannel = res.second; + us4oem = std::get<0>(res); + dstFrame = std::get<1>(res); + dstFrameChannel = std::get<2>(res); } outFcBuilder.setChannelMapping(frameIdx, activeRxChIdx + op.getRxPadding()[0], - destFrame + frameOffsets[dstModule], destFrameChannel); + us4oem, dstFrame, dstFrameChannel); } ++activeRxChIdx; } @@ -317,7 +312,7 @@ void ProbeAdapterImpl::registerOutputBuffer(Us4ROutputBuffer *bufferDst, const U bool isTriggerRequired = workMode == Scheme::WorkMode::HOST; size_t nRepeats = nElementsDst/nElementsSrc; uint16 startFiring = 0; - for(int i = 0; i < bufferSrc.getNumberOfElements(); ++i) { + for(size_t i = 0; i < bufferSrc.getNumberOfElements(); ++i) { auto &srcElement = bufferSrc.getElement(i); uint16 endFiring = srcElement.getFiring(); for(size_t j = 0; j < nRepeats; ++j) { diff --git a/arrus/core/devices/us4r/probeadapter/ProbeAdapterImplTest.cpp b/arrus/core/devices/us4r/probeadapter/ProbeAdapterImplTest.cpp index 05906c7bd..6242f9aa7 100644 --- a/arrus/core/devices/us4r/probeadapter/ProbeAdapterImplTest.cpp +++ b/arrus/core/devices/us4r/probeadapter/ProbeAdapterImplTest.cpp @@ -65,14 +65,13 @@ std::vector getDefaultTxDelays(ChannelIdx nchannels) { return getNTimes(0.0f, nchannels); } -std::tuple< - Us4OEMBuffer, - FrameChannelMapping::Handle> -createEmptySetTxRxResult(FrameChannelMapping::FrameNumber nFrames, ChannelIdx nChannels) { +std::tuple +createEmptySetTxRxResult(FrameChannelMapping::Us4OEMNumber us4oem, FrameChannelMapping::FrameNumber nFrames, + ChannelIdx nChannels) { FrameChannelMappingBuilder builder(nFrames, nChannels); for(int i = 0; i < nFrames; ++i) { for(int j = 0; j < nChannels; ++j) { - builder.setChannelMapping(i, j, i, j); + builder.setChannelMapping(i, j, us4oem, i, j); } } Us4OEMBuffer buffer({Us4OEMBufferElement(0, 10, 0, arrus::Tuple({1, 1}), NdArray::DataType::INT16)}, {}); @@ -92,7 +91,7 @@ class MockUs4OEM : public Us4OEMImplBase { bool triggerSync), (override)); MOCK_METHOD(Interval, getAcceptedVoltageRange, (), (override)); - MOCK_METHOD(double, getSamplingFrequency, (), (override)); + MOCK_METHOD(float, getSamplingFrequency, (), (override)); MOCK_METHOD(void, startTrigger, (), (override)); MOCK_METHOD(void, stopTrigger, (), (override)); MOCK_METHOD(void, start, (), (override)); @@ -104,6 +103,7 @@ class MockUs4OEM : public Us4OEMImplBase { MOCK_METHOD(void, enableSequencer, (), (override)); MOCK_METHOD(std::vector, getChannelMapping, (), (override)); MOCK_METHOD(float, getFPGATemperature, (), (override)); + MOCK_METHOD(void, setTestPattern, (Us4OEMImpl::RxTestPattern), (override)); }; class AbstractProbeAdapterImplTest : public ::testing::Test { @@ -215,7 +215,7 @@ class ProbeAdapterChannelMapping1Test : public AbstractProbeAdapterImplTest { do { \ \ EXPECT_CALL(*(us4oems[deviceId].get()), setTxRxSequence(matcher, _, _, _, _, _)) \ - .WillOnce(Return(ByMove(createEmptySetTxRxResult(nFrames, 32)))); \ + .WillOnce(Return(ByMove(createEmptySetTxRxResult(deviceId, nFrames, 32)))); \ } while(0) #define EXPECT_SEQUENCE_PROPERTY(deviceId, matcher) \ @@ -508,7 +508,6 @@ TEST_F(ProbeAdapterChannelMappingEsaote3Test, DistributesAperturesCorrectlyForMu expectedRxAp01[18 + 32] = false; EXPECT_SEQUENCE_PROPERTY_NFRAMES( - 0, // Tx aperture should stay the same. // Rx aperture should be adjusted appropriately. @@ -699,16 +698,16 @@ TEST_F(ProbeAdapterChannelMappingEsaote3Test, ProducesCorrectFCMSingleDistribute FrameChannelMappingBuilder builder0(1, Us4OEMImpl::N_RX_CHANNELS); for(int i = 0; i < 32; ++i) { if(i < 24) { - builder0.setChannelMapping(0, i, 0, i); + builder0.setChannelMapping(0, i, 0, 0, i); } else { - builder0.setChannelMapping(0, i, 0, -1); + builder0.setChannelMapping(0, i, 0, 0, -1); } } auto fcm0 = builder0.build(); FrameChannelMappingBuilder builder1(1, Us4OEMImpl::N_RX_CHANNELS); for(int i = 0; i < 32; ++i) { - builder1.setChannelMapping(0, i, 0, i); + builder1.setChannelMapping(0, i, 1, 0, i); } auto fcm1 = builder1.build(); Us4OEMBuffer @@ -728,19 +727,22 @@ TEST_F(ProbeAdapterChannelMappingEsaote3Test, ProducesCorrectFCMSingleDistribute EXPECT_EQ(72 - 16, fcm->getNumberOfLogicalChannels()); for(int i = 0; i < 16; ++i) { - auto[frame, channel] = fcm->getLogical(0, i); + auto[us4oem, frame, channel] = fcm->getLogical(0, i); + EXPECT_EQ(0, us4oem); EXPECT_EQ(0, frame); EXPECT_EQ(channel, i); } for(int i = 16; i < 16 + 32; ++i) { - auto[frame, channel] = fcm->getLogical(0, i); - EXPECT_EQ(1, frame); + auto[us4oem, frame, channel] = fcm->getLogical(0, i); + EXPECT_EQ(1, us4oem); + EXPECT_EQ(0, frame); EXPECT_EQ(channel, i - 16); } for(int i = 16 + 32; i < 56; ++i) { - auto[frame, channel] = fcm->getLogical(0, i); + auto[us4oem, frame, channel] = fcm->getLogical(0, i); + EXPECT_EQ(0, us4oem); EXPECT_EQ(0, frame); EXPECT_EQ(channel, i - 32); } @@ -749,7 +751,7 @@ TEST_F(ProbeAdapterChannelMappingEsaote3Test, ProducesCorrectFCMSingleDistribute TEST_F(ProbeAdapterChannelMappingEsaote3Test, ProducesCorrectFCMSingleDistributedOperationWithGaps) { BitMask rxAperture(getNChannels(), false); ::arrus::setValuesInRange(rxAperture, 16, 73, true); - // Channels 20, 30 and 40 were masked for given us4oem and data is missing + // Channels 20, 30 and 40 were masked for given us4oem and data is missing. // Still, the input rx aperture stays as is. std::vector seq = { @@ -769,7 +771,7 @@ TEST_F(ProbeAdapterChannelMappingEsaote3Test, ProducesCorrectFCMSingleDistribute if(i != 20 - 16 && i != 30 - 16 && i <= 25) { currentJ = ++j; } - builder0.setChannelMapping(0, i, 0, currentJ); + builder0.setChannelMapping(0, i, 0, 0, currentJ); } auto fcm0 = builder0.build(); @@ -779,20 +781,18 @@ TEST_F(ProbeAdapterChannelMappingEsaote3Test, ProducesCorrectFCMSingleDistribute if(i != 40 - 32) { currentJ = ++j; } - builder1.setChannelMapping(0, i, 0, currentJ); + builder1.setChannelMapping(0, i, 1, 0, currentJ); } auto fcm1 = builder1.build(); - Us4OEMBuffer - us4oemBuffer({Us4OEMBufferElement(0, 10, 0, arrus::Tuple({1, 1}), NdArray::DataType::INT16)}, {}); + Us4OEMBuffer us4oemBuffer({Us4OEMBufferElement(0, 10, 0, arrus::Tuple({1, 1}), NdArray::DataType::INT16)}, + {}); std::tuple res0(us4oemBuffer, std::move(fcm0)); std::tuple res1(us4oemBuffer, std::move(fcm1)); - EXPECT_CALL(*(us4oems[0].get()), US4OEM_MOCK_SET_TX_RX_SEQUENCE()) - .WillOnce(Return(ByMove(std::move(res0)))); - EXPECT_CALL(*(us4oems[1].get()), US4OEM_MOCK_SET_TX_RX_SEQUENCE()) - .WillOnce(Return(ByMove(std::move(res1)))); + EXPECT_CALL(*(us4oems[0].get()), US4OEM_MOCK_SET_TX_RX_SEQUENCE()).WillOnce(Return(ByMove(std::move(res0)))); + EXPECT_CALL(*(us4oems[1].get()), US4OEM_MOCK_SET_TX_RX_SEQUENCE()).WillOnce(Return(ByMove(std::move(res1)))); auto[buffer, fcm] = SET_TX_RX_SEQUENCE(probeAdapter, seq); @@ -800,13 +800,17 @@ TEST_F(ProbeAdapterChannelMappingEsaote3Test, ProducesCorrectFCMSingleDistribute EXPECT_EQ(73 - 16, fcm->getNumberOfLogicalChannels()); std::vector expectedFrames; + std::vector expectedUs4oems; for(int i = 16; i < 32; ++i) { + expectedUs4oems.push_back(0); expectedFrames.push_back(0); } for(int i = 32; i < 64; ++i) { - expectedFrames.push_back(1); + expectedUs4oems.push_back(1); + expectedFrames.push_back(0); } for(int i = 64; i < 73; ++i) { + expectedUs4oems.push_back(0); expectedFrames.push_back(0); } std::vector expectedChannels = { @@ -818,7 +822,8 @@ TEST_F(ProbeAdapterChannelMappingEsaote3Test, ProducesCorrectFCMSingleDistribute }; for(int i = 0; i < 73 - 16; ++i) { - auto[frame, channel] = fcm->getLogical(0, i); + auto[us4oem, frame, channel] = fcm->getLogical(0, i); + EXPECT_EQ(expectedUs4oems[i], us4oem); EXPECT_EQ(expectedFrames[i], frame); EXPECT_EQ(expectedChannels[i], channel); } @@ -846,7 +851,7 @@ TEST_F(ProbeAdapterChannelMappingEsaote3Test, ProducesCorrectFCMForMultiOpRxAper FrameChannelMappingBuilder builder0(1, Us4OEMImpl::N_RX_CHANNELS); // The second op is Rx NOP. for(int i = 0; i < 32; ++i) { - builder0.setChannelMapping(0, i, 0, i); + builder0.setChannelMapping(0, i, 0, 0, i); } auto fcm0 = builder0.build(); @@ -856,54 +861,55 @@ TEST_F(ProbeAdapterChannelMappingEsaote3Test, ProducesCorrectFCMForMultiOpRxAper int currentJ = -1; if(i != 16 + 3) { currentJ = ++j; - builder1.setChannelMapping(0, i, 0, currentJ); + builder1.setChannelMapping(0, i, 1, 0, currentJ); } else { - builder1.setChannelMapping(0, i, 0, FrameChannelMapping::UNAVAILABLE); + builder1.setChannelMapping(0, i, 1, 0, FrameChannelMapping::UNAVAILABLE); } } // Second frame: for(int i = 0; i < 32; ++i) { if(i < 16) { - builder1.setChannelMapping(1, i, 1, i); + builder1.setChannelMapping(1, i, 1, 1, i); } else { - builder1.setChannelMapping(1, i, 1, FrameChannelMapping::UNAVAILABLE); + builder1.setChannelMapping(1, i, 1, 1, FrameChannelMapping::UNAVAILABLE); } } auto fcm1 = builder1.build(); - Us4OEMBuffer us4oemBuffer({ - Us4OEMBufferElement(0, 10, 0, arrus::Tuple({1, 1}), - ::arrus::framework::NdArray::DataType::INT16)}, {}); + Us4OEMBuffer us4oemBuffer({Us4OEMBufferElement(0, 10, 0, arrus::Tuple({1, 1}), NdArray::DataType::INT16)}, + {}); std::tuple res0(us4oemBuffer, std::move(fcm0)); std::tuple res1(us4oemBuffer, std::move(fcm1)); - EXPECT_CALL(*(us4oems[0].get()), US4OEM_MOCK_SET_TX_RX_SEQUENCE()) - .WillOnce(Return(ByMove(std::move(res0)))); - EXPECT_CALL(*(us4oems[1].get()), US4OEM_MOCK_SET_TX_RX_SEQUENCE()) - .WillOnce(Return(ByMove(std::move(res1)))); + EXPECT_CALL(*(us4oems[0].get()), US4OEM_MOCK_SET_TX_RX_SEQUENCE()).WillOnce(Return(ByMove(std::move(res0)))); + EXPECT_CALL(*(us4oems[1].get()), US4OEM_MOCK_SET_TX_RX_SEQUENCE()).WillOnce(Return(ByMove(std::move(res1)))); auto[buffer, fcm] = SET_TX_RX_SEQUENCE(probeAdapter, seq); EXPECT_EQ(1, fcm->getNumberOfLogicalFrames()); EXPECT_EQ(128 - 48, fcm->getNumberOfLogicalChannels()); + std::vector expectedUs4oems; std::vector expectedFrames; std::vector expectedChannels; - // Us4OEM:1, frame 1, channels 0-16 + // Us4OEM:1, frame 0, channels 0-16 for(int i = 48; i < 64; ++i) { - expectedFrames.push_back(1); + expectedUs4oems.push_back(1); + expectedFrames.push_back(0); expectedChannels.push_back(i - 48); } // Us4OEM:0 for(int i = 64; i < 96; ++i) { + expectedUs4oems.push_back(0); expectedFrames.push_back(0); expectedChannels.push_back(i - 64); } - // Us4OEM:1, frame 1, channels 16-32 + // Us4OEM:1, frame 0, channels 16-32 for(int i = 96; i < 96 + 15; ++i) { // 15 because there will be one -1 - expectedFrames.push_back(1); + expectedUs4oems.push_back(1); + expectedFrames.push_back(0); if(i == 99 && expectedChannels[expectedChannels.size() - 1] != FrameChannelMapping::UNAVAILABLE) { expectedChannels.push_back(FrameChannelMapping::UNAVAILABLE); --i; @@ -911,15 +917,17 @@ TEST_F(ProbeAdapterChannelMappingEsaote3Test, ProducesCorrectFCMForMultiOpRxAper expectedChannels.push_back(i - 96 + 16); } } - // Us4OEM:1, frame 2 + // Us4OEM:1, frame 1 for(int i = 96 + 16; i < 128; ++i) { - expectedFrames.push_back(2); + expectedUs4oems.push_back(1); + expectedFrames.push_back(1); expectedChannels.push_back(i - (96 + 16)); } // VALIDATE for(int i = 0; i < 128 - 48; ++i) { - auto[frame, channel] = fcm->getLogical(0, i); + auto[us4oem, frame, channel] = fcm->getLogical(0, i); + EXPECT_EQ(expectedUs4oems[i], us4oem); EXPECT_EQ(expectedFrames[i], frame); EXPECT_EQ(expectedChannels[i], channel); } @@ -943,9 +951,9 @@ TEST_F(ProbeAdapterChannelMappingEsaote3Test, AppliesPaddingToFCMCorrectly) { FrameChannelMappingBuilder builder0(1, Us4OEMImpl::N_RX_CHANNELS); for(int i = 0; i < 32; ++i) { if(i < 16) { - builder0.setChannelMapping(0, i, 0, i); + builder0.setChannelMapping(0, i, 0, 0, i); } else { - builder0.setChannelMapping(0, i, 0, -1); + builder0.setChannelMapping(0, i, 0, 0, -1); } } auto fcm0 = builder0.build(); @@ -959,10 +967,8 @@ TEST_F(ProbeAdapterChannelMappingEsaote3Test, AppliesPaddingToFCMCorrectly) { std::tuple res0(us4oemBuffer, std::move(fcm0)); std::tuple res1(us4oemBuffer, std::move(fcm1)); - EXPECT_CALL(*(us4oems[0].get()), US4OEM_MOCK_SET_TX_RX_SEQUENCE()) - .WillOnce(Return(ByMove(std::move(res0)))); - EXPECT_CALL(*(us4oems[1].get()), US4OEM_MOCK_SET_TX_RX_SEQUENCE()) - .WillOnce(Return(ByMove(std::move(res1)))); + EXPECT_CALL(*(us4oems[0].get()), US4OEM_MOCK_SET_TX_RX_SEQUENCE()).WillOnce(Return(ByMove(std::move(res0)))); + EXPECT_CALL(*(us4oems[1].get()), US4OEM_MOCK_SET_TX_RX_SEQUENCE()).WillOnce(Return(ByMove(std::move(res1)))); auto[buffer, fcm] = SET_TX_RX_SEQUENCE(probeAdapter, seq); @@ -970,13 +976,14 @@ TEST_F(ProbeAdapterChannelMappingEsaote3Test, AppliesPaddingToFCMCorrectly) { EXPECT_EQ(32, fcm->getNumberOfLogicalChannels()); // 16 active + 16 rx padding for(int i = 0; i < 16; ++i) { - auto[frame, channel] = fcm->getLogical(0, i); + auto[us4oem, frame, channel] = fcm->getLogical(0, i); ASSERT_EQ(0, frame); ASSERT_EQ(channel, FrameChannelMapping::UNAVAILABLE); } for(int i = 16; i < 32; ++i) { - auto[frame, channel] = fcm->getLogical(0, i); + auto[us4oem, frame, channel] = fcm->getLogical(0, i); + ASSERT_EQ(0, us4oem); ASSERT_EQ(0, frame); ASSERT_EQ(channel, i - 16); } @@ -999,16 +1006,16 @@ TEST_F(ProbeAdapterChannelMappingEsaote3Test, AppliesPaddingToFCMCorrectlyRxAper }; FrameChannelMappingBuilder builder0(1, Us4OEMImpl::N_RX_CHANNELS); for(int i = 0; i < 32; ++i) { - builder0.setChannelMapping(0, i, 0, i); + builder0.setChannelMapping(0, i, 0, 0, i); } auto fcm0 = builder0.build(); FrameChannelMappingBuilder builder1(1, Us4OEMImpl::N_RX_CHANNELS); for(int i = 0; i < 32; ++i) { if(i < 17) { - builder1.setChannelMapping(0, i, 0, i); + builder1.setChannelMapping(0, i, 1, 0, i); } else { - builder1.setChannelMapping(0, i, 0, FrameChannelMapping::UNAVAILABLE); + builder1.setChannelMapping(0, i, 1, 0, FrameChannelMapping::UNAVAILABLE); } } auto fcm1 = builder1.build(); @@ -1018,10 +1025,8 @@ TEST_F(ProbeAdapterChannelMappingEsaote3Test, AppliesPaddingToFCMCorrectlyRxAper std::tuple res0(us4oemBuffer, std::move(fcm0)); std::tuple res1(us4oemBuffer, std::move(fcm1)); - EXPECT_CALL(*(us4oems[0].get()), US4OEM_MOCK_SET_TX_RX_SEQUENCE()) - .WillOnce(Return(ByMove(std::move(res0)))); - EXPECT_CALL(*(us4oems[1].get()), US4OEM_MOCK_SET_TX_RX_SEQUENCE()) - .WillOnce(Return(ByMove(std::move(res1)))); + EXPECT_CALL(*(us4oems[0].get()), US4OEM_MOCK_SET_TX_RX_SEQUENCE()).WillOnce(Return(ByMove(std::move(res0)))); + EXPECT_CALL(*(us4oems[1].get()), US4OEM_MOCK_SET_TX_RX_SEQUENCE()).WillOnce(Return(ByMove(std::move(res1)))); auto[buffer, fcm] = SET_TX_RX_SEQUENCE(probeAdapter, seq); @@ -1029,17 +1034,19 @@ TEST_F(ProbeAdapterChannelMappingEsaote3Test, AppliesPaddingToFCMCorrectlyRxAper EXPECT_EQ(64, fcm->getNumberOfLogicalChannels()); // 49 active + 15 rx padding for(int i = 0; i < 15; ++i) { - auto[frame, channel] = fcm->getLogical(0, i); + auto[us4oem, frame, channel] = fcm->getLogical(0, i); ASSERT_EQ(channel, FrameChannelMapping::UNAVAILABLE); } for(int i = 15; i < 15 + 32; ++i) { - auto[frame, channel] = fcm->getLogical(0, i); + auto[us4oem, frame, channel] = fcm->getLogical(0, i); + ASSERT_EQ(0, us4oem); ASSERT_EQ(0, frame); ASSERT_EQ(channel, i - 15); } for(int i = 15 + 32; i < 64; ++i) { - auto[frame, channel] = fcm->getLogical(0, i); - ASSERT_EQ(1, frame); + auto[us4oem, frame, channel] = fcm->getLogical(0, i); + ASSERT_EQ(1, us4oem); + ASSERT_EQ(0, frame); ASSERT_EQ(channel, i - (15 + 32)); } } @@ -1065,9 +1072,9 @@ TEST_F(ProbeAdapterChannelMappingEsaote3Test, AppliesPaddingToFCMCorrectlyRightS FrameChannelMappingBuilder builder1(1, Us4OEMImpl::N_RX_CHANNELS); for(int i = 0; i < 32; ++i) { if(i < 16) { - builder1.setChannelMapping(0, i, 0, i); + builder1.setChannelMapping(0, i, 1, 0, i); } else { - builder1.setChannelMapping(0, i, 0, FrameChannelMapping::UNAVAILABLE); + builder1.setChannelMapping(0, i, 1, 0, FrameChannelMapping::UNAVAILABLE); } } auto fcm1 = builder1.build(); @@ -1078,10 +1085,8 @@ TEST_F(ProbeAdapterChannelMappingEsaote3Test, AppliesPaddingToFCMCorrectlyRightS std::tuple res0(us4oemBuffer, std::move(fcm0)); std::tuple res1(us4oemBuffer, std::move(fcm1)); - EXPECT_CALL(*(us4oems[0].get()), US4OEM_MOCK_SET_TX_RX_SEQUENCE()) - .WillOnce(Return(ByMove(std::move(res0)))); - EXPECT_CALL(*(us4oems[1].get()), US4OEM_MOCK_SET_TX_RX_SEQUENCE()) - .WillOnce(Return(ByMove(std::move(res1)))); + EXPECT_CALL(*(us4oems[0].get()), US4OEM_MOCK_SET_TX_RX_SEQUENCE()).WillOnce(Return(ByMove(std::move(res0)))); + EXPECT_CALL(*(us4oems[1].get()), US4OEM_MOCK_SET_TX_RX_SEQUENCE()).WillOnce(Return(ByMove(std::move(res1)))); auto[buffer, fcm] = SET_TX_RX_SEQUENCE(probeAdapter, seq); @@ -1089,12 +1094,13 @@ TEST_F(ProbeAdapterChannelMappingEsaote3Test, AppliesPaddingToFCMCorrectlyRightS EXPECT_EQ(32, fcm->getNumberOfLogicalChannels()); // 16 active + 16 rx padding for(int i = 0; i < 16; ++i) { - auto[frame, channel] = fcm->getLogical(0, i); + auto[us4oem, frame, channel] = fcm->getLogical(0, i); + ASSERT_EQ(1, us4oem); ASSERT_EQ(0, frame); ASSERT_EQ(channel, i); } for(int i = 16; i < 32; ++i) { - auto[frame, channel] = fcm->getLogical(0, i); + auto[us4oem, frame, channel] = fcm->getLogical(0, i); ASSERT_EQ(channel, FrameChannelMapping::UNAVAILABLE); } } diff --git a/arrus/core/devices/us4r/us4oem/Us4OEMImpl.cpp b/arrus/core/devices/us4r/us4oem/Us4OEMImpl.cpp index e722115e7..f3d0e4acb 100644 --- a/arrus/core/devices/us4r/us4oem/Us4OEMImpl.cpp +++ b/arrus/core/devices/us4r/us4oem/Us4OEMImpl.cpp @@ -402,9 +402,7 @@ Us4OEMImpl::setRxMappings(const std::vector &seq) { for(const auto isOn : op.getRxAperture()) { if(isOn) { isRxNop = false; - ARRUS_REQUIRES_TRUE_E( - onChannel < N_RX_CHANNELS, - ArrusException("Up to 32 active rx channels can be set.")); + ARRUS_REQUIRES_TRUE_E(onChannel < N_RX_CHANNELS, ArrusException("Up to 32 active rx channels can be set.")); // Physical channel number, values 0-31 auto rxChannel = channelMapping[channel]; @@ -427,7 +425,7 @@ Us4OEMImpl::setRxMappings(const std::vector &seq) { frameNumber = opId; } fcmBuilder.setChannelMapping(frameNumber, onChannel, - frameNumber, (int8) (mapping.size() - 1)); + getDeviceId().getOrdinal(), frameNumber, (int8) (mapping.size() - 1)); ++onChannel; } ++channel; @@ -463,10 +461,9 @@ Us4OEMImpl::setRxMappings(const std::vector &seq) { result.emplace(opId, rxMapId); // Set channel mapping ARRUS_REQUIRES_TRUE(rxMapping.size() == N_RX_CHANNELS, - arrus::format("Invalid size of the RX channel mapping to set: {}", rxMapping.size())); - ARRUS_REQUIRES_TRUE( - rxMapId < 128, - arrus::format("128 different rx mappings can be loaded only, deviceId: {}.", getDeviceId().toString())); + format("Invalid size of the RX channel mapping to set: {}", rxMapping.size())); + ARRUS_REQUIRES_TRUE(rxMapId < 128, + format("128 different rx mappings can be loaded only, deviceId: {}.", getDeviceId().toString())); ius4oem->SetRxChannelMapping(rxMapping, rxMapId); ++rxMapId; } else { diff --git a/arrus/core/devices/us4r/us4oem/Us4OEMImplTest.cpp b/arrus/core/devices/us4r/us4oem/Us4OEMImplTest.cpp index 794b1de1b..cf2eab9ff 100644 --- a/arrus/core/devices/us4r/us4oem/Us4OEMImplTest.cpp +++ b/arrus/core/devices/us4r/us4oem/Us4OEMImplTest.cpp @@ -573,7 +573,8 @@ TEST_F(Us4OEMImplEsaote3LikeTest, TestFrameChannelMappingForNonconflictingRxMapp EXPECT_EQ(fcm->getNumberOfLogicalFrames(), 1); for(size_t i = 0; i < Us4OEMImpl::N_RX_CHANNELS; ++i) { - auto[dstFrame, dstChannel] = fcm->getLogical(0, i); + auto[us4oem, dstFrame, dstChannel] = fcm->getLogical(0, i); + EXPECT_EQ(us4oem, 0); EXPECT_EQ(dstChannel, i); EXPECT_EQ(dstFrame, 0); } @@ -594,7 +595,8 @@ TEST_F(Us4OEMImplEsaote3LikeTest, TestFrameChannelMappingForNonconflictingRxMapp EXPECT_EQ(fcm->getNumberOfLogicalFrames(), 1); for(size_t i = 0; i < Us4OEMImpl::N_RX_CHANNELS; ++i) { - auto[dstFrame, dstChannel] = fcm->getLogical(0, i); + auto[us4oem, dstFrame, dstChannel] = fcm->getLogical(0, i); + EXPECT_EQ(us4oem, 0); EXPECT_EQ(dstChannel, i); EXPECT_EQ(dstFrame, 0); } @@ -617,7 +619,8 @@ TEST_F(Us4OEMImplEsaote3LikeTest, TestFrameChannelMappingIncompleteRxAperture) { EXPECT_EQ(fcm->getNumberOfLogicalFrames(), 1); for(size_t i = 0; i < 30; ++i) { - auto[dstFrame, dstChannel] = fcm->getLogical(0, i); + auto[us4oem, dstFrame, dstChannel] = fcm->getLogical(0, i); + EXPECT_EQ(us4oem, 0); EXPECT_EQ(dstChannel, i); EXPECT_EQ(dstFrame, 0); } @@ -640,7 +643,7 @@ TEST_F(Us4OEMImplConflictingChannelsTest, TestFrameChannelMappingForConflictingM auto [buffer, fcm] = SET_TX_RX_SEQUENCE(us4oem, seq); for(size_t i = 0; i < Us4OEMImpl::N_RX_CHANNELS; ++i) { - auto[dstfr, dstch] = fcm->getLogical(0, i); + auto[us4oem, dstfr, dstch] = fcm->getLogical(0, i); std::cerr << (int16) dstch << ", "; } std::cerr << std::endl; @@ -655,7 +658,8 @@ TEST_F(Us4OEMImplConflictingChannelsTest, TestFrameChannelMappingForConflictingM }; for(size_t i = 0; i < Us4OEMImpl::N_RX_CHANNELS; ++i) { - auto[dstFrame, dstChannel] = fcm->getLogical(0, i); + auto[us4oem, dstFrame, dstChannel] = fcm->getLogical(0, i); + EXPECT_EQ(us4oem, 0); EXPECT_EQ(dstChannel, expectedDstChannels[i]); EXPECT_EQ(dstFrame, 0); } @@ -786,7 +790,8 @@ TEST_F(Us4OEMImplEsaote3ChannelsMaskTest, MasksProperlyASingleChannel) { expectedSrcChannels[3] = 3; for(int i = 0; i < Us4OEMImpl::N_RX_CHANNELS; ++i) { - auto[srcFrame, srcChannel] = fcm->getLogical(0, i); + auto[us4oem, srcFrame, srcChannel] = fcm->getLogical(0, i); + EXPECT_EQ(us4oem, 0); EXPECT_EQ(srcFrame, 0); ASSERT_EQ(srcChannel, expectedSrcChannels[i]); } @@ -916,7 +921,8 @@ TEST_F(Us4OEMImplEsaote3ChannelsMaskTest, MasksProperlyASingleChannelForAllOpera expectedSrcChannels[3] = 3; for(int i = 0; i < Us4OEMImpl::N_RX_CHANNELS; ++i) { - auto [srcFrame, srcChannel] = fcm->getLogical(0, i); + auto [us4oem, srcFrame, srcChannel] = fcm->getLogical(0, i); + EXPECT_EQ(us4oem, 0); EXPECT_EQ(srcFrame, 0); ASSERT_EQ(srcChannel, expectedSrcChannels[i]); } @@ -928,8 +934,9 @@ TEST_F(Us4OEMImplEsaote3ChannelsMaskTest, MasksProperlyASingleChannelForAllOpera ChannelIdx rxChannelNumber = 0; for(auto bit : rxApertures[frame]) { if(bit) { - auto [srcFrame, srcChannel] = fcm->getLogical(frame, i); + auto [us4oem, srcFrame, srcChannel] = fcm->getLogical(frame, i); std::cerr << frame << ", " << (int)i << ", " << srcFrame << ", " << (int)srcChannel << std::endl; + ASSERT_EQ(us4oem, 0); ASSERT_EQ(srcFrame, frame); ASSERT_EQ(srcChannel, i++); } From 360d9828dacd6cd8b7d47ddac3c921e9137e7885 Mon Sep 17 00:00:00 2001 From: Piotr Jarosik Date: Mon, 1 Nov 2021 09:29:46 +0100 Subject: [PATCH 2/8] From now on FCM gives information about the TARGET US4OEM, channel and frame. --- api/python/arrus/devices/us4r.py | 5 +++++ api/python/arrus/session.py | 9 +++++++-- api/python/arrus/utils/core.py | 15 +++++++++++---- api/python/arrus/utils/imaging.py | 5 +++++ api/python/wrappers/core.i | 2 +- .../api/devices/us4r/FrameChannelMapping.h | 12 ++++++++++++ .../devices/us4r/FrameChannelMappingImpl.cpp | 18 ++++++++++++++---- .../devices/us4r/FrameChannelMappingImpl.h | 8 +++++++- .../us4r/probeadapter/ProbeAdapterImpl.cpp | 5 +++++ .../us4r/probeadapter/ProbeAdapterImplTest.cpp | 15 +++++++++++++++ 10 files changed, 82 insertions(+), 12 deletions(-) diff --git a/api/python/arrus/devices/us4r.py b/api/python/arrus/devices/us4r.py index 0e82f36a9..6d169eacc 100644 --- a/api/python/arrus/devices/us4r.py +++ b/api/python/arrus/devices/us4r.py @@ -28,9 +28,14 @@ class FrameChannelMapping: :param frames: a mapping: (logical frame, logical channel) -> physical frame :param channels: a mapping: (logical frame, logical channel) -> physical channel + :param us4oems: a mapping: (logical frame, logical channel) -> us4OEM number + :param frame_offsets: frame starting number for each us4OEM available in the system + :param batch_size: number of sequences in a single batch """ frames: np.ndarray channels: np.ndarray + us4oems: np.ndarray + frame_offsets: np.ndarray batch_size: int = 1 diff --git a/api/python/arrus/session.py b/api/python/arrus/session.py index 193373b9b..ff669ef95 100644 --- a/api/python/arrus/session.py +++ b/api/python/arrus/session.py @@ -105,9 +105,14 @@ def upload(self, scheme: arrus.ops.us4r.Scheme): ### # -- Constant metadata # --- FCM - fcm_frame, fcm_channel = arrus.utils.core.convert_fcm_to_np_arrays(fcm) + fcm_us4oems, fcm_frame, fcm_channel, frame_offsets = \ + arrus.utils.core.convert_fcm_to_np_arrays(fcm, us_device.n_us4oems) fcm = arrus.devices.us4r.FrameChannelMapping( - frames=fcm_frame, channels=fcm_channel, batch_size=1) + us4oems=fcm_us4oems, + frames=fcm_frame, + channels=fcm_channel, + frame_offsets=frame_offsets, + batch_size=1) # --- Frame acquisition context fac = self._create_frame_acquisition_context(seq, raw_seq, us_device_dto, medium) diff --git a/api/python/arrus/utils/core.py b/api/python/arrus/utils/core.py index 3ee616a10..c7b4b5006 100644 --- a/api/python/arrus/utils/core.py +++ b/api/python/arrus/utils/core.py @@ -62,27 +62,34 @@ def convert_to_core_sequence(seq): return core_seq -def convert_fcm_to_np_arrays(fcm): +def convert_fcm_to_np_arrays(fcm, n_us4oems): """ Converts frame channel mapping to a tupple of numpy arrays. :param fcm: arrus.core.FrameChannelMapping :return: a pair of numpy arrays: fcm_frame, fcm_channel """ + fcm_us4oem = np.zeros( + (fcm.getNumberOfLogicalFrames(), fcm.getNumberOfLogicalChannels()), + dtype=np.uint8) fcm_frame = np.zeros( (fcm.getNumberOfLogicalFrames(), fcm.getNumberOfLogicalChannels()), dtype=np.int16) fcm_channel = np.zeros( (fcm.getNumberOfLogicalFrames(), fcm.getNumberOfLogicalChannels()), dtype=np.int8) + frame_offsets = np.zeros(n_us4oems, dtype=np.uint32) for frame in range(fcm.getNumberOfLogicalFrames()): for channel in range(fcm.getNumberOfLogicalChannels()): frame_channel = fcm.getLogical(frame, channel) - src_frame = frame_channel[0] - src_channel = frame_channel[1] + src_us4oem = frame_channel[0] + src_frame = frame_channel[1] + src_channel = frame_channel[2] + fcm_us4oem[frame, channel] = src_us4oem fcm_frame[frame, channel] = src_frame fcm_channel[frame, channel] = src_channel - return fcm_frame, fcm_channel + frame_offsets = [fcm.getFirstFrame(i) for i in range(n_us4oems)] + return fcm_us4oem, fcm_frame, fcm_channel, frame_offsets def convert_to_py_probe_model(core_model): diff --git a/api/python/arrus/utils/imaging.py b/api/python/arrus/utils/imaging.py index 5cde98fa1..554693d6f 100644 --- a/api/python/arrus/utils/imaging.py +++ b/api/python/arrus/utils/imaging.py @@ -1733,6 +1733,11 @@ def prepare(self, const_metadata: arrus.metadata.ConstMetadata): self._fcm_frames = cp.asarray(fcm.frames) self._fcm_channels = cp.asarray(fcm.channels) self._fcm_us4oems = cp.asarray(fcm.us4oems) + # 32 - number of us4OEM rx channels + frame_offsets = [fcm.frame_offsets*n_samples*32 + for us4oem in range(us4r.n_us4oems)] + self._frame_offsets = cp.asarray(frame_offsets) + self.grid_size, self.block_size = get_default_grid_block_size( self._fcm_frames, n_samples, batch_size diff --git a/api/python/wrappers/core.i b/api/python/wrappers/core.i index 4ec8ee085..f5a64238c 100644 --- a/api/python/wrappers/core.i +++ b/api/python/wrappers/core.i @@ -146,7 +146,7 @@ using namespace arrus::devices; %shared_ptr(arrus::framework::DataBuffer); namespace std { - %template(FrameChannelMappingElement) pair; + %template(FrameChannelMappingElement) tuple; }; namespace arrus { %template(TupleUint32) Tuple; diff --git a/arrus/core/api/devices/us4r/FrameChannelMapping.h b/arrus/core/api/devices/us4r/FrameChannelMapping.h index dfbc8d24c..4933213c2 100644 --- a/arrus/core/api/devices/us4r/FrameChannelMapping.h +++ b/arrus/core/api/devices/us4r/FrameChannelMapping.h @@ -30,6 +30,18 @@ class FrameChannelMapping { */ virtual std::tuple getLogical(FrameNumber frame, ChannelIdx channel) = 0; + /** + * Returns the number of frame where the given us4OEM data starts. + * The frame number is computed taking into account the batch size and the number of frames in the + * sequence of data produced by preceding us4OEM modules. That is, assuming the same number of samples + * is acquired in each RF frame, you can get the address where us4oem data starts using the following + * formula: the frame number * number of samples * 32 (number of us4OEM RX channels). + * + * @param us4oem us4oem ordinal number (0, 1, ...) + * @return the number of frame, which starts portion of data acquired by the given us4OEM. + */ + virtual arrus::uint32 getFirstFrame(arrus::uint8 us4oem) = 0; + virtual FrameNumber getNumberOfLogicalFrames() = 0; virtual ChannelIdx getNumberOfLogicalChannels() = 0; diff --git a/arrus/core/devices/us4r/FrameChannelMappingImpl.cpp b/arrus/core/devices/us4r/FrameChannelMappingImpl.cpp index 5d8a55362..7e1778e4d 100644 --- a/arrus/core/devices/us4r/FrameChannelMappingImpl.cpp +++ b/arrus/core/devices/us4r/FrameChannelMappingImpl.cpp @@ -8,9 +8,10 @@ namespace arrus::devices { FrameChannelMappingImpl::FrameChannelMappingImpl( - Us4OEMMapping &us4oemMapping, FrameMapping &frameMapping, ChannelMapping &channelMapping) + Us4OEMMapping &us4oemMapping, FrameMapping &frameMapping, ChannelMapping &channelMapping, + std::vector frameOffsets) : us4oemMapping(std::move(us4oemMapping)), frameMapping(std::move(frameMapping)), - channelMapping(std::move(channelMapping)) { + channelMapping(std::move(channelMapping)), frameOffsets(std::move(frameOffsets)) { ARRUS_REQUIRES_TRUE_E(frameMapping.rows() == channelMapping.rows() && frameMapping.cols() == channelMapping.cols() @@ -19,6 +20,8 @@ FrameChannelMappingImpl::FrameChannelMappingImpl( ArrusException("All channel mapping structures should have the same shape")); } +FrameChannelMappingImpl::~FrameChannelMappingImpl() = default; + std::tuple FrameChannelMappingImpl::getLogical(FrameNumber frame, ChannelIdx channel) { auto us4oem = us4oemMapping(frame, channel); @@ -39,7 +42,9 @@ ChannelIdx FrameChannelMappingImpl::getNumberOfLogicalChannels() { return static_cast(frameMapping.cols()); } -FrameChannelMappingImpl::~FrameChannelMappingImpl() = default; +uint32 FrameChannelMappingImpl::getFirstFrame(uint8 us4oem) { + return frameOffsets[us4oem]; +} void FrameChannelMappingBuilder::setChannelMapping(FrameNumber logicalFrame, ChannelIdx logicalChannel, uint8 us4oem, FrameNumber physicalFrame, int8 physicalChannel) { @@ -49,7 +54,8 @@ void FrameChannelMappingBuilder::setChannelMapping(FrameNumber logicalFrame, Cha } FrameChannelMappingImpl::Handle FrameChannelMappingBuilder::build() { - return std::make_unique(this->us4oemMapping, this->frameMapping, this->channelMapping); + return std::make_unique(this->us4oemMapping, this->frameMapping, this->channelMapping, + this->frameOffsets); } FrameChannelMappingBuilder::FrameChannelMappingBuilder(FrameNumber nFrames, ChannelIdx nChannels) @@ -62,5 +68,9 @@ FrameChannelMappingBuilder::FrameChannelMappingBuilder(FrameNumber nFrames, Chan channelMapping.fill(FrameChannelMapping::UNAVAILABLE); } +void FrameChannelMappingBuilder::setFrameOffsets(const std::vector &offsets) { + FrameChannelMappingBuilder::frameOffsets = offsets; +} + } diff --git a/arrus/core/devices/us4r/FrameChannelMappingImpl.h b/arrus/core/devices/us4r/FrameChannelMappingImpl.h index 2d1bd26f7..b5f6a6dbc 100644 --- a/arrus/core/devices/us4r/FrameChannelMappingImpl.h +++ b/arrus/core/devices/us4r/FrameChannelMappingImpl.h @@ -22,10 +22,13 @@ class FrameChannelMappingImpl : public FrameChannelMapping { /** * Takes ownership for the provided frames. */ - FrameChannelMappingImpl(Us4OEMMapping &us4oemMapping, FrameMapping &frameMapping, ChannelMapping &channelMapping); + FrameChannelMappingImpl(Us4OEMMapping &us4oemMapping, FrameMapping &frameMapping, ChannelMapping &channelMapping, + std::vector frameOffsets = {0}); std::tuple getLogical(FrameNumber frame, ChannelIdx channel) override; + uint32 getFirstFrame(uint8 us4oem) override; + FrameNumber getNumberOfLogicalFrames() override; ChannelIdx getNumberOfLogicalChannels() override; @@ -37,6 +40,7 @@ class FrameChannelMappingImpl : public FrameChannelMapping { Us4OEMMapping us4oemMapping; FrameMapping frameMapping; ChannelMapping channelMapping; + std::vector frameOffsets; }; class FrameChannelMappingBuilder { @@ -50,12 +54,14 @@ class FrameChannelMappingBuilder { Us4OEMNumber us4oem, FrameNumber physicalFrame, int8 physicalChannel); FrameChannelMappingImpl::Handle build(); + void setFrameOffsets(const std::vector &frameOffsets); private: // logical (frame, number) -> physical (frame, number) FrameChannelMappingImpl::Us4OEMMapping us4oemMapping; FrameChannelMappingImpl::FrameMapping frameMapping; FrameChannelMappingImpl::ChannelMapping channelMapping; + std::vector frameOffsets = {0}; }; } diff --git a/arrus/core/devices/us4r/probeadapter/ProbeAdapterImpl.cpp b/arrus/core/devices/us4r/probeadapter/ProbeAdapterImpl.cpp index 90ea1e66f..ade945b32 100644 --- a/arrus/core/devices/us4r/probeadapter/ProbeAdapterImpl.cpp +++ b/arrus/core/devices/us4r/probeadapter/ProbeAdapterImpl.cpp @@ -187,12 +187,16 @@ ProbeAdapterImpl::setTxRxSequence(const std::vector &seq, std::vector fcMappings; // section -> us4oem -> transfer std::vector> outputTransfers; + uint32 currentFrameOffset = 0; + std::vector frameOffsets{static_cast(us4oems.size()), 0}; Us4RBufferBuilder us4RBufferBuilder; for(Ordinal us4oemOrdinal = 0; us4oemOrdinal < us4oems.size(); ++us4oemOrdinal) { auto &us4oem = us4oems[us4oemOrdinal]; auto[buffer, fcMapping] = us4oem->setTxRxSequence(splittedOps[us4oemOrdinal], tgcSamples, rxBufferSize, batchSize, sri, triggerSync); + frameOffsets[us4oemOrdinal] = currentFrameOffset; + currentFrameOffset += fcMapping->getNumberOfLogicalFrames()*batchSize; fcMappings.push_back(std::move(fcMapping)); // fcMapping is not valid anymore here us4RBufferBuilder.pushBack(buffer); @@ -246,6 +250,7 @@ ProbeAdapterImpl::setTxRxSequence(const std::vector &seq, } ++frameIdx; } + outFcBuilder.setFrameOffsets(frameOffsets); return {us4RBufferBuilder.build(), outFcBuilder.build()}; } diff --git a/arrus/core/devices/us4r/probeadapter/ProbeAdapterImplTest.cpp b/arrus/core/devices/us4r/probeadapter/ProbeAdapterImplTest.cpp index 6242f9aa7..8ab915ed4 100644 --- a/arrus/core/devices/us4r/probeadapter/ProbeAdapterImplTest.cpp +++ b/arrus/core/devices/us4r/probeadapter/ProbeAdapterImplTest.cpp @@ -746,6 +746,10 @@ TEST_F(ProbeAdapterChannelMappingEsaote3Test, ProducesCorrectFCMSingleDistribute EXPECT_EQ(0, frame); EXPECT_EQ(channel, i - 32); } + + // Make sure the correct frame offsets are set. + EXPECT_EQ(0, fcm->getFirstFrame(0)); // Us4OEM:0 + EXPECT_EQ(1, fcm->getFirstFrame(1)); // Us4OEM:1 } TEST_F(ProbeAdapterChannelMappingEsaote3Test, ProducesCorrectFCMSingleDistributedOperationWithGaps) { @@ -827,6 +831,9 @@ TEST_F(ProbeAdapterChannelMappingEsaote3Test, ProducesCorrectFCMSingleDistribute EXPECT_EQ(expectedFrames[i], frame); EXPECT_EQ(expectedChannels[i], channel); } + // Make sure the correct frame offsets are set. + EXPECT_EQ(0, fcm->getFirstFrame(0)); // Us4OEM:0 + EXPECT_EQ(1, fcm->getFirstFrame(1)); // Us4OEM:1 } TEST_F(ProbeAdapterChannelMappingEsaote3Test, ProducesCorrectFCMForMultiOpRxAperture) { @@ -931,6 +938,9 @@ TEST_F(ProbeAdapterChannelMappingEsaote3Test, ProducesCorrectFCMForMultiOpRxAper EXPECT_EQ(expectedFrames[i], frame); EXPECT_EQ(expectedChannels[i], channel); } + // Make sure the correct frame offsets are set. + EXPECT_EQ(0, fcm->getFirstFrame(0)); // Us4OEM:0 + EXPECT_EQ(1, fcm->getFirstFrame(1)); // Us4OEM:1 } // Currently padding impacts the output frame channel mapping @@ -987,6 +997,8 @@ TEST_F(ProbeAdapterChannelMappingEsaote3Test, AppliesPaddingToFCMCorrectly) { ASSERT_EQ(0, frame); ASSERT_EQ(channel, i - 16); } + // Make sure the correct frame offsets are set. + EXPECT_EQ(0, fcm->getFirstFrame(0)); // Us4OEM:0 } // The same as above, but with aperture using two modules @@ -1049,6 +1061,8 @@ TEST_F(ProbeAdapterChannelMappingEsaote3Test, AppliesPaddingToFCMCorrectlyRxAper ASSERT_EQ(0, frame); ASSERT_EQ(channel, i - (15 + 32)); } + EXPECT_EQ(0, fcm->getFirstFrame(0)); // Us4OEM:0 + EXPECT_EQ(1, fcm->getFirstFrame(1)); // Us4OEM:1 } TEST_F(ProbeAdapterChannelMappingEsaote3Test, AppliesPaddingToFCMCorrectlyRightSide) { @@ -1103,6 +1117,7 @@ TEST_F(ProbeAdapterChannelMappingEsaote3Test, AppliesPaddingToFCMCorrectlyRightS auto[us4oem, frame, channel] = fcm->getLogical(0, i); ASSERT_EQ(channel, FrameChannelMapping::UNAVAILABLE); } + EXPECT_EQ(0, fcm->getFirstFrame(1)); // Us4OEM:1 } // ------------------------------------------ TODO Test that all other parameters are passed unmodified } From b73c6c1d412905528f8edb8a6b9c019a97abad88 Mon Sep 17 00:00:00 2001 From: Piotr Jarosik Date: Mon, 1 Nov 2021 11:17:08 +0100 Subject: [PATCH 3/8] Wraped the output of the getLogical function into a class FrameChannelMappingAddress. --- api/python/arrus/utils/core.py | 9 ++-- api/python/arrus/utils/imaging.py | 10 ++--- api/python/wrappers/core.i | 3 -- .../api/devices/us4r/FrameChannelMapping.h | 41 ++++++++++++++++++- arrus/core/devices/probe/ProbeImplTest.cpp | 33 +++++++-------- .../devices/us4r/FrameChannelMappingImpl.cpp | 4 +- .../devices/us4r/FrameChannelMappingImpl.h | 34 ++++++++++++++- arrus/core/devices/us4r/us4oem/Us4OEMImpl.cpp | 3 +- .../devices/us4r/us4oem/Us4OEMImplTest.cpp | 1 + 9 files changed, 105 insertions(+), 33 deletions(-) diff --git a/api/python/arrus/utils/core.py b/api/python/arrus/utils/core.py index c7b4b5006..4b683d77a 100644 --- a/api/python/arrus/utils/core.py +++ b/api/python/arrus/utils/core.py @@ -82,13 +82,14 @@ def convert_fcm_to_np_arrays(fcm, n_us4oems): for frame in range(fcm.getNumberOfLogicalFrames()): for channel in range(fcm.getNumberOfLogicalChannels()): frame_channel = fcm.getLogical(frame, channel) - src_us4oem = frame_channel[0] - src_frame = frame_channel[1] - src_channel = frame_channel[2] + src_us4oem = frame_channel.getUs4oem() + src_frame = frame_channel.getFrame() + src_channel = frame_channel.getChannel() fcm_us4oem[frame, channel] = src_us4oem fcm_frame[frame, channel] = src_frame fcm_channel[frame, channel] = src_channel - frame_offsets = [fcm.getFirstFrame(i) for i in range(n_us4oems)] + frame_offsets = [fcm.getStartFrame(i) for i in range(n_us4oems)] + frame_offsets = np.array(frame_offsets) return fcm_us4oem, fcm_frame, fcm_channel, frame_offsets diff --git a/api/python/arrus/utils/imaging.py b/api/python/arrus/utils/imaging.py index 554693d6f..0e026eafd 100644 --- a/api/python/arrus/utils/imaging.py +++ b/api/python/arrus/utils/imaging.py @@ -1733,11 +1733,10 @@ def prepare(self, const_metadata: arrus.metadata.ConstMetadata): self._fcm_frames = cp.asarray(fcm.frames) self._fcm_channels = cp.asarray(fcm.channels) self._fcm_us4oems = cp.asarray(fcm.us4oems) - # 32 - number of us4OEM rx channels - frame_offsets = [fcm.frame_offsets*n_samples*32 - for us4oem in range(us4r.n_us4oems)] + # 32 - number of us4OEM rx channels, 2 - number of bytes per sample + frame_offsets = fcm.frame_offsets*n_samples*32*2 + # TODO constant memory self._frame_offsets = cp.asarray(frame_offsets) - self.grid_size, self.block_size = get_default_grid_block_size( self._fcm_frames, n_samples, batch_size @@ -1746,7 +1745,8 @@ def prepare(self, const_metadata: arrus.metadata.ConstMetadata): def gpu_remap_fn(data): run_remap(self.grid_size, self.block_size, [self._output_buffer, data, - self._fcm_frames, self._fcm_channels, + self._fcm_frames, self._fcm_channels, self._fcm_us4oems, + self._frame_offsets, n_frames, n_samples, n_channels]) self._remap_fn = gpu_remap_fn diff --git a/api/python/wrappers/core.i b/api/python/wrappers/core.i index f5a64238c..88c972ff9 100644 --- a/api/python/wrappers/core.i +++ b/api/python/wrappers/core.i @@ -145,9 +145,6 @@ using namespace arrus::devices; %shared_ptr(arrus::framework::BufferElement); %shared_ptr(arrus::framework::DataBuffer); -namespace std { - %template(FrameChannelMappingElement) tuple; -}; namespace arrus { %template(TupleUint32) Tuple; }; diff --git a/arrus/core/api/devices/us4r/FrameChannelMapping.h b/arrus/core/api/devices/us4r/FrameChannelMapping.h index 4933213c2..6b320050f 100644 --- a/arrus/core/api/devices/us4r/FrameChannelMapping.h +++ b/arrus/core/api/devices/us4r/FrameChannelMapping.h @@ -7,6 +7,45 @@ namespace arrus::devices { + +/** + * A tuple that describes position of a given channel in a sequence + * of frames produced by us4OEM modules. + */ +class FrameChannelMappingAddress { + +public: + FrameChannelMappingAddress(uint8 us4oem, unsigned short frame, int8 channel) + : us4oem(us4oem), frame(frame), channel(channel) {} + + uint8 getUs4oem() const { + return us4oem; + } + + unsigned short getFrame() const { + return frame; + } + + int8 getChannel() const { + return channel; + } + + bool operator==(const FrameChannelMappingAddress &rhs) const { + return us4oem == rhs.us4oem && + frame == rhs.frame && + channel == rhs.channel; + } + + bool operator!=(const FrameChannelMappingAddress &rhs) const { + return !(rhs == *this); + } + +private: + arrus::uint8 us4oem; + unsigned short frame; + arrus::int8 channel; +}; + /** * Frame channel mapping: logical (frame, channel) -> physical (frame, channel) */ @@ -28,7 +67,7 @@ class FrameChannelMapping { * @param channel logical channel number * @return a tuple: us4oem module number, frame number (within a single sequence), channel number */ - virtual std::tuple getLogical(FrameNumber frame, ChannelIdx channel) = 0; + virtual FrameChannelMappingAddress getLogical(FrameNumber frame, ChannelIdx channel) = 0; /** * Returns the number of frame where the given us4OEM data starts. diff --git a/arrus/core/devices/probe/ProbeImplTest.cpp b/arrus/core/devices/probe/ProbeImplTest.cpp index 97bff84c5..c7026831f 100644 --- a/arrus/core/devices/probe/ProbeImplTest.cpp +++ b/arrus/core/devices/probe/ProbeImplTest.cpp @@ -7,6 +7,7 @@ namespace { using namespace arrus; using namespace arrus::devices; +using ::arrus::devices::FrameChannelMappingAddress; class ProbeImplFcmRemapTest : public ::testing::Test { protected: @@ -54,15 +55,15 @@ TEST_F(ProbeImplFcmRemapTest, OneToOne) { EXPECT_EQ(actualNFrames, N_FRAMES); EXPECT_EQ(actualNChannels, N_CHANNELS); - EXPECT_EQ(actualFcm->getLogical(0, 0), (std::tuple(0, 0, 0))); - EXPECT_EQ(actualFcm->getLogical(0, 1), (std::tuple(0, 0, 1))); - EXPECT_EQ(actualFcm->getLogical(0, 2), (std::tuple(0, 1, 0))); - EXPECT_EQ(actualFcm->getLogical(0, 3), (std::tuple(0, 1, 1))); + EXPECT_EQ(actualFcm->getLogical(0, 0), (FrameChannelMappingAddress(0, 0, 0))); + EXPECT_EQ(actualFcm->getLogical(0, 1), (FrameChannelMappingAddress(0, 0, 1))); + EXPECT_EQ(actualFcm->getLogical(0, 2), (FrameChannelMappingAddress(0, 1, 0))); + EXPECT_EQ(actualFcm->getLogical(0, 3), (FrameChannelMappingAddress(0, 1, 1))); - EXPECT_EQ(actualFcm->getLogical(1, 0), (std::tuple(0, 2, 0))); - EXPECT_EQ(actualFcm->getLogical(1, 1), (std::tuple(0, 2, 1))); - EXPECT_EQ(actualFcm->getLogical(1, 2), (std::tuple(0, 3, 0))); - EXPECT_EQ(actualFcm->getLogical(1, 3), (std::tuple(0, 3, 1))); + EXPECT_EQ(actualFcm->getLogical(1, 0), (FrameChannelMappingAddress(0, 2, 0))); + EXPECT_EQ(actualFcm->getLogical(1, 1), (FrameChannelMappingAddress(0, 2, 1))); + EXPECT_EQ(actualFcm->getLogical(1, 2), (FrameChannelMappingAddress(0, 3, 0))); + EXPECT_EQ(actualFcm->getLogical(1, 3), (FrameChannelMappingAddress(0, 3, 1))); } TEST_F(ProbeImplFcmRemapTest, NonStandard) { @@ -81,16 +82,16 @@ TEST_F(ProbeImplFcmRemapTest, NonStandard) { EXPECT_EQ(actualNFrames, N_FRAMES); EXPECT_EQ(actualNChannels, N_CHANNELS); - EXPECT_EQ(actualFcm->getLogical(0, 0), (std::tuple(0, 0, 0))); - EXPECT_EQ(actualFcm->getLogical(0, 1), (std::tuple(0, 0, 1))); - EXPECT_EQ(actualFcm->getLogical(0, 3), (std::tuple(0, 1, 0))); - EXPECT_EQ(actualFcm->getLogical(0, 2), (std::tuple(0, 1, 1))); + EXPECT_EQ(actualFcm->getLogical(0, 0), (FrameChannelMappingAddress(0, 0, 0))); + EXPECT_EQ(actualFcm->getLogical(0, 1), (FrameChannelMappingAddress(0, 0, 1))); + EXPECT_EQ(actualFcm->getLogical(0, 3), (FrameChannelMappingAddress(0, 1, 0))); + EXPECT_EQ(actualFcm->getLogical(0, 2), (FrameChannelMappingAddress(0, 1, 1))); // Change - EXPECT_EQ(actualFcm->getLogical(1, 3), (std::tuple(0, 2, 0))); - EXPECT_EQ(actualFcm->getLogical(1, 1), (std::tuple(0, 2, 1))); - EXPECT_EQ(actualFcm->getLogical(1, 2), (std::tuple(0, 3, 0))); - EXPECT_EQ(actualFcm->getLogical(1, 0), (std::tuple(0, 3, 1))); + EXPECT_EQ(actualFcm->getLogical(1, 3), (FrameChannelMappingAddress(0, 2, 0))); + EXPECT_EQ(actualFcm->getLogical(1, 1), (FrameChannelMappingAddress(0, 2, 1))); + EXPECT_EQ(actualFcm->getLogical(1, 2), (FrameChannelMappingAddress(0, 3, 0))); + EXPECT_EQ(actualFcm->getLogical(1, 0), (FrameChannelMappingAddress(0, 3, 1))); } } diff --git a/arrus/core/devices/us4r/FrameChannelMappingImpl.cpp b/arrus/core/devices/us4r/FrameChannelMappingImpl.cpp index 7e1778e4d..aee890562 100644 --- a/arrus/core/devices/us4r/FrameChannelMappingImpl.cpp +++ b/arrus/core/devices/us4r/FrameChannelMappingImpl.cpp @@ -22,12 +22,12 @@ FrameChannelMappingImpl::FrameChannelMappingImpl( FrameChannelMappingImpl::~FrameChannelMappingImpl() = default; -std::tuple +FrameChannelMappingAddress FrameChannelMappingImpl::getLogical(FrameNumber frame, ChannelIdx channel) { auto us4oem = us4oemMapping(frame, channel); auto physicalFrame = frameMapping(frame, channel); auto physicalChannel = channelMapping(frame, channel); - return {us4oem, physicalFrame, physicalChannel}; + return FrameChannelMappingAddress{us4oem, physicalFrame, physicalChannel}; } FrameChannelMapping::FrameNumber FrameChannelMappingImpl::getNumberOfLogicalFrames() { diff --git a/arrus/core/devices/us4r/FrameChannelMappingImpl.h b/arrus/core/devices/us4r/FrameChannelMappingImpl.h index b5f6a6dbc..2cfba7300 100644 --- a/arrus/core/devices/us4r/FrameChannelMappingImpl.h +++ b/arrus/core/devices/us4r/FrameChannelMappingImpl.h @@ -3,12 +3,44 @@ #include #include +#include #include #include "arrus/core/api/devices/DeviceId.h" #include "arrus/core/api/devices/us4r/FrameChannelMapping.h" +// Make the FrameChannelMappingAddress available for structure binding +namespace std { + +template<> struct tuple_size<::arrus::devices::FrameChannelMappingAddress>: integral_constant {}; + +template<> +struct tuple_element<0, ::arrus::devices::FrameChannelMappingAddress> { + using type = ::arrus::devices::FrameChannelMapping::Us4OEMNumber; +}; + +template<> +struct tuple_element<1, ::arrus::devices::FrameChannelMappingAddress> { + using type = ::arrus::devices::FrameChannelMapping::FrameNumber; +}; + +template<> +struct tuple_element<2, ::arrus::devices::FrameChannelMappingAddress> { + using type = int8_t; +}; + +template +std::tuple_element_t get( + const ::arrus::devices::FrameChannelMappingAddress& address) +{ + static_assert(Index < 3, "Index out of bounds for FrameChannelMappingAddress"); + if constexpr (Index == 0) return address.getUs4oem(); + if constexpr (Index == 1) return address.getFrame(); + if constexpr (Index == 2) return address.getChannel(); +} + +} namespace arrus::devices { @@ -25,7 +57,7 @@ class FrameChannelMappingImpl : public FrameChannelMapping { FrameChannelMappingImpl(Us4OEMMapping &us4oemMapping, FrameMapping &frameMapping, ChannelMapping &channelMapping, std::vector frameOffsets = {0}); - std::tuple getLogical(FrameNumber frame, ChannelIdx channel) override; + FrameChannelMappingAddress getLogical(FrameNumber frame, ChannelIdx channel) override; uint32 getFirstFrame(uint8 us4oem) override; diff --git a/arrus/core/devices/us4r/us4oem/Us4OEMImpl.cpp b/arrus/core/devices/us4r/us4oem/Us4OEMImpl.cpp index f3d0e4acb..603ecf95c 100644 --- a/arrus/core/devices/us4r/us4oem/Us4OEMImpl.cpp +++ b/arrus/core/devices/us4r/us4oem/Us4OEMImpl.cpp @@ -425,7 +425,8 @@ Us4OEMImpl::setRxMappings(const std::vector &seq) { frameNumber = opId; } fcmBuilder.setChannelMapping(frameNumber, onChannel, - getDeviceId().getOrdinal(), frameNumber, (int8) (mapping.size() - 1)); + static_cast(getDeviceId().getOrdinal()), + frameNumber, (int8) (mapping.size() - 1)); ++onChannel; } ++channel; diff --git a/arrus/core/devices/us4r/us4oem/Us4OEMImplTest.cpp b/arrus/core/devices/us4r/us4oem/Us4OEMImplTest.cpp index cf2eab9ff..1e76738ae 100644 --- a/arrus/core/devices/us4r/us4oem/Us4OEMImplTest.cpp +++ b/arrus/core/devices/us4r/us4oem/Us4OEMImplTest.cpp @@ -8,6 +8,7 @@ #include "arrus/core/devices/us4r/tests/MockIUs4OEM.h" #include "arrus/common/logging/impl/Logging.h" #include "arrus/core/api/ops/us4r/tgc.h" +#include "arrus/core/devices/us4r/FrameChannelMappingImpl.h" namespace { using namespace arrus; From d2a02d31d7d385d5cb85587d0e9a0c67d0b1c712 Mon Sep 17 00:00:00 2001 From: Piotr Jarosik Date: Mon, 1 Nov 2021 14:28:49 +0100 Subject: [PATCH 4/8] Fixed single sequence acquisition. --- api/python/arrus/session.py | 4 +--- api/python/arrus/utils/core.py | 4 ++-- api/python/arrus/utils/imaging.py | 7 +++---- api/python/arrus/utils/us4r_remap_gpu.py | 9 +++++---- .../core/api/devices/us4r/FrameChannelMapping.h | 13 +++++++++---- arrus/core/devices/probe/ProbeImpl.cpp | 3 +-- .../devices/us4r/FrameChannelMappingImpl.cpp | 15 ++++++++++----- .../core/devices/us4r/FrameChannelMappingImpl.h | 17 +++++++++++++---- arrus/core/session/SessionImpl.cpp | 1 - 9 files changed, 44 insertions(+), 29 deletions(-) diff --git a/api/python/arrus/session.py b/api/python/arrus/session.py index ff669ef95..36c13d46b 100644 --- a/api/python/arrus/session.py +++ b/api/python/arrus/session.py @@ -126,9 +126,9 @@ def upload(self, scheme: arrus.ops.us4r.Scheme): "Currently only a sequence with constant number of samples " "can be accepted.") n_samples = next(iter(n_samples)) - input_shape = self._get_physical_frame_shape(fcm, n_samples, rx_batch_size=batch_size) buffer = arrus.framework.DataBuffer(buffer_handle) + input_shape = buffer.elements[0].data.shape const_metadata = arrus.metadata.ConstMetadata( context=fac, data_desc=echo_data_description, @@ -142,7 +142,6 @@ def upload(self, scheme: arrus.ops.us4r.Scheme): raise ValueError("Currently only arrus.utils.imaging.Pipeline " "processing is supported only.") import cupy as cp - out_metadata = processing.prepare(const_metadata) self.gpu_buffer = arrus.utils.imaging.Buffer(n_elements=4, shape=const_metadata.input_shape, @@ -172,7 +171,6 @@ def buffer_callback(elements): print(f"Exception: {type(e)}") except: print("Unknown exception") - pipeline_wrapper = arrus.utils.imaging.PipelineRunner( buffer, self.gpu_buffer, self.out_buffer, processing, buffer_callback) diff --git a/api/python/arrus/utils/core.py b/api/python/arrus/utils/core.py index 4b683d77a..6885fce6e 100644 --- a/api/python/arrus/utils/core.py +++ b/api/python/arrus/utils/core.py @@ -88,8 +88,8 @@ def convert_fcm_to_np_arrays(fcm, n_us4oems): fcm_us4oem[frame, channel] = src_us4oem fcm_frame[frame, channel] = src_frame fcm_channel[frame, channel] = src_channel - frame_offsets = [fcm.getStartFrame(i) for i in range(n_us4oems)] - frame_offsets = np.array(frame_offsets) + frame_offsets = [fcm.getFirstFrame(i) for i in range(n_us4oems)] + frame_offsets = np.array(frame_offsets, dtype=np.uint32) return fcm_us4oem, fcm_frame, fcm_channel, frame_offsets diff --git a/api/python/arrus/utils/imaging.py b/api/python/arrus/utils/imaging.py index 0e026eafd..a62a47692 100644 --- a/api/python/arrus/utils/imaging.py +++ b/api/python/arrus/utils/imaging.py @@ -1734,23 +1734,22 @@ def prepare(self, const_metadata: arrus.metadata.ConstMetadata): self._fcm_channels = cp.asarray(fcm.channels) self._fcm_us4oems = cp.asarray(fcm.us4oems) # 32 - number of us4OEM rx channels, 2 - number of bytes per sample - frame_offsets = fcm.frame_offsets*n_samples*32*2 + frame_offsets = fcm.frame_offsets*n_samples*32 + # TODO constant memory self._frame_offsets = cp.asarray(frame_offsets) self.grid_size, self.block_size = get_default_grid_block_size( self._fcm_frames, n_samples, batch_size ) - def gpu_remap_fn(data): run_remap(self.grid_size, self.block_size, [self._output_buffer, data, self._fcm_frames, self._fcm_channels, self._fcm_us4oems, self._frame_offsets, - n_frames, n_samples, n_channels]) + batch_size, n_frames, n_samples, n_channels]) self._remap_fn = gpu_remap_fn - return const_metadata.copy(input_shape=self.output_shape) def process(self, data): diff --git a/api/python/arrus/utils/us4r_remap_gpu.py b/api/python/arrus/utils/us4r_remap_gpu.py index 450ad14ab..6e84782dd 100644 --- a/api/python/arrus/utils/us4r_remap_gpu.py +++ b/api/python/arrus/utils/us4r_remap_gpu.py @@ -2,6 +2,7 @@ # TODO strategy for case batch size == 1 + _arrus_remap_str = r''' // Naive implementation of data remapping (physical -> logical order). extern "C" @@ -9,7 +10,7 @@ const short* fcmFrames, const char* fcmChannels, const unsigned char *fcmUs4oems, - const int frameOffsets, + const unsigned int *frameOffsets, // Output shape const unsigned nSequences, const unsigned nFrames, const unsigned nSamples, const unsigned nChannels) { @@ -17,8 +18,8 @@ int sample = blockIdx.y * 32 + threadIdx.y; // logical sample int frame = blockIdx.z; // logical frame, global in the whole batch of sequences // Determine sequence number (in batch) and frame number (within sequence) - int sequence = frame / batchSize; - int localFrame = frame % batchSize; + int sequence = frame / nFrames; + int localFrame = frame % nFrames; if(channel >= nChannels || sample >= nSamples || localFrame >= nFrames || sequence >= nSequences) { // outside the range return; @@ -43,7 +44,7 @@ int us4oem = fcmUs4oems[channel + nChannels*localFrame]; int us4oemOffset = frameOffsets[us4oem]; - int indexIn = us4oemOffset // nbytes + int indexIn = us4oemOffset // number of samples // physicalFrame should be calculated relative to the us4oem module begin (first acquired frame should be 0) + physicalFrame*nSamples*32 + sample*32 diff --git a/arrus/core/api/devices/us4r/FrameChannelMapping.h b/arrus/core/api/devices/us4r/FrameChannelMapping.h index 6b320050f..b06a21c3b 100644 --- a/arrus/core/api/devices/us4r/FrameChannelMapping.h +++ b/arrus/core/api/devices/us4r/FrameChannelMapping.h @@ -67,7 +67,7 @@ class FrameChannelMapping { * @param channel logical channel number * @return a tuple: us4oem module number, frame number (within a single sequence), channel number */ - virtual FrameChannelMappingAddress getLogical(FrameNumber frame, ChannelIdx channel) = 0; + virtual FrameChannelMappingAddress getLogical(FrameNumber frame, ChannelIdx channel) const = 0; /** * Returns the number of frame where the given us4OEM data starts. @@ -79,10 +79,15 @@ class FrameChannelMapping { * @param us4oem us4oem ordinal number (0, 1, ...) * @return the number of frame, which starts portion of data acquired by the given us4OEM. */ - virtual arrus::uint32 getFirstFrame(arrus::uint8 us4oem) = 0; + virtual arrus::uint32 getFirstFrame(arrus::uint8 us4oem) const = 0; - virtual FrameNumber getNumberOfLogicalFrames() = 0; - virtual ChannelIdx getNumberOfLogicalChannels() = 0; + /** + * Returns the list of frame offsets ('position of first us4oem frame'). See `getFirstFrame` for more information. + */ + virtual const std::vector & getFrameOffsets() const = 0; + + virtual FrameNumber getNumberOfLogicalFrames() const = 0; + virtual ChannelIdx getNumberOfLogicalChannels() const = 0; /** * Returns true if the given PHYSICAL channel number is unavailable. diff --git a/arrus/core/devices/probe/ProbeImpl.cpp b/arrus/core/devices/probe/ProbeImpl.cpp index 3291db95d..a1c6fe1fd 100644 --- a/arrus/core/devices/probe/ProbeImpl.cpp +++ b/arrus/core/devices/probe/ProbeImpl.cpp @@ -152,8 +152,7 @@ FrameChannelMapping::Handle ProbeImpl::remapFcm(const FrameChannelMapping::Handl if (adapterFcm->getNumberOfLogicalFrames() != nOps) { throw std::runtime_error("Inconsistent mapping and op number of probe's Rx apertures"); } - FrameChannelMappingBuilder builder(adapterFcm->getNumberOfLogicalFrames(), - adapterFcm->getNumberOfLogicalChannels()); + FrameChannelMappingBuilder builder = FrameChannelMappingBuilder::like(*adapterFcm); unsigned short frameNumber = 0; for (const auto &mapping : adapterActiveChannels) { diff --git a/arrus/core/devices/us4r/FrameChannelMappingImpl.cpp b/arrus/core/devices/us4r/FrameChannelMappingImpl.cpp index aee890562..441a83062 100644 --- a/arrus/core/devices/us4r/FrameChannelMappingImpl.cpp +++ b/arrus/core/devices/us4r/FrameChannelMappingImpl.cpp @@ -1,6 +1,7 @@ #include "FrameChannelMappingImpl.h" #include +#include #include "arrus/common/asserts.h" #include "arrus/core/api/common/exceptions.h" @@ -23,29 +24,33 @@ FrameChannelMappingImpl::FrameChannelMappingImpl( FrameChannelMappingImpl::~FrameChannelMappingImpl() = default; FrameChannelMappingAddress -FrameChannelMappingImpl::getLogical(FrameNumber frame, ChannelIdx channel) { +FrameChannelMappingImpl::getLogical(FrameNumber frame, ChannelIdx channel) const { auto us4oem = us4oemMapping(frame, channel); auto physicalFrame = frameMapping(frame, channel); auto physicalChannel = channelMapping(frame, channel); return FrameChannelMappingAddress{us4oem, physicalFrame, physicalChannel}; } -FrameChannelMapping::FrameNumber FrameChannelMappingImpl::getNumberOfLogicalFrames() { +FrameChannelMapping::FrameNumber FrameChannelMappingImpl::getNumberOfLogicalFrames() const { ARRUS_REQUIRES_TRUE(frameMapping.rows() >= 0 && frameMapping.rows() <= std::numeric_limits::max(), "FCM number of logical frames exceeds the maximum number of frames (uint16::max)."); return static_cast(frameMapping.rows()); } -ChannelIdx FrameChannelMappingImpl::getNumberOfLogicalChannels() { +ChannelIdx FrameChannelMappingImpl::getNumberOfLogicalChannels() const { ARRUS_REQUIRES_TRUE(frameMapping.cols() >= 0 && frameMapping.cols() <= std::numeric_limits::max(), "FCM number of logical channels exceeds the maximum number of channels (uint16::max)."); return static_cast(frameMapping.cols()); } -uint32 FrameChannelMappingImpl::getFirstFrame(uint8 us4oem) { +uint32 FrameChannelMappingImpl::getFirstFrame(uint8 us4oem) const { return frameOffsets[us4oem]; } +const std::vector & FrameChannelMappingImpl::getFrameOffsets() const { + return frameOffsets; +} + void FrameChannelMappingBuilder::setChannelMapping(FrameNumber logicalFrame, ChannelIdx logicalChannel, uint8 us4oem, FrameNumber physicalFrame, int8 physicalChannel) { us4oemMapping(logicalFrame, logicalChannel) = us4oem; @@ -69,7 +74,7 @@ FrameChannelMappingBuilder::FrameChannelMappingBuilder(FrameNumber nFrames, Chan } void FrameChannelMappingBuilder::setFrameOffsets(const std::vector &offsets) { - FrameChannelMappingBuilder::frameOffsets = offsets; + this->frameOffsets = offsets; } } diff --git a/arrus/core/devices/us4r/FrameChannelMappingImpl.h b/arrus/core/devices/us4r/FrameChannelMappingImpl.h index 2cfba7300..09b784cd4 100644 --- a/arrus/core/devices/us4r/FrameChannelMappingImpl.h +++ b/arrus/core/devices/us4r/FrameChannelMappingImpl.h @@ -57,13 +57,15 @@ class FrameChannelMappingImpl : public FrameChannelMapping { FrameChannelMappingImpl(Us4OEMMapping &us4oemMapping, FrameMapping &frameMapping, ChannelMapping &channelMapping, std::vector frameOffsets = {0}); - FrameChannelMappingAddress getLogical(FrameNumber frame, ChannelIdx channel) override; + FrameChannelMappingAddress getLogical(FrameNumber frame, ChannelIdx channel) const override; - uint32 getFirstFrame(uint8 us4oem) override; + uint32 getFirstFrame(uint8 us4oem) const override; - FrameNumber getNumberOfLogicalFrames() override; + const std::vector & getFrameOffsets() const override; - ChannelIdx getNumberOfLogicalChannels() override; + FrameNumber getNumberOfLogicalFrames() const override; + + ChannelIdx getNumberOfLogicalChannels() const override; ~FrameChannelMappingImpl() override; @@ -80,6 +82,13 @@ class FrameChannelMappingBuilder { using FrameNumber = FrameChannelMapping::FrameNumber; using Us4OEMNumber = FrameChannelMapping::Us4OEMNumber; + static FrameChannelMappingBuilder like(const FrameChannelMapping &mapping) { + FrameChannelMappingBuilder builder{mapping.getNumberOfLogicalFrames(), + mapping.getNumberOfLogicalChannels()}; + builder.setFrameOffsets(mapping.getFrameOffsets()); + return builder; + } + FrameChannelMappingBuilder(FrameNumber nFrames, ChannelIdx nChannels); void setChannelMapping(FrameNumber logicalFrame, ChannelIdx logicalChannel, // -> diff --git a/arrus/core/session/SessionImpl.cpp b/arrus/core/session/SessionImpl.cpp index b772353d5..ac99ec9a2 100644 --- a/arrus/core/session/SessionImpl.cpp +++ b/arrus/core/session/SessionImpl.cpp @@ -146,7 +146,6 @@ void SessionImpl::startScheme() { void SessionImpl::stopScheme() { std::lock_guard guard(stateMutex); - ASSERT_STATE(State::STARTED); auto us4r = (::arrus::devices::Us4R *) getDevice(DeviceId(DeviceType::Us4R, 0)); us4r->stop(); state = State::STOPPED; From 617b7819fe5760dc36af45496d719faacaf55819 Mon Sep 17 00:00:00 2001 From: Piotr Jarosik Date: Tue, 2 Nov 2021 10:47:55 +0100 Subject: [PATCH 5/8] Work in progres. --- api/python/arrus/kernels/imaging.py | 3 +- api/python/arrus/ops/imaging.py | 3 ++ api/python/arrus/utils/imaging.py | 46 +++++++++++++++++++++++++++++ 3 files changed, 51 insertions(+), 1 deletion(-) diff --git a/api/python/arrus/kernels/imaging.py b/api/python/arrus/kernels/imaging.py index 129d02dc1..62e60c745 100644 --- a/api/python/arrus/kernels/imaging.py +++ b/api/python/arrus/kernels/imaging.py @@ -111,7 +111,8 @@ def process_simple_tx_rx_sequence(context): rx = Rx(rx_aperture, sample_range, op.downsampling_factor, padding=rx_padding) txrx.append(TxRx(tx, rx, op.pri)) - return TxRxSequence(txrx, tgc_curve=tgc_curve, sri=op.sri) + return TxRxSequence(txrx, tgc_curve=tgc_curve, sri=op.sri, + n_repeats=op.n_repeats) def get_aperture_center(tx_aperture_center_element, probe): diff --git a/api/python/arrus/ops/imaging.py b/api/python/arrus/ops/imaging.py index 9a86682a0..0d8b2fcdc 100644 --- a/api/python/arrus/ops/imaging.py +++ b/api/python/arrus/ops/imaging.py @@ -60,6 +60,8 @@ class SimpleTxRxSequence: :param sri: sequence repetition interval - the time between consecutive RF \ frames. When None, the time between consecutive RF frames is determined \ by the total pri only. [s] + :param n_repeats: size of a single batch -- how many times this sequence should be \ + repeated before data is transferred to computer (integer) """ pulse: arrus.ops.us4r.Pulse rx_sample_range: tuple @@ -78,6 +80,7 @@ class SimpleTxRxSequence: tgc_start: float = None tgc_slope: float = None tgc_curve: list = None + n_repeats: int = 1 def __post_init__(self): # Validation diff --git a/api/python/arrus/utils/imaging.py b/api/python/arrus/utils/imaging.py index a62a47692..50a8afdaf 100644 --- a/api/python/arrus/utils/imaging.py +++ b/api/python/arrus/utils/imaging.py @@ -1184,6 +1184,49 @@ def _put_ignore_full(self, data): return data +class SelectSequence(Operation): + """ + Selects frames for a given sequence for further processing. + """ + + def __init__(self, frames): + """ + Constructor. + + :param frames: frames to select + """ + self.frames = frames + + def set_pkgs(self, **kwargs): + pass + + def prepare(self, const_metadata): + input_shape = const_metadata.input_shape + context = const_metadata.context + seq = context.sequence + n_frames = len(self.frames) + + input_n_frames, d2, d3 = input_shape + output_shape = (n_frames, d2, d3) + # TODO make this op less prone to changes in op implementation + if isinstance(seq, arrus.ops.imaging.PwiSequence): + # select appropriate angles + output_angles = seq.angles[self.frames] + new_seq = dataclasses.replace(seq, angles=output_angles) + new_context = const_metadata.context + new_context = arrus.metadata.FrameAcquisitionContext( + device=new_context.device, sequence=new_seq, + raw_sequence=new_context.raw_sequence, + medium=new_context.medium, custom_data=new_context.custom_data) + return const_metadata.copy(input_shape=output_shape, + context=new_context) + else: + return const_metadata.copy(input_shape=output_shape) + + def process(self, data): + return data[self.frames] + + class SelectFrames(Operation): """ Selects frames for a given sequence for further processing. @@ -1230,6 +1273,9 @@ def prepare(self, const_metadata): def process(self, data): return data[self.frames] +# Alias +SelectFrame = SelectFrames + class Squeeze(Operation): """ From 36662d57b139b0e67893c8e3af6071f586d39381 Mon Sep 17 00:00:00 2001 From: Piotr Jarosik Date: Mon, 8 Nov 2021 14:21:36 +0100 Subject: [PATCH 6/8] Fixed expected output batch size. --- api/python/arrus/session.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/api/python/arrus/session.py b/api/python/arrus/session.py index 36c13d46b..eff47f256 100644 --- a/api/python/arrus/session.py +++ b/api/python/arrus/session.py @@ -112,7 +112,7 @@ def upload(self, scheme: arrus.ops.us4r.Scheme): frames=fcm_frame, channels=fcm_channel, frame_offsets=frame_offsets, - batch_size=1) + batch_size=batch_size) # --- Frame acquisition context fac = self._create_frame_acquisition_context(seq, raw_seq, us_device_dto, medium) From 760956ad4cfec52787364f46fcf849d4abb2a842 Mon Sep 17 00:00:00 2001 From: Piotr Jarosik Date: Tue, 9 Nov 2021 10:30:02 +0100 Subject: [PATCH 7/8] Minor fixes in the remapping code. --- api/python/arrus/utils/imaging.py | 23 ++++++++++++---- api/python/arrus/utils/us4r_remap_gpu.py | 34 ++++++++++++++---------- 2 files changed, 38 insertions(+), 19 deletions(-) diff --git a/api/python/arrus/utils/imaging.py b/api/python/arrus/utils/imaging.py index 50a8afdaf..3d873cb24 100644 --- a/api/python/arrus/utils/imaging.py +++ b/api/python/arrus/utils/imaging.py @@ -267,7 +267,8 @@ def process(self, data): outputs.appendleft(output) else: data = step.process(data) - outputs.appendleft(data) + if not self._is_last_endpoint: + outputs.appendleft(data) return outputs def __initialize(self, const_metadata): @@ -298,6 +299,9 @@ def prepare(self, const_metadata): self.__initialize(const_metadata) if not isinstance(self.steps[-1], Pipeline): metadatas.appendleft(current_metadata) + self._is_last_endpoint = False + else: + self._is_last_endpoint = True return metadatas def set_placement(self, device): @@ -1779,11 +1783,19 @@ def prepare(self, const_metadata: arrus.metadata.ConstMetadata): self._fcm_frames = cp.asarray(fcm.frames) self._fcm_channels = cp.asarray(fcm.channels) self._fcm_us4oems = cp.asarray(fcm.us4oems) - # 32 - number of us4OEM rx channels, 2 - number of bytes per sample - frame_offsets = fcm.frame_offsets*n_samples*32 - - # TODO constant memory + frame_offsets = fcm.frame_offsets + # TODO constant memory self._frame_offsets = cp.asarray(frame_offsets) + # For each us4OEM, get number of physical frames this us4OEM gathers. + # Note: this is the max number of us4OEM IN USE. + n_us4oems = cp.max(self._fcm_us4oems).get()+1 + n_frames_us4oems = [] + for us4oem in range(n_us4oems): + n_frames_us4oem = cp.max(self._fcm_frames[self._fcm_us4oems == us4oem]) + n_frames_us4oems.append(n_frames_us4oem) + + # TODO constant memory + self._n_frames_us4oems = cp.asarray(n_frames_us4oems, dtype=cp.uint32)+1 self.grid_size, self.block_size = get_default_grid_block_size( self._fcm_frames, n_samples, batch_size @@ -1793,6 +1805,7 @@ def gpu_remap_fn(data): [self._output_buffer, data, self._fcm_frames, self._fcm_channels, self._fcm_us4oems, self._frame_offsets, + self._n_frames_us4oems, batch_size, n_frames, n_samples, n_channels]) self._remap_fn = gpu_remap_fn diff --git a/api/python/arrus/utils/us4r_remap_gpu.py b/api/python/arrus/utils/us4r_remap_gpu.py index 6e84782dd..c49f7b23a 100644 --- a/api/python/arrus/utils/us4r_remap_gpu.py +++ b/api/python/arrus/utils/us4r_remap_gpu.py @@ -10,45 +10,51 @@ const short* fcmFrames, const char* fcmChannels, const unsigned char *fcmUs4oems, + // Number of sample, that starts given us4oEM data const unsigned int *frameOffsets, + const unsigned int *nFramesUs4OEM, // Output shape - const unsigned nSequences, const unsigned nFrames, const unsigned nSamples, const unsigned nChannels) + const unsigned nSequences, const unsigned nFrames, + const unsigned nSamples, const unsigned nChannels) { - int channel = blockIdx.x * 32 + threadIdx.x; // logical channel - int sample = blockIdx.y * 32 + threadIdx.y; // logical sample + int channel = blockIdx.x*32 + threadIdx.x; // logical channel + int sample = blockIdx.y*32 + threadIdx.y; // logical sample int frame = blockIdx.z; // logical frame, global in the whole batch of sequences // Determine sequence number (in batch) and frame number (within sequence) int sequence = frame / nFrames; int localFrame = frame % nFrames; + if(channel >= nChannels || sample >= nSamples || localFrame >= nFrames || sequence >= nSequences) { // outside the range return; } + // FCM describes here a single sequence int physicalChannel = fcmChannels[channel + nChannels*localFrame]; if(physicalChannel < 0) { // channel is turned off return; } - // [sequence, frame, sample, channel] - int indexOut = sequence*nFrames*nSamples*nChannels - + frame*nSamples*nChannels - + sample*nChannels - + channel; - + size_t indexOut = sequence*nFrames*nSamples*nChannels + + localFrame*nSamples*nChannels + + sample*nChannels + + channel; + + // FCM describes here a single sequence int physicalFrame = fcmFrames[channel + nChannels*localFrame]; // 32 - number of channels in the physical mapping // [us4oem, sequence, physicalFrame, sample, physicalChannel] int us4oem = fcmUs4oems[channel + nChannels*localFrame]; int us4oemOffset = frameOffsets[us4oem]; + int nPhysicalFrames = nFramesUs4OEM[us4oem]; - int indexIn = us4oemOffset // number of samples - // physicalFrame should be calculated relative to the us4oem module begin (first acquired frame should be 0) - + physicalFrame*nSamples*32 - + sample*32 - + physicalChannel; + size_t indexIn = us4oemOffset*nSamples*32 + + sequence*nPhysicalFrames*nSamples*32 + + physicalFrame*nSamples*32 + + sample*32 + + physicalChannel; out[indexOut] = in[indexIn]; }''' From 5e66ce10d8ff23b470d8a020fe63d2e53df710be Mon Sep 17 00:00:00 2001 From: Piotr Jarosik Date: Tue, 9 Nov 2021 21:20:28 +0100 Subject: [PATCH 8/8] Implemented SelectSeqeunceRaw, reduced the size of GPU buffers. --- api/python/arrus/session.py | 4 +- api/python/arrus/utils/imaging.py | 140 +++++++++++++++++++++++------- 2 files changed, 113 insertions(+), 31 deletions(-) diff --git a/api/python/arrus/session.py b/api/python/arrus/session.py index eff47f256..06140f7de 100644 --- a/api/python/arrus/session.py +++ b/api/python/arrus/session.py @@ -143,12 +143,12 @@ def upload(self, scheme: arrus.ops.us4r.Scheme): "processing is supported only.") import cupy as cp out_metadata = processing.prepare(const_metadata) - self.gpu_buffer = arrus.utils.imaging.Buffer(n_elements=4, + self.gpu_buffer = arrus.utils.imaging.Buffer(n_elements=2, shape=const_metadata.input_shape, dtype=const_metadata.dtype, math_pkg=cp, type="locked") - self.out_buffer = [arrus.utils.imaging.Buffer(n_elements=4, + self.out_buffer = [arrus.utils.imaging.Buffer(n_elements=2, shape=m.input_shape, dtype=m.dtype, math_pkg=np, type="locked") diff --git a/api/python/arrus/utils/imaging.py b/api/python/arrus/utils/imaging.py index 8429e5c5b..612017eec 100644 --- a/api/python/arrus/utils/imaging.py +++ b/api/python/arrus/utils/imaging.py @@ -18,7 +18,7 @@ import cupy if cupy.__version__ < "9.0.0": raise Exception(f"The version of cupy module is too low. " - f"Try install the version ''9.0.0'' or higher.") + f"Use version ''9.0.0'' or higher.") def get_extent(x_grid, z_grid): @@ -1193,18 +1193,105 @@ def _put_ignore_full(self, data): return data +class SelectSequenceRaw(Operation): + + def __init__(self, sequence): + if isinstance(sequence, Iterable) and len(sequence) > 1: + raise ValueError("Only a single sequence can be selected") + self.sequence = sequence + self.output = None + self.num_pkg = None + self.positions = None + + def set_pkgs(self, num_pkg, **kwargs): + self.num_pkg = num_pkg + + def prepare(self, const_metadata): + context = const_metadata.context + seq = context.sequence + raw_seq = context.raw_sequence + n_seq = len(self.sequence) + + # For each us4oem, compute tuples: (src_start, dst_start, src_end, dst_end) + # Where each value is the number of rows (we assume 32 columns, i.e. RX channels) + n_samples_set = {op.rx.get_n_samples() for op in raw_seq.ops} + + if len(n_samples_set) > 1: + raise arrus.exceptions.IllegalArgumentError( + f"Each tx/rx in the sequence should acquire the same number of " + f"samples (actual: {n_samples_set})") + n_samples = next(iter(n_samples_set)) + + fcm = const_metadata.data_description.custom["frame_channel_mapping"] + fcm_us4oems = fcm.us4oems + fcm_frames = fcm.frames + # TODO update frame offsets + us4oems = set(fcm.us4oems.flatten().tolist()) + sorted(us4oems) + + self.positions = [] + dst_start = 0 + dst_end = 0 + frame_offsets = [] + current_frame = 0 # Current physical frame. + for us4oem in us4oems: + n_frames = self.num_pkg.max(fcm_frames[fcm_us4oems == us4oem])+1 + us4oem_offset = fcm.frame_offsets[us4oem] + # NOTE: below we use only a single sequence + src_start = us4oem_offset*n_samples+self.sequence[0]*n_frames*n_samples + src_end = src_start+n_frames*n_samples + dst_end = dst_start+n_frames*n_samples + self.positions.append((src_start, dst_start, src_end, dst_end)) + frame_offsets.append(current_frame) + current_frame += n_frames + dst_start = dst_end + + output_shape = (dst_end, 32) + self.output = self.num_pkg.zeros(output_shape, dtype=np.int16) + + # Update const metadata + new_seq = dataclasses.replace(seq, n_repeats=n_seq) + new_raw_seq = dataclasses.replace(raw_seq, n_repeats=n_seq) + new_context = arrus.metadata.FrameAcquisitionContext( + device=context.device, sequence=new_seq, + raw_sequence=new_raw_seq, medium=context.medium, + custom_data=context.custom_data) + + # Update FCM (change the batch_size) + data_desc = const_metadata.data_description + data_desc_custom = data_desc.custom + new_data_desc_custom = data_desc_custom.copy() + fcm = data_desc_custom["frame_channel_mapping"] + new_fcm = dataclasses.replace(fcm, batch_size=1, + frame_offsets=frame_offsets) + new_data_desc_custom["frame_channel_mapping"] = new_fcm + new_data_desc = dataclasses.replace(data_desc, custom=new_data_desc_custom) + + return const_metadata.copy(input_shape=output_shape, + context=new_context, + data_desc=new_data_desc) + + def process(self, data): + for src_start, dst_start, src_end, dst_end in self.positions: + self.output[dst_start:dst_end, :] = data[src_start:src_end, :] + return self.output + + class SelectSequence(Operation): """ - Selects frames for a given sequence for further processing. - """ + Selects sequences for a given batch for further processing. - def __init__(self, frames): - """ - Constructor. + This operator modifies input context so the appropriate + number of sequences is properly set. - :param frames: frames to select - """ - self.frames = frames + :param frames: sequences to select + """ + + def __init__(self, sequence): + if not isinstance(sequence, Iterable): + # Wrap into an array + sequence = [sequence] + self.sequence = sequence def set_pkgs(self, **kwargs): pass @@ -1213,27 +1300,22 @@ def prepare(self, const_metadata): input_shape = const_metadata.input_shape context = const_metadata.context seq = context.sequence - n_frames = len(self.frames) - - input_n_frames, d2, d3 = input_shape - output_shape = (n_frames, d2, d3) - # TODO make this op less prone to changes in op implementation - if isinstance(seq, arrus.ops.imaging.PwiSequence): - # select appropriate angles - output_angles = seq.angles[self.frames] - new_seq = dataclasses.replace(seq, angles=output_angles) - new_context = const_metadata.context - new_context = arrus.metadata.FrameAcquisitionContext( - device=new_context.device, sequence=new_seq, - raw_sequence=new_context.raw_sequence, - medium=new_context.medium, custom_data=new_context.custom_data) - return const_metadata.copy(input_shape=output_shape, - context=new_context) - else: - return const_metadata.copy(input_shape=output_shape) + raw_seq = context.raw_sequence + n_seq = len(self.sequence) + + output_shape = input_shape[1:] + output_shape = (n_seq, ) + output_shape + new_seq = dataclasses.replace(seq, n_repeats=n_seq) + new_raw_seq = dataclasses.replace(raw_seq, n_repeats=n_seq) + new_context = arrus.metadata.FrameAcquisitionContext( + device=context.device, sequence=new_seq, + raw_sequence=new_raw_seq, medium=context.medium, + custom_data=context.custom_data) + return const_metadata.copy(input_shape=output_shape, + context=new_context) def process(self, data): - return data[self.frames] + return data[self.sequence] class SelectFrames(Operation): @@ -1299,7 +1381,7 @@ def process(self, data): def _limit_params(self, value, frames): if value is not None and hasattr(value, "__len__") and len(value) > 1: - return value[frames] + return np.array(value)[frames] else: return value