From 906f931fa60e0a32f05823ae8fcd4651fbea2cec Mon Sep 17 00:00:00 2001
From: Piotr Jarosik <piotr.jarosik@us4us.eu>
Date: Fri, 12 Nov 2021 09:46:01 +0100
Subject: [PATCH 01/21] Improved iq raw 2 LRI performance.

Moved the probe geometry to constant memory, changed kernel  grid of threads dimensions.
---
 api/python/arrus/utils/imaging.py      |  47 +++++--
 api/python/arrus/utils/iq_raw_2_lri.cu | 174 ++++++++++++-------------
 2 files changed, 122 insertions(+), 99 deletions(-)

diff --git a/api/python/arrus/utils/imaging.py b/api/python/arrus/utils/imaging.py
index 33a006ac5..5d63244e3 100644
--- a/api/python/arrus/utils/imaging.py
+++ b/api/python/arrus/utils/imaging.py
@@ -1435,18 +1435,21 @@ def __init__(self, x_grid, z_grid, rx_tang_limits=None):
 
     def set_pkgs(self, num_pkg, **kwargs):
         if num_pkg is np:
-            raise ValueError("Currently reconstructLri operation is "
-                             "implemented for GPU only.")
+            raise ValueError("ReconstructLri operation is implemented for GPU only.")
 
     def prepare(self, const_metadata):
         from pathlib import Path
         import os
+        import cupy as cp
+
         current_dir = os.path.dirname(os.path.join(os.path.abspath(__file__)))
         _kernel_source = Path(os.path.join(current_dir, "iq_raw_2_lri.cu")).read_text()
-        self._kernel = self.num_pkg.RawKernel(_kernel_source, "iqRaw2Lri")
+        self._kernel_module = self.num_pkg.RawModule(code=_kernel_source)
+        self._kernel = self._kernel_module.get_function("iqRaw2Lri")
+        self._z_elem_const = self._kernel_module.get_global("zElemConst")
+        self._tang_elem_const = self._kernel_module.get_global("tangElemConst")
 
         # INPUT PARAMETERS.
-
         # Input data shape.
         self.n_tx, self.n_rx, self.n_samples = const_metadata.input_shape
 
@@ -1461,10 +1464,13 @@ def prepare(self, const_metadata):
         self.output_buffer = self.num_pkg.zeros(output_shape, dtype=self.num_pkg.complex64)
         x_block_size = min(self.x_size, 16)
         z_block_size = min(self.z_size, 16)
-        self.block_size = (z_block_size, x_block_size, 1)
+        tx_block_size = min(self.n_tx, 4)
+        self.block_size = (z_block_size, x_block_size, tx_block_size)
         self.grid_size = (int((self.z_size-1)//z_block_size + 1),
                           int((self.x_size-1)//x_block_size + 1),
-                          self.n_tx)
+                          int((self.n_tx-1)//tx_block_size + 1))
+        print(self.block_size)
+        print(self.grid_size)
         self.x_pix = self.num_pkg.asarray(self.x_grid, dtype=self.num_pkg.float32)
         self.z_pix = self.num_pkg.asarray(self.z_grid, dtype=self.num_pkg.float32)
 
@@ -1478,11 +1484,23 @@ def prepare(self, const_metadata):
         element_pos_x = probe_model.element_pos_x
         element_pos_z = probe_model.element_pos_z
         element_angle_tang = np.tan(probe_model.element_angle)
-        self.x_elem = self.num_pkg.asarray(element_pos_x, dtype=self.num_pkg.float32)
-        self.z_elem = self.num_pkg.asarray(element_pos_z, dtype=self.num_pkg.float32)
-        self.tang_elem = self.num_pkg.asarray(element_angle_tang, dtype=self.num_pkg.float32)
+
         self.n_elements = probe_model.n_elements
 
+        device_props = cp.cuda.runtime.getDeviceProperties(0)
+        if device_props["totalConstMem"] < 256*3*4: # 3 float32 arrays, 256 elements max
+            raise ValueError("There is not enough constant memory available!")
+
+        x_elem = np.asarray(element_pos_x, dtype=self.num_pkg.float32)
+        self._x_elem_const = _get_const_memory_array(
+            self._kernel_module, name="xElemConst", input_array=x_elem)
+        z_elem = np.asarray(element_pos_z, dtype=self.num_pkg.float32)
+        self._z_elem_const = _get_const_memory_array(
+            self._kernel_module, name="zElemConst", input_array=z_elem)
+        tang_elem = np.asarray(element_angle_tang, dtype=self.num_pkg.float32)
+        self._tang_elem_const = _get_const_memory_array(
+            self._kernel_module, name="tangElemConst", input_array=tang_elem)
+
         # TX aperture description
         # Convert the sequence to the positions of the aperture centers
         tx_rx_params = arrus.kernels.imaging.compute_tx_rx_params(
@@ -1530,8 +1548,7 @@ def process(self, data):
         params = (
             self.output_buffer,
             data,
-            self.x_elem, self.z_elem, self.tang_elem, self.n_elements, # DONE
-            self.n_tx, self.n_samples,
+            self.n_elements, self.n_tx, self.n_samples,
             self.z_pix, self.z_size,
             self.x_pix, self.x_size,
             self.sos, self.fs, self.fn,
@@ -1757,3 +1774,11 @@ def process(self, data):
         self._remap_fn(data)
         return self._output_buffer
 
+
+def _get_const_memory_array(module, name, input_array):
+    import cupy as cp
+    const_arr_ptr = module.get_global(name)
+    const_arr = cp.ndarray(shape=input_array.shape, dtype=input_array.dtype,
+                           memptr=const_arr_ptr)
+    const_arr.set(input_array)
+    return const_arr
\ No newline at end of file
diff --git a/api/python/arrus/utils/iq_raw_2_lri.cu b/api/python/arrus/utils/iq_raw_2_lri.cu
index 4b5bcdc28..45bfed761 100644
--- a/api/python/arrus/utils/iq_raw_2_lri.cu
+++ b/api/python/arrus/utils/iq_raw_2_lri.cu
@@ -2,11 +2,14 @@
 
 #define CUDART_PI_F 3.141592654f
 
+__constant__ float zElemConst[256];
+__constant__ float xElemConst[256];
+__constant__ float tangElemConst[256];
+
 extern "C"
 __global__ void
 iqRaw2Lri(complex<float> *iqLri, const complex<float> *iqRaw,
-          const float *xElem, const float *zElem, const float *tangElem, const int nElem,
-          const int nTx, const int nSamp,
+          const int nElem, const int nTx, const int nSamp,
           const float *zPix, const int nZPix,
           const float *xPix, const int nXPix,
           float const sos, float const fs, float const fn,
@@ -19,8 +22,9 @@ iqRaw2Lri(complex<float> *iqLri, const complex<float> *iqRaw,
 
     int z = blockIdx.x * blockDim.x + threadIdx.x;
     int x = blockIdx.y * blockDim.y + threadIdx.y;
+    int iTx = blockIdx.z * blockDim.z + threadIdx.z;
 
-    if (z >= nZPix || x >= nXPix) {
+    if(z >= nZPix || x >= nXPix || iTx >= nTx) {
         return;
     }
 
@@ -36,93 +40,87 @@ iqRaw2Lri(complex<float> *iqLri, const complex<float> *iqRaw,
     const float centRxTang = (maxRxTang + minRxTang) * 0.5f;
     complex<float> pix(0.0f, 0.0f), samp(0.0f, 0.0f), modFactor;
 
-    for (int iTx = 0; iTx < nTx; ++iTx) {
-        int txOffset = iTx*nSamp*nRx;
-
-        if (!isinf(txFoc[iTx])) {
-            /* STA */
-            float zFoc = txApCentZ[iTx] + txFoc[iTx]*cosf(txAngZX[iTx]);
-            float xFoc = txApCentX[iTx] + txFoc[iTx]*sinf(txAngZX[iTx]);
-
-            float pixFocArrang;
-
-            if (txFoc[iTx] <= 0.0f) {
-                /* Virtual Point Source BEHIND probe surface */
-                // Valid pixels are assumed to be always in front of the focal point (VSP)
-                pixFocArrang = 1.0f;
-            }
-            else {
-                /* Virtual Point Source IN FRONT OF probe surface */
-                // Projection of the Foc-Pix vector on the ApCent-Foc vector (dot product) ...
-                // to determine if the pixel is behind (-) or in front of (+) the focal point (VSP).
-                pixFocArrang = (((zPix[z]-zFoc)*(zFoc-txApCentZ[iTx]) +
-                                 (xPix[x]-xFoc)*(xFoc-txApCentX[iTx])) >= 0.f) ? 1.f : -1.f;
-            }
-            txDist	= hypotf(zPix[z] - zFoc, xPix[x] - xFoc);
-            txDist *= pixFocArrang; // Compensation for the Pix-Foc arrangement
-            txDist += txFoc[iTx]; // Compensation for the reference time being the moment when txApCent fires.
-
-            // Projections of Foc-Pix vector on the rotated Foc-ApEdge vectors (dot products) ...
-            // to determine if the pixel is in the sonified area (dot product >= 0).
-            // Foc-ApEdgeFst vector is rotated left, Foc-ApEdgeLst vector is rotated right.
-            txApod = ( ( (-(xElem[txApFstElem[iTx]] - xFoc)*(zPix[z] - zFoc) +
-                          (zElem[txApFstElem[iTx]] - zFoc)*(xPix[x] - xFoc))*pixFocArrang >= 0.f ) &&
-                       ( ( (xElem[txApLstElem[iTx]] - xFoc)*(zPix[z] - zFoc) -
-                           (zElem[txApLstElem[iTx]] - zFoc)*(xPix[x] - xFoc))*pixFocArrang >= 0.f ) ) ? 1.f : 0.f;
-        }
-        else {
-            /* PWI */
-            txDist = (zPix[z] - txApCentZ[iTx]) * cosf(txAngZX[iTx]) +
-                     (xPix[x] - txApCentX[iTx]) * sinf(txAngZX[iTx]);
-
-            // Projections of ApEdge-Pix vector on the rotated unit vector of tx direction (dot products) ...
-            // to determine if the pixel is in the sonified area (dot product >= 0).
-            // For ApEdgeFst, the vector is rotated left, for ApEdgeLst the vector is rotated right.
-            txApod = (((-(zPix[z] - zElem[txApFstElem[iTx]])*sinf(txAngZX[iTx]) +
-                        (xPix[x] - xElem[txApFstElem[iTx]])*cosf(txAngZX[iTx])) >= 0.f) &&
-                      (((zPix[z] - zElem[txApLstElem[iTx]])*sinf(txAngZX[iTx]) -
-                        (xPix[x] - xElem[txApLstElem[iTx]])*cosf(txAngZX[iTx])) >= 0.f)) ? 1.f : 0.f;
+    int txOffset = iTx * nSamp * nRx;
+
+    if(!isinf(txFoc[iTx])) {
+        /* STA */
+        float zFoc = txApCentZ[iTx] + txFoc[iTx] * cosf(txAngZX[iTx]);
+        float xFoc = txApCentX[iTx] + txFoc[iTx] * sinf(txAngZX[iTx]);
+
+        float pixFocArrang;
+
+        if(txFoc[iTx] <= 0.0f) {
+            /* Virtual Point Source BEHIND probe surface */
+            // Valid pixels are assumed to be always in front of the focal point (VSP)
+            pixFocArrang = 1.0f;
+        } else {
+            /* Virtual Point Source IN FRONT OF probe surface */
+            // Projection of the Foc-Pix vector on the ApCent-Foc vector (dot product) ...
+            // to determine if the pixel is behind (-) or in front of (+) the focal point (VSP).
+            pixFocArrang = (((zPix[z] - zFoc) * (zFoc - txApCentZ[iTx]) +
+                             (xPix[x] - xFoc) * (xFoc - txApCentX[iTx])) >= 0.f) ? 1.f : -1.f;
         }
-        pixWgh = 0.0f;
-        pix.real(0.0f);
-        pix.imag(0.0f);
-
-        if (txApod != 0.0f) {
-            for (int iRx = 0; iRx < nRx; iRx++) {
-                iElem = iRx + rxApOrigElem[iTx];
-                if (iElem < 0 || iElem >= nElem) continue;
-
-                rxDist = hypotf(xPix[x] - xElem[iElem], zPix[z] - zElem[iElem]);
-                rxTang = __fdividef(xPix[x] - xElem[iElem], zPix[z] - zElem[iElem]);
-                rxTang = __fdividef(rxTang - tangElem[iElem], 1.f + rxTang*tangElem[iElem]);
-                if (rxTang < minRxTang || rxTang > maxRxTang) continue;
-
-                rxApod = (rxTang - centRxTang) * rngRxTangInv;
-                rxApod = __expf(-rxApod * rxApod * twoSigSqrInv);
-
-                time = (txDist + rxDist)*sosInv + initDel;
-                iSamp = time * fs;
-                if (iSamp < 0.0f || iSamp >= static_cast<float>(nSamp - 1)) {
-                    continue;
-                }
-                offset = txOffset + iRx*nSamp;
-                interpWgh = modff(iSamp, &iSamp);
-                int intSamp = int(iSamp);
-
-                __sincosf(omega*time, &modSin, &modCos);
-                complex<float> modFactor = complex<float>(modCos, modSin);
-
-                samp = iqRaw[offset+intSamp]*(1-interpWgh) + iqRaw[offset+intSamp+1]*interpWgh;
-                pix += samp*modFactor*rxApod;
-                pixWgh += rxApod;
+        txDist = hypotf(zPix[z] - zFoc, xPix[x] - xFoc);
+        txDist *= pixFocArrang; // Compensation for the Pix-Foc arrangement
+        txDist += txFoc[iTx]; // Compensation for the reference time being the moment when txApCent fires.
+
+        // Projections of Foc-Pix vector on the rotated Foc-ApEdge vectors (dot products) ...
+        // to determine if the pixel is in the sonified area (dot product >= 0).
+        // Foc-ApEdgeFst vector is rotated left, Foc-ApEdgeLst vector is rotated right.
+        txApod = (((-(xElemConst[txApFstElem[iTx]] - xFoc) * (zPix[z] - zFoc) +
+                    (zElemConst[txApFstElem[iTx]] - zFoc) * (xPix[x] - xFoc)) * pixFocArrang >= 0.f) &&
+                  (((xElemConst[txApLstElem[iTx]] - xFoc) * (zPix[z] - zFoc) -
+                    (zElemConst[txApLstElem[iTx]] - zFoc) * (xPix[x] - xFoc)) * pixFocArrang >= 0.f)) ? 1.f : 0.f;
+    } else {
+        /* PWI */
+        txDist = (zPix[z] - txApCentZ[iTx]) * cosf(txAngZX[iTx]) +
+                 (xPix[x] - txApCentX[iTx]) * sinf(txAngZX[iTx]);
+
+        // Projections of ApEdge-Pix vector on the rotated unit vector of tx direction (dot products) ...
+        // to determine if the pixel is in the sonified area (dot product >= 0).
+        // For ApEdgeFst, the vector is rotated left, for ApEdgeLst the vector is rotated right.
+        txApod = (((-(zPix[z] - zElemConst[txApFstElem[iTx]]) * sinf(txAngZX[iTx]) +
+                    (xPix[x] - xElemConst[txApFstElem[iTx]]) * cosf(txAngZX[iTx])) >= 0.f) &&
+                  (((zPix[z] - zElemConst[txApLstElem[iTx]]) * sinf(txAngZX[iTx]) -
+                    (xPix[x] - xElemConst[txApLstElem[iTx]]) * cosf(txAngZX[iTx])) >= 0.f)) ? 1.f : 0.f;
+    }
+    pixWgh = 0.0f;
+    pix.real(0.0f);
+    pix.imag(0.0f);
+
+    if(txApod != 0.0f) {
+        for(int iRx = 0; iRx < nRx; iRx++) {
+            iElem = iRx + rxApOrigElem[iTx];
+            if(iElem < 0 || iElem >= nElem) continue;
+
+            rxDist = hypotf(xPix[x] - xElemConst[iElem], zPix[z] - zElemConst[iElem]);
+            rxTang = __fdividef(xPix[x] - xElemConst[iElem], zPix[z] - zElemConst[iElem]);
+            rxTang = __fdividef(rxTang - tangElemConst[iElem], 1.f + rxTang * tangElemConst[iElem]);
+            if(rxTang < minRxTang || rxTang > maxRxTang) continue;
+
+            rxApod = (rxTang - centRxTang) * rngRxTangInv;
+            rxApod = __expf(-rxApod * rxApod * twoSigSqrInv);
+
+            time = (txDist + rxDist) * sosInv + initDel;
+            iSamp = time * fs;
+            if(iSamp < 0.0f || iSamp >= static_cast<float>(nSamp - 1)) {
+                continue;
             }
-        }
-        if(pixWgh == 0.0f) {
-            iqLri[z + x*nZPix + iTx*nZPix*nXPix] = complex<float>(0.0f, 0.0f);
-        }
-        else {
-            iqLri[z + x * nZPix + iTx * nZPix * nXPix] = pix/pixWgh*txApod;
+            offset = txOffset + iRx * nSamp;
+            interpWgh = modff(iSamp, &iSamp);
+            int intSamp = int(iSamp);
+
+            __sincosf(omega * time, &modSin, &modCos);
+            complex<float> modFactor = complex<float>(modCos, modSin);
+
+            samp = iqRaw[offset + intSamp] * (1 - interpWgh) + iqRaw[offset + intSamp + 1] * interpWgh;
+            pix += samp * modFactor * rxApod;
+            pixWgh += rxApod;
         }
     }
-
+    if(pixWgh == 0.0f) {
+        iqLri[z + x * nZPix + iTx * nZPix * nXPix] = complex<float>(0.0f, 0.0f);
+    } else {
+        iqLri[z + x * nZPix + iTx * nZPix * nXPix] = pix / pixWgh * txApod;
+    }
 }
\ No newline at end of file

From f6d7c495c6784b4f16cd5a1094d5b80b1fd159fa Mon Sep 17 00:00:00 2001
From: Piotr Jarosik <piotr.jarosik@us4us.eu>
Date: Sat, 13 Nov 2021 18:34:16 +0100
Subject: [PATCH 02/21] Added phased-scanning for linear probes in Python.

---
 api/python/arrus/devices/us4r.py         |   4 +-
 api/python/arrus/utils/imaging.py        | 264 ++++++++++++++++++++---
 api/python/arrus/utils/rx_beamforming.cu |  84 ++++++++
 3 files changed, 317 insertions(+), 35 deletions(-)
 create mode 100644 api/python/arrus/utils/rx_beamforming.cu

diff --git a/api/python/arrus/devices/us4r.py b/api/python/arrus/devices/us4r.py
index 694d6a9ef..ba27e8999 100644
--- a/api/python/arrus/devices/us4r.py
+++ b/api/python/arrus/devices/us4r.py
@@ -4,8 +4,6 @@
 import ctypes
 import collections.abc
 
-import arrus.utils.core
-import arrus.logging
 from arrus.devices.device import Device, DeviceId, DeviceType
 import arrus.exceptions
 import arrus.devices.probe
@@ -104,10 +102,12 @@ def get_probe_model(self):
         """
         Returns probe model description.
         """
+        import arrus.utils.core
         return arrus.utils.core.convert_to_py_probe_model(
             core_model=self._handle.getProbe(0).getModel())
 
     def _get_dto(self):
+        import arrus.utils.core
         probe_model = arrus.utils.core.convert_to_py_probe_model(
             core_model=self._handle.getProbe(0).getModel())
         probe_dto = arrus.devices.probe.ProbeDTO(model=probe_model)
diff --git a/api/python/arrus/utils/imaging.py b/api/python/arrus/utils/imaging.py
index 5d63244e3..4c380399e 100644
--- a/api/python/arrus/utils/imaging.py
+++ b/api/python/arrus/utils/imaging.py
@@ -14,6 +14,8 @@
 import threading
 from collections import deque
 from collections.abc import Iterable
+from pathlib import Path
+import os
 
 
 def get_extent(x_grid, z_grid):
@@ -706,11 +708,140 @@ def _legacy_decimate(self, data):
 class RxBeamforming(Operation):
     """
     Classical rx beamforming (reconstructing scanline by scanline).
+    This operator implements beamforming for linear scanning (element by element)
+    and phased scanning (angle by angle).
+    """
+    def __init__(self, num_pkg=None):
+        # Actual implementation of the operator.
+        self._op = None
+        self.xp = None
+
+    def set_pkgs(self, num_pkg, **kwargs):
+        self.xp = num_pkg
+
+    def prepare(self, const_metadata):
+        seq = const_metadata.context.sequence
+        # Determine scanning type based on the sequence of parameters.
+        tx_centers = seq.tx_aperture_center_element
+        if tx_centers is None:
+            tx_centers = seq.tx_aperture_center
+        tx_centers = set(np.atleast_1d(tx_centers))
+        tx_angles = set(np.atleast_1d(seq.angles))
+        # Phased array scanning:
+        # - single TX/RX aperture position
+        # - multiple different angles
+        if len(tx_centers) == 1 and len(tx_angles) > 1:
+            self._op = RxBeamformingPhasedScanning(num_pkg=self.xp)
+        # Linear array scanning:
+        # - single transmit angle (equal 0)
+        # - multiple different aperture positions
+        elif len(tx_centers) > 1 and len(tx_angles) == 1:
+            self._op = RxBeamformingLin(num_pkg=self.xp)
+        # Otherwise: unsupported scanning method (linear/phased)
+        else:
+            raise ValueError("RX beamformer does not support parameters of "
+                             "the provided TX/RX sequence.")
+        return self._op.prepare(const_metadata)
+
+    def process(self, data):
+        return self._op.process(data)
 
-    Expected input data shape: n_emissions, n_rx, n_samples
 
-    Currently the beamforming op works only for LIN sequence output data.
+class RxBeamformingPhasedScanning(Operation):
+    """
+    Classical beamforming for phase array scanning.
     """
+    def __init__(self, num_pkg=None):
+        self.num_pkg = num_pkg
+
+    def prepare(self, const_metadata):
+        import cupy as cp
+        if self.num_pkg != cp:
+            raise ValueError("Phased scanning is implemented for GPU only.")
+        probe_model = const_metadata.context.device.probe.model
+        if probe_model.is_convex_array():
+            raise ValueError("Phased array scanning is implemented for "
+                             "linear phased arrays only.")
+
+        self._kernel_module = _read_kernel_module("rx_beamforming.cu")
+        self._kernel = self._kernel_module.get_function("beamformPhasedArray")
+
+        self.n_tx, self.n_rx, self.n_samples = const_metadata.input_shape
+        self.output_buffer = cp.zeros((self.n_tx, self.n_samples), dtype=cp.complex64)
+
+        seq = const_metadata.context.sequence
+        self.tx_angles = cp.asarray(seq.angles, dtype=cp.float32)
+
+        device_fs = const_metadata.context.device.sampling_frequency
+        acq_fs = (device_fs/seq.downsampling_factor)
+        fs = const_metadata.data_description.sampling_frequency
+        fc = seq.pulse.center_frequency
+        n_periods = seq.pulse.n_periods
+        medium = const_metadata.context.medium
+        if seq.speed_of_sound is not None:
+            c = seq.speed_of_sound
+        else:
+            c = medium.speed_of_sound
+        start_sample, end_sample = seq.rx_sample_range
+        initial_delay = - start_sample / acq_fs
+        if seq.init_delay == "tx_start":
+            burst_factor = n_periods / (2 * fc)
+            tx_rx_params = arrus.kernels.imaging.compute_tx_rx_params(
+                probe_model, seq, c)
+            tx_center_delay = tx_rx_params["tx_center_delay"]
+            initial_delay += tx_center_delay + burst_factor
+        elif not seq.init_delay == "tx_center":
+            raise ValueError(f"Unrecognized init_delay value: {initial_delay}")
+        lambd = c / fc
+        max_tang = abs(math.tan(
+            math.asin(min(1, 2 / 3 * lambd / probe_model.pitch))))
+        print(f"MAX TANG: {max_tang}")
+
+        self.fc = cp.float32(fc)
+        self.fs = cp.float32(fs)
+        self.c = cp.float32(c)
+        # Note: start sample has to be appropriately adjusted for
+        # the ACQ sampling frequency.
+        self.start_time = cp.float32(start_sample/acq_fs)
+        self.init_delay = cp.float32(initial_delay)
+        self.max_tang = cp.float32(max_tang)
+        scanline_block_size = min(self.n_tx, 16)
+        sample_block_size = min(self.n_samples, 16)
+        self.block_size = (sample_block_size, scanline_block_size, 1)
+        self.grid_size = (int((self.n_samples-1)//sample_block_size + 1),
+                          int((self.n_tx-1)//scanline_block_size + 1),
+                          1)
+        # xElemConst
+        # Get aperture origin (for the given aperture center element/aperture center)
+        tx_rx_params = arrus.kernels.imaging.preprocess_sequence_parameters(probe_model, seq)
+        # There is a single TX and RX aperture center for all TX/RXs
+        rx_aperture_center_element = np.array(tx_rx_params["rx_ap_cent"])[0]
+        rx_aperture_origin = _get_rx_aperture_origin(
+            rx_aperture_center_element, seq.rx_aperture_size)
+        rx_aperture_offset = rx_aperture_center_element-rx_aperture_origin
+        x_elem = (np.arange(0, self.n_rx)-rx_aperture_offset) * probe_model.pitch
+        x_elem = x_elem.astype(np.float32)
+        self.x_elem_const = _get_const_memory_array(
+            self._kernel_module, "xElemConst", x_elem)
+        return const_metadata.copy(input_shape=self.output_buffer.shape)
+
+    def process(self, data):
+        data = self.num_pkg.ascontiguousarray(data)
+        params = (
+            self.output_buffer, data,
+            self.delays,
+            self.n_tx, self.n_rx, self.n_samples,
+            self.tx_angles,
+            self.init_delay, self.start_time,
+            self.c, self.fs, self.fc, self.max_tang)
+        self._kernel(self.grid_size, self.block_size, params)
+        # import matplotlib.pyplot as plt
+        # plt.imshow(self.delays.get())
+        # plt.show()
+        return self.output_buffer
+
+
+class RxBeamformingLin(Operation):
 
     def __init__(self, num_pkg=None):
         self.delays = None
@@ -719,8 +850,7 @@ def __init__(self, num_pkg=None):
         self.xp = num_pkg
         self.interp1d_func = None
 
-    def set_pkgs(self, num_pkg, **kwargs):
-        self.xp = num_pkg
+    def _set_interpolator(self, **kwargs):
         if self.xp is np:
             import scipy.interpolate
 
@@ -744,14 +874,14 @@ def numpy_interp1d(input, samples, output):
             self.interp1d_func = arrus.utils.interpolate.interp1d
 
     def prepare(self, const_metadata: arrus.metadata.ConstMetadata):
-        # TODO verify that all angles, focal points are the same
-        # TODO make sure start_sample is computed appropriately
+        self._set_interpolator()
         context = const_metadata.context
         probe_model = const_metadata.context.device.probe.model
         seq = const_metadata.context.sequence
         raw_seq = const_metadata.context.raw_sequence
         medium = const_metadata.context.medium
-        rx_aperture_center_element = np.array(seq.rx_aperture_center_element)
+        tx_rx_params = arrus.kernels.imaging.preprocess_sequence_parameters(probe_model, seq)
+        rx_aperture_center_element = np.array(tx_rx_params["tx_ap_cent"])
 
         self.n_tx, self.n_rx, self.n_samples = const_metadata.input_shape
         self.is_iq = const_metadata.is_iq_data
@@ -774,11 +904,9 @@ def prepare(self, const_metadata: arrus.metadata.ConstMetadata):
             c = seq.speed_of_sound
         else:
             c = medium.speed_of_sound
-        tx_angle = 0  # TODO use appropriate tx angle
+        tx_angle = 0
         start_sample = seq.rx_sample_range[0]
-        rx_aperture_origin = _get_rx_aperture_origin(seq)
-
-
+        rx_aperture_origin = _get_rx_aperture_origin(rx_aperture_center_element, seq.rx_aperture_size)
         # -start_sample compensates the fact, that the data indices always
         # start from 0
         initial_delay = - start_sample / acq_fs
@@ -797,7 +925,7 @@ def prepare(self, const_metadata: arrus.metadata.ConstMetadata):
         z_distance = radial_distance * np.cos(tx_angle).reshape(1, -1)
 
         origin_offset = (rx_aperture_origin[0]
-                         - (seq.rx_aperture_center_element[0]))
+                         - (rx_aperture_center_element[0]))
         # New coordinate system: origin: rx aperture center
         element_position = ((np.arange(0, self.n_rx) + origin_offset)
                             * probe_model.pitch)
@@ -940,6 +1068,7 @@ def __init__(self, x_grid, z_grid):
     def set_pkgs(self, num_pkg, **kwargs):
         if num_pkg != np:
             self.is_gpu = True
+        self.num_pkg = num_pkg
         # Ignoring provided num. package - currently CPU implementation is
         # available only.
 
@@ -949,9 +1078,29 @@ def prepare(self, const_metadata: arrus.metadata.ConstMetadata):
             self.process = self._process_convex
             return self._prepare_convex(const_metadata)
         else:
-            # linear array
-            self.process = self._process_linear_array
-            return self._prepare_linear_array(const_metadata)
+            # linear array or phased array
+            seq = const_metadata.context.sequence
+            # Determine scanning type based on the sequence of parameters.
+            tx_centers = seq.tx_aperture_center_element
+            if tx_centers is None:
+                tx_centers = seq.tx_aperture_center
+            tx_centers = set(np.atleast_1d(tx_centers))
+            tx_angles = set(np.atleast_1d(seq.angles))
+            # Phased array scanning:
+            # - single TX/RX aperture position
+            # - multiple different angles
+            if len(tx_centers) == 1 and len(tx_angles) > 1:
+                self.process = self._process_phased_array
+                return self._prepare_phased_array(const_metadata)
+            # Linear array scanning:
+            # - single transmit angle (equal 0)
+            # - multiple different aperture positions
+            elif len(tx_centers) > 1 and len(tx_angles) == 1:
+                self.process = self._process_linear_array
+                return self._prepare_phased_array(const_metadata)
+            else:
+                raise ValueError("The given combination of TX/RX parameters is "
+                                 "not supported by ScanConversion")
 
     def _prepare_linear_array(self, const_metadata: arrus.metadata.ConstMetadata):
         # Determine interpolation function.
@@ -961,15 +1110,14 @@ def _prepare_linear_array(self, const_metadata: arrus.metadata.ConstMetadata):
         import cupy as cp
         import cupyx.scipy.ndimage
         self.interp_function = cupyx.scipy.ndimage.map_coordinates
-
         n_samples, n_scanlines = const_metadata.input_shape
         seq = const_metadata.context.sequence
         if not isinstance(seq, arrus.ops.imaging.LinSequence):
             raise ValueError("Scan conversion works only with LinSequence.")
-
         medium = const_metadata.context.medium
         probe = const_metadata.context.device.probe.model
-
+        tx_rx_params = arrus.kernels.imaging.preprocess_sequence_parameters(probe, seq)
+        tx_aperture_center_element = tx_rx_params["tx_ap_cent"]
         n_elements = probe.n_elements
         if n_elements % 2 != 0:
             raise ValueError("Even number of probe elements is required.")
@@ -979,16 +1127,15 @@ def _prepare_linear_array(self, const_metadata: arrus.metadata.ConstMetadata):
             c = seq.speed_of_sound
         else:
             c = medium.speed_of_sound
-        tx_center_elements = seq.tx_aperture_center_element
-        tx_center_diff = set(np.diff(tx_center_elements))
+        tx_center_diff = set(np.diff(tx_aperture_center_element))
         if len(tx_center_diff) != 1:
             raise ValueError("Transmits should be done by consecutive "
                              "center elements (got tx center elements: "
-                             f"{tx_center_elements}")
+                             f"{tx_aperture_center_element}")
         tx_center_diff = next(iter(tx_center_diff))
         # Determine input grid.
         input_x_grid_diff = tx_center_diff*pitch
-        input_x_grid_origin = tx_center_elements[0]-(n_elements-1)/2*pitch
+        input_x_grid_origin = tx_aperture_center_element[0]-(n_elements-1)/2*pitch
         acq_fs = (const_metadata.context.device.sampling_frequency
                   / seq.downsampling_factor)
         fs = data_desc.sampling_frequency
@@ -1055,6 +1202,48 @@ def _process_convex(self, data):
             bounds_error=False, fill_value=0)
         return self.interpolator(self.dst_points).reshape(self.dst_shape)
 
+    def _prepare_phased_array(self, const_metadata: arrus.metadata.ConstMetadata):
+        probe = const_metadata.context.device.probe.model
+        data_desc = const_metadata.data_description
+
+        n_samples, _ = const_metadata.input_shape
+        seq = const_metadata.context.sequence
+        fs = const_metadata.context.device.sampling_frequency
+        acq_fs = fs / seq.downsampling_factor
+        fs = data_desc.sampling_frequency
+        start_sample, _ = seq.rx_sample_range
+        start_time = start_sample/acq_fs
+        c = _get_speed_of_sound(const_metadata.context)
+        tx_rx_params = arrus.kernels.imaging.preprocess_sequence_parameters(probe, seq)
+        tx_ap_cent_elem = np.array(tx_rx_params["tx_ap_cent"])[0]
+        tx_ap_cent_ang, tx_ap_cent_x, tx_ap_cent_z = arrus.kernels.imaging.get_aperture_center(
+            tx_ap_cent_elem, probe)
+
+        # There is a single position of TX aperture.
+        tx_ap_cent_x = tx_ap_cent_x.squeeze().item()
+        tx_ap_cent_z = tx_ap_cent_z.squeeze().item()
+        tx_ap_cent_ang = tx_ap_cent_ang.squeeze().item()
+
+        self.radGridIn = (start_time + np.arange(0, n_samples)/fs)*c/2
+        self.azimuthGridIn = seq.angles + tx_ap_cent_ang
+        azimuthGridOut = np.arctan2((self.x_grid-tx_ap_cent_x), (self.z_grid.T-tx_ap_cent_z))
+        radGridOut = np.sqrt((self.x_grid-tx_ap_cent_x)**2 + (self.z_grid.T-tx_ap_cent_z)**2)
+        dst_points = np.dstack((radGridOut, azimuthGridOut))
+        w, h, d = dst_points.shape
+        self.dst_points = dst_points.reshape((w * h, d))
+        self.dst_shape = len(self.z_grid.squeeze()), len(self.x_grid.squeeze())
+        return const_metadata.copy(input_shape=self.dst_shape)
+
+    def _process_phased_array(self, data):
+        if self.is_gpu:
+            data = data.get()
+        data[np.isnan(data)] = 0.0
+        self.interpolator = scipy.interpolate.RegularGridInterpolator(
+            (self.radGridIn, self.azimuthGridIn), data, method="linear",
+            bounds_error=False, fill_value=0)
+        result = self.interpolator(self.dst_points).reshape(self.dst_shape)
+        return self.num_pkg.asarray(result)
+
 
 class LogCompression(Operation):
     """
@@ -1438,8 +1627,7 @@ def set_pkgs(self, num_pkg, **kwargs):
             raise ValueError("ReconstructLri operation is implemented for GPU only.")
 
     def prepare(self, const_metadata):
-        from pathlib import Path
-        import os
+
         import cupy as cp
 
         current_dir = os.path.dirname(os.path.join(os.path.abspath(__file__)))
@@ -1469,8 +1657,6 @@ def prepare(self, const_metadata):
         self.grid_size = (int((self.z_size-1)//z_block_size + 1),
                           int((self.x_size-1)//x_block_size + 1),
                           int((self.n_tx-1)//tx_block_size + 1))
-        print(self.block_size)
-        print(self.grid_size)
         self.x_pix = self.num_pkg.asarray(self.x_grid, dtype=self.num_pkg.float32)
         self.z_pix = self.num_pkg.asarray(self.z_grid, dtype=self.num_pkg.float32)
 
@@ -1610,12 +1796,8 @@ def process(self, data):
         return self.num_pkg.mean(data, axis=self.axis)
 
 
-def _get_rx_aperture_origin(sequence):
-    rx_aperture_size = sequence.rx_aperture_size
-    rx_aperture_center_element = np.array(sequence.rx_aperture_center_element)
-    rx_aperture_origin = np.round(rx_aperture_center_element -
-                               (rx_aperture_size - 1) / 2 + 1e-9)
-    return rx_aperture_origin
+def _get_rx_aperture_origin(aperture_center_element, aperture_size):
+    return np.round(aperture_center_element-(aperture_size-1)/2+1e-9)
 
 
 # -------------------------------------------- RF frame remapping.
@@ -1781,4 +1963,20 @@ def _get_const_memory_array(module, name, input_array):
     const_arr = cp.ndarray(shape=input_array.shape, dtype=input_array.dtype,
                            memptr=const_arr_ptr)
     const_arr.set(input_array)
-    return const_arr
\ No newline at end of file
+    return const_arr
+
+
+def _read_kernel_module(path):
+    import cupy as cp
+    current_dir = os.path.dirname(os.path.join(os.path.abspath(__file__)))
+    kernel_src = Path(os.path.join(current_dir, path)).read_text()
+    return cp.RawModule(code=kernel_src)
+
+
+def _get_speed_of_sound(context):
+    seq = context.sequence
+    medium = context.medium
+    if seq.speed_of_sound is not None:
+        return seq.speed_of_sound
+    else:
+        return medium.speed_of_sound
\ No newline at end of file
diff --git a/api/python/arrus/utils/rx_beamforming.cu b/api/python/arrus/utils/rx_beamforming.cu
new file mode 100644
index 000000000..354b3209d
--- /dev/null
+++ b/api/python/arrus/utils/rx_beamforming.cu
@@ -0,0 +1,84 @@
+#include <cupy/complex.cuh>
+
+#define CUDART_PI_F 3.141592654f
+
+__constant__ float xElemConst[256]; // [m]
+
+// Assumptions:
+// - TX and RX apertures have the same center position
+// - x/z/angle/Elem refers to the RX aperture (i.e. the first value is the position of first aperture element, relative
+// to the center of TX/RX aperture
+extern "C"
+__global__ void beamformPhasedArray(complex<float> *output, const complex<float> *input,
+                                    float* delays, // DEBUG
+                                    const unsigned nTx, const unsigned nRx, const unsigned nSamples,
+                                    const float *txAngles, // [rad]
+                                    const float initDelay, const float startTime,
+                                    const float c, const float fs, const float fc, float maxApodTang) {
+    complex<float> a, b;
+    float elementX, elementZ;
+    float rxTang, pixWgh = 0;
+    float txDistance, rxDistance, time, s, txAngleSin, txAngleCos;
+    float modSin, modCos;
+    unsigned signalOffset;
+    int sInt;
+    unsigned sample = blockIdx.x * blockDim.x + threadIdx.x;
+    unsigned scanline = blockIdx.y * blockDim.y + threadIdx.y;
+    complex<float> result = complex<float>(0.0f, 0.0f);
+    complex<float> currentResult;
+    complex<float> modFactor;
+
+    if(sample >= nSamples || scanline >= nTx) {
+        return;
+    }
+
+    float txAngle = txAngles[scanline];
+
+    float r = (sample/fs + startTime)*c/2;
+    __sincosf(txAngle, &txAngleSin, &txAngleCos);
+
+    // Note: relative to the center of aperture.
+    float pointX = r*txAngleSin;
+    float pointZ = r*txAngleCos;
+    txDistance = r;
+
+    unsigned txOffset = scanline*nSamples*nRx;
+    float cInv = 1/c;
+
+    for(unsigned element = 0; element < nRx; ++element) {
+        // Note: relative to the center of aperture.
+        elementX = xElemConst[element];
+        elementZ = 0; // Linear array.
+
+        // RX apodization.
+        rxTang = (pointX-elementX)/(pointZ-elementZ);
+        if(fabs(rxTang) > maxApodTang) {
+            continue;
+        }
+        // RX distance and sample number for given RX element.
+        rxDistance = hypotf(elementX-pointX, elementZ-pointZ);
+        time = (txDistance + rxDistance) * cInv + initDelay;
+        s = time * fs;
+        sInt = (int) s;
+
+        signalOffset = txOffset + element*nSamples;
+        if(sInt >= 0 && sInt < nSamples-1) {
+            float ratio = s - sInt;
+            a = input[signalOffset + sInt];
+            b = input[signalOffset + sInt + 1];
+            currentResult = (1.0f - ratio) * a + ratio * b;
+        }
+        else {
+            continue;
+        }
+        __sincosf(2.0f * CUDART_PI_F*fc*time, &modSin, &modCos);
+        modFactor = complex<float>(modCos, modSin);
+        result = currentResult*modFactor;
+        ++pixWgh;
+    }
+    if(pixWgh == 0.0f) {
+        output[sample + scanline*nSamples] = complex<float>(0.0f, 0.0f);
+    } else {
+        output[sample + scanline*nSamples] = result;
+    }
+}
\ No newline at end of file

From c8ba29200821d64521a9c3d2e8f448e81695d264 Mon Sep 17 00:00:00 2001
From: Piotr Jarosik <piotr.jarosik@us4us.eu>
Date: Sat, 13 Nov 2021 19:45:40 +0100
Subject: [PATCH 03/21] Minor fixes.

---
 api/python/arrus/ops/imaging.py          | 1 -
 api/python/arrus/utils/imaging.py        | 7 +------
 api/python/arrus/utils/rx_beamforming.cu | 1 -
 3 files changed, 1 insertion(+), 8 deletions(-)

diff --git a/api/python/arrus/ops/imaging.py b/api/python/arrus/ops/imaging.py
index 9a86682a0..64e873541 100644
--- a/api/python/arrus/ops/imaging.py
+++ b/api/python/arrus/ops/imaging.py
@@ -167,7 +167,6 @@ def __post_init__(self):
         super().__post_init__()
         if self.tx_focus <= 0 or np.isinf(self.tx_focus):
             raise ValueError("TX focus has to be a positive value.")
-        assert_is_scalar("angles", self.angles)
 
 
 @dataclasses.dataclass(frozen=True)
diff --git a/api/python/arrus/utils/imaging.py b/api/python/arrus/utils/imaging.py
index 4c380399e..2644e7db8 100644
--- a/api/python/arrus/utils/imaging.py
+++ b/api/python/arrus/utils/imaging.py
@@ -795,7 +795,6 @@ def prepare(self, const_metadata):
         lambd = c / fc
         max_tang = abs(math.tan(
             math.asin(min(1, 2 / 3 * lambd / probe_model.pitch))))
-        print(f"MAX TANG: {max_tang}")
 
         self.fc = cp.float32(fc)
         self.fs = cp.float32(fs)
@@ -829,15 +828,11 @@ def process(self, data):
         data = self.num_pkg.ascontiguousarray(data)
         params = (
             self.output_buffer, data,
-            self.delays,
             self.n_tx, self.n_rx, self.n_samples,
             self.tx_angles,
             self.init_delay, self.start_time,
             self.c, self.fs, self.fc, self.max_tang)
         self._kernel(self.grid_size, self.block_size, params)
-        # import matplotlib.pyplot as plt
-        # plt.imshow(self.delays.get())
-        # plt.show()
         return self.output_buffer
 
 
@@ -1242,7 +1237,7 @@ def _process_phased_array(self, data):
             (self.radGridIn, self.azimuthGridIn), data, method="linear",
             bounds_error=False, fill_value=0)
         result = self.interpolator(self.dst_points).reshape(self.dst_shape)
-        return self.num_pkg.asarray(result)
+        return self.num_pkg.asarray(result).astype(np.float32)
 
 
 class LogCompression(Operation):
diff --git a/api/python/arrus/utils/rx_beamforming.cu b/api/python/arrus/utils/rx_beamforming.cu
index 354b3209d..f13b232a0 100644
--- a/api/python/arrus/utils/rx_beamforming.cu
+++ b/api/python/arrus/utils/rx_beamforming.cu
@@ -10,7 +10,6 @@ __constant__ float xElemConst[256]; // [m]
 // to the center of TX/RX aperture
 extern "C"
 __global__ void beamformPhasedArray(complex<float> *output, const complex<float> *input,
-                                    float* delays, // DEBUG
                                     const unsigned nTx, const unsigned nRx, const unsigned nSamples,
                                     const float *txAngles, // [rad]
                                     const float initDelay, const float startTime,

From df28f6e876de8ffb24c1754df31ac0425b9bcbf1 Mon Sep 17 00:00:00 2001
From: Piotr Jarosik <piotr.jarosik@us4us.eu>
Date: Sun, 14 Nov 2021 09:53:54 +0100
Subject: [PATCH 04/21] Minor fixes in the PA scanning beamformer.

---
 api/python/arrus/utils/rx_beamforming.cu | 9 ++++++---
 1 file changed, 6 insertions(+), 3 deletions(-)

diff --git a/api/python/arrus/utils/rx_beamforming.cu b/api/python/arrus/utils/rx_beamforming.cu
index f13b232a0..3d3c163f3 100644
--- a/api/python/arrus/utils/rx_beamforming.cu
+++ b/api/python/arrus/utils/rx_beamforming.cu
@@ -67,17 +67,20 @@ __global__ void beamformPhasedArray(complex<float> *output, const complex<float>
             b = input[signalOffset + sInt + 1];
             currentResult = (1.0f - ratio) * a + ratio * b;
         }
+        else if(sInt == nSamples-1) {
+            currentResult = input[signalOffset + sInt];
+        }
         else {
             continue;
         }
-        __sincosf(2.0f * CUDART_PI_F*fc*time, &modSin, &modCos);
+        __sincosf(2.0f*CUDART_PI_F*fc*time, &modSin, &modCos);
         modFactor = complex<float>(modCos, modSin);
-        result = currentResult*modFactor;
+        result += currentResult*modFactor;
         ++pixWgh;
     }
     if(pixWgh == 0.0f) {
         output[sample + scanline*nSamples] = complex<float>(0.0f, 0.0f);
     } else {
-        output[sample + scanline*nSamples] = result;
+        output[sample + scanline*nSamples] = result/pixWgh;
     }
 }
\ No newline at end of file

From cd82c6b22fb2719767a786914d3120aa1c09a43f Mon Sep 17 00:00:00 2001
From: Piotr Jarosik <piotr.jarosik@us4us.eu>
Date: Sun, 14 Nov 2021 11:59:30 +0100
Subject: [PATCH 05/21] Minor fixes in the implementation of classical
 beamforming processsing.

---
 api/python/arrus/utils/imaging.py | 17 ++++++++---------
 1 file changed, 8 insertions(+), 9 deletions(-)

diff --git a/api/python/arrus/utils/imaging.py b/api/python/arrus/utils/imaging.py
index 2644e7db8..e7b42e48a 100644
--- a/api/python/arrus/utils/imaging.py
+++ b/api/python/arrus/utils/imaging.py
@@ -1092,7 +1092,7 @@ def prepare(self, const_metadata: arrus.metadata.ConstMetadata):
             # - multiple different aperture positions
             elif len(tx_centers) > 1 and len(tx_angles) == 1:
                 self.process = self._process_linear_array
-                return self._prepare_phased_array(const_metadata)
+                return self._prepare_linear_array(const_metadata)
             else:
                 raise ValueError("The given combination of TX/RX parameters is "
                                  "not supported by ScanConversion")
@@ -1118,16 +1118,14 @@ def _prepare_linear_array(self, const_metadata: arrus.metadata.ConstMetadata):
             raise ValueError("Even number of probe elements is required.")
         pitch = probe.pitch
         data_desc = const_metadata.data_description
-        if seq.speed_of_sound is not None:
-            c = seq.speed_of_sound
-        else:
-            c = medium.speed_of_sound
-        tx_center_diff = set(np.diff(tx_aperture_center_element))
-        if len(tx_center_diff) != 1:
+        c = _get_speed_of_sound(const_metadata.context)
+        tx_center_diff = np.diff(tx_aperture_center_element)
+        # Check if tx aperture centers are evenly spaced.
+        if not np.allclose(tx_center_diff, [tx_center_diff[0]]*len(tx_center_diff)):
             raise ValueError("Transmits should be done by consecutive "
                              "center elements (got tx center elements: "
                              f"{tx_aperture_center_element}")
-        tx_center_diff = next(iter(tx_center_diff))
+        tx_center_diff = tx_center_diff[0]
         # Determine input grid.
         input_x_grid_diff = tx_center_diff*pitch
         input_x_grid_origin = tx_aperture_center_element[0]-(n_elements-1)/2*pitch
@@ -1195,7 +1193,8 @@ def _process_convex(self, data):
         self.interpolator = scipy.interpolate.RegularGridInterpolator(
             (self.radGridIn, self.azimuthGridIn), data, method="linear",
             bounds_error=False, fill_value=0)
-        return self.interpolator(self.dst_points).reshape(self.dst_shape)
+        result = self.interpolator(self.dst_points).reshape(self.dst_shape)
+        return self.num_pkg.asarray(result).astype(np.float32)
 
     def _prepare_phased_array(self, const_metadata: arrus.metadata.ConstMetadata):
         probe = const_metadata.context.device.probe.model

From 1a80e2d8992bb5a35716b5842755832098afdca6 Mon Sep 17 00:00:00 2001
From: Piotr Jarosik <piotr.jarosik@us4us.eu>
Date: Sun, 14 Nov 2021 12:37:45 +0100
Subject: [PATCH 06/21] Added adapter settings for ATL/Phiplips-us4r4 adapter.

---
 arrus/core/io/test-data/dictionary.prototxt | 41 +++++++++++++++++++++
 1 file changed, 41 insertions(+)

diff --git a/arrus/core/io/test-data/dictionary.prototxt b/arrus/core/io/test-data/dictionary.prototxt
index a88a676fe..93bb3f65a 100644
--- a/arrus/core/io/test-data/dictionary.prototxt
+++ b/arrus/core/io/test-data/dictionary.prototxt
@@ -225,6 +225,47 @@ probe_adapter_models: [
 			   120, 121, 122, 123, 124, 125, 126, 127]
             }
         ]
+    },
+    {
+        id: {
+            manufacturer: "us4us"
+            name: "atl/philips-us4r4"
+        }
+        n_channels: 128
+        channel_mapping_regions: [
+        {
+            # BLUE, A
+            us4oem: 0
+            channels: [31, 30, 29, 28, 27, 26, 25, 24,
+                       23, 22, 21, 20, 19, 18, 17, 15,
+                       16, 14, 13, 12, 11, 10, 9,  8,
+                       7,  6,  5,  4,  3,  2,  1,  0]
+        },
+        {
+            # BLUE, C
+            us4oem: 1
+            channels: [31, 30, 29, 28, 27, 26, 25, 24,
+                       23, 22, 21, 20, 19, 18, 17, 15,
+                       16, 14, 13, 12, 11, 10, 9,  8,
+                       7,  6,  5,  4,  3,  2,  1,  0]
+        },
+        {
+            # GREEN, H
+            us4oem: 2
+            channels: [31, 30, 29, 28, 27, 26, 25, 24,
+                       23, 22, 21, 20, 19, 18, 17, 15,
+                       16, 14, 13, 12, 11, 10, 9,  8,
+                       7,  6,  5,  4,  3,  2,  1,  0]
+            },
+            {
+            # GREEN, F
+            us4oem: 3
+            channels: [31, 30, 29, 28, 27, 26, 25, 24,
+                       23, 22, 21, 20, 19, 18, 17, 15,
+                       16, 14, 13, 12, 11, 10, 9,  8,
+                       7,  6,  5,  4,  3,  2,  1,  0]
+            }
+         ]
     }
 ]
 probe_models: [

From 562f9197812c271815f7579c5804dc58d37b28af Mon Sep 17 00:00:00 2001
From: Piotr Jarosik <piotr.jarosik@us4us.eu>
Date: Sun, 14 Nov 2021 16:31:27 +0100
Subject: [PATCH 07/21] Added adapter settings for ATL/Phiplips-us4r4 adapter.

---
 arrus/cfg/default.dict                      | 45 +++++++++++++++++++++
 arrus/core/io/test-data/dictionary.prototxt | 41 -------------------
 2 files changed, 45 insertions(+), 41 deletions(-)

diff --git a/arrus/cfg/default.dict b/arrus/cfg/default.dict
index ae46f3409..9bcb7a64e 100644
--- a/arrus/cfg/default.dict
+++ b/arrus/cfg/default.dict
@@ -310,6 +310,47 @@ probe_adapter_models: [
 			   120, 121, 122, 123, 124, 125, 126, 127]
             }
         ]
+    },
+    {
+        id: {
+            manufacturer: "us4us"
+            name: "atl/philips-us4r4"
+        }
+        n_channels: 128
+        channel_mapping_regions: [
+        {
+            # BLUE, A
+            us4oem: 0
+            channels: [31, 30, 29, 28, 27, 26, 25, 24,
+                       23, 22, 21, 20, 19, 18, 17, 15,
+                       16, 14, 13, 12, 11, 10, 9,  8,
+                       7,  6,  5,  4,  3,  2,  1,  0]
+        },
+        {
+            # BLUE, C
+            us4oem: 1
+            channels: [31, 30, 29, 28, 27, 26, 25, 24,
+                       23, 22, 21, 20, 19, 18, 17, 15,
+                       16, 14, 13, 12, 11, 10, 9,  8,
+                       7,  6,  5,  4,  3,  2,  1,  0]
+        },
+        {
+            # GREEN, H
+            us4oem: 2
+            channels: [31, 30, 29, 28, 27, 26, 25, 24,
+                       23, 22, 21, 20, 19, 18, 17, 15,
+                       16, 14, 13, 12, 11, 10, 9,  8,
+                       7,  6,  5,  4,  3,  2,  1,  0]
+        },
+        {
+            # GREEN, F
+            us4oem: 3
+            channels: [31, 30, 29, 28, 27, 26, 25, 24,
+                       23, 22, 21, 20, 19, 18, 17, 15,
+                       16, 14, 13, 12, 11, 10, 9,  8,
+                       7,  6,  5,  4,  3,  2,  1,  0]
+        }
+        ]
     }
 ]
 probe_models: [
@@ -721,6 +762,10 @@ probe_to_adapter_connections: [
             {
                 manufacturer: "us4us"
                 name: "atl/philips"
+            },
+            {
+                manufacturer: "us4us"
+                name: "atl/philips-us4r4"
             }
         ]
         channel_mapping_ranges: {
diff --git a/arrus/core/io/test-data/dictionary.prototxt b/arrus/core/io/test-data/dictionary.prototxt
index 93bb3f65a..a88a676fe 100644
--- a/arrus/core/io/test-data/dictionary.prototxt
+++ b/arrus/core/io/test-data/dictionary.prototxt
@@ -225,47 +225,6 @@ probe_adapter_models: [
 			   120, 121, 122, 123, 124, 125, 126, 127]
             }
         ]
-    },
-    {
-        id: {
-            manufacturer: "us4us"
-            name: "atl/philips-us4r4"
-        }
-        n_channels: 128
-        channel_mapping_regions: [
-        {
-            # BLUE, A
-            us4oem: 0
-            channels: [31, 30, 29, 28, 27, 26, 25, 24,
-                       23, 22, 21, 20, 19, 18, 17, 15,
-                       16, 14, 13, 12, 11, 10, 9,  8,
-                       7,  6,  5,  4,  3,  2,  1,  0]
-        },
-        {
-            # BLUE, C
-            us4oem: 1
-            channels: [31, 30, 29, 28, 27, 26, 25, 24,
-                       23, 22, 21, 20, 19, 18, 17, 15,
-                       16, 14, 13, 12, 11, 10, 9,  8,
-                       7,  6,  5,  4,  3,  2,  1,  0]
-        },
-        {
-            # GREEN, H
-            us4oem: 2
-            channels: [31, 30, 29, 28, 27, 26, 25, 24,
-                       23, 22, 21, 20, 19, 18, 17, 15,
-                       16, 14, 13, 12, 11, 10, 9,  8,
-                       7,  6,  5,  4,  3,  2,  1,  0]
-            },
-            {
-            # GREEN, F
-            us4oem: 3
-            channels: [31, 30, 29, 28, 27, 26, 25, 24,
-                       23, 22, 21, 20, 19, 18, 17, 15,
-                       16, 14, 13, 12, 11, 10, 9,  8,
-                       7,  6,  5,  4,  3,  2,  1,  0]
-            }
-         ]
     }
 ]
 probe_models: [

From 118e66a9477b0758416634140d1384a8a72e9ea5 Mon Sep 17 00:00:00 2001
From: Piotr Jarosik <piotr.jarosik@us4us.eu>
Date: Mon, 15 Nov 2021 09:26:19 +0100
Subject: [PATCH 08/21] Version 0.6.3.

---
 CMakeLists.txt | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/CMakeLists.txt b/CMakeLists.txt
index b5d3c5863..ef252b70d 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -3,7 +3,7 @@ cmake_minimum_required(VERSION 3.17.0)
 project(arrus LANGUAGES CXX)
 set(CMAKE_CXX_STANDARD 17)
 # Version
-set(PROJECT_VERSION 0.6.2)
+set(PROJECT_VERSION 0.6.3)
 set(ARRUS_PROJECT_VERSION "${PROJECT_VERSION}")
 
 option(ARRUS_DEVELOP_VERSION "Build develop version." ON)

From 00939ca3ce854697a0800ed0e037056a30a95e86 Mon Sep 17 00:00:00 2001
From: Piotr Jarosik <piotr.jarosik@us4us.eu>
Date: Fri, 19 Nov 2021 09:18:41 +0100
Subject: [PATCH 09/21] Added an example for phased array scanning.

---
 api/python/examples/phased_array_scanning.py | 77 ++++++++++++++++++++
 1 file changed, 77 insertions(+)
 create mode 100644 api/python/examples/phased_array_scanning.py

diff --git a/api/python/examples/phased_array_scanning.py b/api/python/examples/phased_array_scanning.py
new file mode 100644
index 000000000..69732f006
--- /dev/null
+++ b/api/python/examples/phased_array_scanning.py
@@ -0,0 +1,77 @@
+"""
+This script acquires and reconstructs RF img for classical imaging scheme
+(scanline by scanline).
+
+GPU is required.
+"""
+import numpy as np
+import arrus
+from arrus.ops.us4r import Scheme, Pulse
+from arrus.ops.imaging import LinSequence
+from arrus.utils.gui import Display2D
+from arrus.utils.imaging import get_bmode_imaging, get_extent
+import pickle
+from arrus.utils.imaging import *
+from collections import deque
+
+arrus.set_clog_level(arrus.logging.INFO)
+arrus.add_log_file("test.log", arrus.logging.INFO)
+
+# Here starts communication with the device.
+with arrus.Session() as sess:
+    us4r = sess.get_device("/Us4R:0")
+    us4r.set_hv_voltage(5)
+
+    sequence = LinSequence(
+        tx_aperture_center=0,
+        tx_aperture_size=96,
+        tx_focus=20e-3,
+        angles=np.linspace(-30, 30, 128)*np.pi/180,
+        pulse=Pulse(center_frequency=2.5e6, n_periods=2, inverse=False),
+        rx_aperture_center=0,
+        rx_aperture_size=64,
+        rx_sample_range=(0, 4096),
+        pri=200e-6,
+        tgc_start=14,
+        tgc_slope=2e2,
+        downsampling_factor=1,
+        speed_of_sound=1450)
+
+    # Imaging output grid.
+    x_grid = np.arange(-18, 18, 0.2) * 1e-3
+    z_grid = np.arange(0, 45, 0.2) * 1e-3
+
+    rf_queue = deque(maxlen=2)
+
+    scheme = Scheme(
+        tx_rx_sequence=sequence,
+        processing=Pipeline(
+            steps=(
+                RemapToLogicalOrder(),
+                Pipeline(
+                    steps=(Lambda(lambda data: (rf_queue.append(data), data)[1]), ),
+                    placement="/GPU:0",
+                ),
+                Transpose(axes=(0, 2, 1)),
+                BandpassFilter(),
+                QuadratureDemodulation(),
+                Decimation(decimation_factor=4, cic_order=2),
+                RxBeamforming(),
+                EnvelopeDetection(),
+                Transpose(),
+                ScanConversion(x_grid, z_grid),
+                LogCompression()
+            ),
+            placement="/GPU:0")
+    )
+    # Upload sequence on the us4r-lite device.
+    buffer, (bmode_metadata, rf_metadata) = sess.upload(scheme)
+    display = Display2D(metadata=bmode_metadata, value_range=(10, 80), cmap="gray",
+                        title="B-mode", xlabel="OX (mm)", ylabel="OZ (mm)",
+                        extent=get_extent(x_grid, z_grid)*1e3,
+                        show_colorbar=True)
+    sess.start_scheme()
+    display.start(buffer)
+
+# When we exit the above scope, the session and scheme is properly closed.
+print("Stopping the example.")

From bcd4fde0d902c3b3a35a35e3e962f12e684704a9 Mon Sep 17 00:00:00 2001
From: Piotr Jarosik <piotr.jarosik@us4us.eu>
Date: Fri, 19 Nov 2021 15:56:51 +0100
Subject: [PATCH 10/21] Fixed the first scanline position in the ScanConverter
 implementation.

---
 api/python/arrus/utils/imaging.py | 5 +++--
 1 file changed, 3 insertions(+), 2 deletions(-)

diff --git a/api/python/arrus/utils/imaging.py b/api/python/arrus/utils/imaging.py
index e7b42e48a..122158d72 100644
--- a/api/python/arrus/utils/imaging.py
+++ b/api/python/arrus/utils/imaging.py
@@ -914,7 +914,7 @@ def prepare(self, const_metadata: arrus.metadata.ConstMetadata):
         elif not seq.init_delay == "tx_center":
             raise ValueError(f"Unrecognized init_delay value: {initial_delay}")
         radial_distance = (
-                (start_sample / acq_fs + np.arange(0, self.n_samples) / fs)* c/2
+                (start_sample / acq_fs + np.arange(0, self.n_samples) / fs) * c/2
         )
         x_distance = (radial_distance * np.sin(tx_angle)).reshape(1, -1)
         z_distance = radial_distance * np.cos(tx_angle).reshape(1, -1)
@@ -1128,13 +1128,14 @@ def _prepare_linear_array(self, const_metadata: arrus.metadata.ConstMetadata):
         tx_center_diff = tx_center_diff[0]
         # Determine input grid.
         input_x_grid_diff = tx_center_diff*pitch
-        input_x_grid_origin = tx_aperture_center_element[0]-(n_elements-1)/2*pitch
+        input_x_grid_origin = (tx_aperture_center_element[0]-(n_elements-1)/2)*pitch
         acq_fs = (const_metadata.context.device.sampling_frequency
                   / seq.downsampling_factor)
         fs = data_desc.sampling_frequency
         start_sample = seq.rx_sample_range[0]
         input_z_grid_origin = start_sample/acq_fs*c/2
         input_z_grid_diff = c/(fs*2)
+        # Map x_grid and z_grid to the RF frame coordinates.
         interp_x_grid = (self.x_grid-input_x_grid_origin)/input_x_grid_diff
         interp_z_grid = (self.z_grid-input_z_grid_origin)/input_z_grid_diff
         self._interp_mesh = cp.asarray(np.meshgrid(interp_z_grid, interp_x_grid, indexing="ij"))

From a657eba418a1ec867346d1c521cd1657b1f40e54 Mon Sep 17 00:00:00 2001
From: Piotr Jarosik <jarosik.piotr@gmail.com>
Date: Sat, 20 Nov 2021 08:35:49 +0100
Subject: [PATCH 11/21] Added phased array example to the arrus package,
 updated release notes

---
 api/python/CMakeLists.txt           |  1 +
 docs/content/misc/release_notes.rst | 18 ++++++++++++++++++
 2 files changed, 19 insertions(+)

diff --git a/api/python/CMakeLists.txt b/api/python/CMakeLists.txt
index 93e530dfa..8c6d33e52 100644
--- a/api/python/CMakeLists.txt
+++ b/api/python/CMakeLists.txt
@@ -198,6 +198,7 @@ install(
         examples/diverging_beams.py
         examples/custom_tx_rx_sequence.py
         examples/custom_callback.py
+	examples/phased_array_scanning.py
         examples/requirements.txt
     examples/requirements.txt
         DESTINATION
diff --git a/docs/content/misc/release_notes.rst b/docs/content/misc/release_notes.rst
index a454d001f..eaaa5a6c6 100644
--- a/docs/content/misc/release_notes.rst
+++ b/docs/content/misc/release_notes.rst
@@ -4,6 +4,24 @@ Release notes
 0.6.x
 -----
 
+
+0.6.4
+.....
+
+- Python API:
+
+    - Fixed linear scanning for tx apertures starting from channels > 0.
+
+0.6.3
+.....
+
+- Python API: 
+
+    - Added phased array scanning & example.
+    - Added definition for the probe adapter atl/philips-us4r4.
+    - Improved IQ raw to LRI CUDA kernel performance.
+    - Increased the maximum allowable voltage for Esaote probes to 90 V.
+
 0.6.2
 `````
 

From a7044557b05c516237cc167e861e40279ceb3b19 Mon Sep 17 00:00:00 2001
From: Piotr Jarosik <jarosik.piotr@gmail.com>
Date: Sat, 20 Nov 2021 08:36:53 +0100
Subject: [PATCH 12/21] Updated release notes.

---
 docs/content/misc/release_notes.rst | 3 ++-
 1 file changed, 2 insertions(+), 1 deletion(-)

diff --git a/docs/content/misc/release_notes.rst b/docs/content/misc/release_notes.rst
index eaaa5a6c6..1143f673a 100644
--- a/docs/content/misc/release_notes.rst
+++ b/docs/content/misc/release_notes.rst
@@ -11,13 +11,14 @@ Release notes
 - Python API:
 
     - Fixed linear scanning for tx apertures starting from channels > 0.
+    - Added phased array scanning example.  
 
 0.6.3
 .....
 
 - Python API: 
 
-    - Added phased array scanning & example.
+    - Added phased array scanning.
     - Added definition for the probe adapter atl/philips-us4r4.
     - Improved IQ raw to LRI CUDA kernel performance.
     - Increased the maximum allowable voltage for Esaote probes to 90 V.

From c723949e9add92143ebc924a014c7fa5ec56f1ed Mon Sep 17 00:00:00 2001
From: Piotr Jarosik <jarosik.piotr@gmail.com>
Date: Sat, 20 Nov 2021 21:41:04 +0100
Subject: [PATCH 13/21] Version 0.6.4

---
 CMakeLists.txt | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/CMakeLists.txt b/CMakeLists.txt
index ef252b70d..e03855daa 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -3,7 +3,7 @@ cmake_minimum_required(VERSION 3.17.0)
 project(arrus LANGUAGES CXX)
 set(CMAKE_CXX_STANDARD 17)
 # Version
-set(PROJECT_VERSION 0.6.3)
+set(PROJECT_VERSION 0.6.4)
 set(ARRUS_PROJECT_VERSION "${PROJECT_VERSION}")
 
 option(ARRUS_DEVELOP_VERSION "Build develop version." ON)

From 102c0dbd38163230dbb7f7f8711e58a7b5a7903d Mon Sep 17 00:00:00 2001
From: Piotr Jarosik <jarosik.piotr@gmail.com>
Date: Mon, 6 Dec 2021 22:11:36 +0100
Subject: [PATCH 14/21] Added CMake for option python-dev lib version.

---
 CMakeLists.txt            | 1 +
 api/python/CMakeLists.txt | 4 ++--
 2 files changed, 3 insertions(+), 2 deletions(-)

diff --git a/CMakeLists.txt b/CMakeLists.txt
index e03855daa..9f6a291dc 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -38,6 +38,7 @@ set(ARRUS_PYTHON_INSTALL_DIR python)
 set(Us4_LIB_DIR ${Us4_ROOT_DIR}/lib64)
 
 option(ARRUS_BUILD_PY "Build python API." OFF)
+option(ARRUS_PY_VERSION "Python version for which the ARRUS package should be built." 3.8)
 option(ARRUS_BUILD_MATLAB "Build MATLAB API." OFF)
 option(ARRUS_BUILD_DOCS "Build documentation." OFF)
 option(ARRUS_RUN_TESTS "Run all tests builded packages." OFF)
diff --git a/api/python/CMakeLists.txt b/api/python/CMakeLists.txt
index 8c6d33e52..93a31a5b7 100644
--- a/api/python/CMakeLists.txt
+++ b/api/python/CMakeLists.txt
@@ -5,8 +5,8 @@
 ################################################################################
 find_package(SWIG REQUIRED)
 include(UseSWIG)
-find_package(PythonInterp REQUIRED)
-find_package(PythonLibs REQUIRED)
+find_package(PythonInterp ${ARRUS_PY_VERSION} REQUIRED)
+find_package(PythonLibs ${ARRUS_PY_VERSION} REQUIRED)
 set_property(SOURCE wrappers/core.i PROPERTY CPLUSPLUS ON)
 if (MSVC)
     set_property(

From 491d4101e2ddeb2d6184bf7fd1c31a6d26260dc6 Mon Sep 17 00:00:00 2001
From: Piotr Jarosik <piotr.jarosik@us4us.eu>
Date: Thu, 16 Dec 2021 18:01:03 +0100
Subject: [PATCH 15/21] Unlock page-locked output buffer when the scheme is
 stopped.

---
 CMakeLists.txt                                |  2 +-
 arrus/core/devices/probe/ProbeImpl.cpp        |  3 ++
 arrus/core/devices/probe/ProbeImpl.h          |  2 +
 arrus/core/devices/probe/ProbeImplBase.h      |  1 +
 arrus/core/devices/us4r/Us4RImpl.cpp          |  5 +-
 arrus/core/devices/us4r/Us4RImpl.h            |  1 +
 .../us4r/probeadapter/ProbeAdapterImpl.cpp    | 53 ++++++++++++++-----
 .../us4r/probeadapter/ProbeAdapterImpl.h      |  4 ++
 .../us4r/probeadapter/ProbeAdapterImplBase.h  |  6 +--
 9 files changed, 60 insertions(+), 17 deletions(-)

diff --git a/CMakeLists.txt b/CMakeLists.txt
index 9f6a291dc..3885dc468 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -38,7 +38,7 @@ set(ARRUS_PYTHON_INSTALL_DIR python)
 set(Us4_LIB_DIR ${Us4_ROOT_DIR}/lib64)
 
 option(ARRUS_BUILD_PY "Build python API." OFF)
-option(ARRUS_PY_VERSION "Python version for which the ARRUS package should be built." 3.8)
+set(ARRUS_PY_VERSION "3.8" CACHE STRING "Python version for which the ARRUS package should be built.")
 option(ARRUS_BUILD_MATLAB "Build MATLAB API." OFF)
 option(ARRUS_BUILD_DOCS "Build documentation." OFF)
 option(ARRUS_RUN_TESTS "Run all tests builded packages." OFF)
diff --git a/arrus/core/devices/probe/ProbeImpl.cpp b/arrus/core/devices/probe/ProbeImpl.cpp
index 982c2aafe..71323b3b5 100644
--- a/arrus/core/devices/probe/ProbeImpl.cpp
+++ b/arrus/core/devices/probe/ProbeImpl.cpp
@@ -142,6 +142,9 @@ void ProbeImpl::registerOutputBuffer(Us4ROutputBuffer *buffer,
                                      bool isTriggerSync) {
     adapter->registerOutputBuffer(buffer, us4rBuffer, isTriggerSync);
 }
+void ProbeImpl::unregisterOutputBuffer(Us4ROutputBuffer *buffer, const Us4RBuffer::Handle &us4RBuffer) {
+    adapter->unregisterOutputBuffer(buffer, us4RBuffer);
+}
 
 // Remaps FCM according to given rx aperture active channels mappings.
 FrameChannelMapping::Handle ProbeImpl::remapFcm(
diff --git a/arrus/core/devices/probe/ProbeImpl.h b/arrus/core/devices/probe/ProbeImpl.h
index bae565554..d8bd25af2 100644
--- a/arrus/core/devices/probe/ProbeImpl.h
+++ b/arrus/core/devices/probe/ProbeImpl.h
@@ -38,6 +38,8 @@ class ProbeImpl : public ProbeImplBase {
 
     void registerOutputBuffer(Us4ROutputBuffer *buffer, const Us4RBuffer::Handle &us4rBuffer, bool isTriggerSync) override;
 
+    void unregisterOutputBuffer(Us4ROutputBuffer *buffer, const Us4RBuffer::Handle &handle) override;
+
     static FrameChannelMapping::Handle remapFcm(const FrameChannelMapping::Handle &adapterFcm,
                                                 const std::vector<std::vector<ChannelIdx>> &adapterActiveChannels,
                                                 const std::vector<ChannelIdx> &rxPaddingLeft,
diff --git a/arrus/core/devices/probe/ProbeImplBase.h b/arrus/core/devices/probe/ProbeImplBase.h
index 4e2a9914d..d5da3d748 100644
--- a/arrus/core/devices/probe/ProbeImplBase.h
+++ b/arrus/core/devices/probe/ProbeImplBase.h
@@ -19,6 +19,7 @@ class ProbeImplBase : public Probe, public UltrasoundDevice {
                     uint16 rxBatchSize, std::optional<float> sri, bool triggerSync) = 0;
 
     virtual void registerOutputBuffer(Us4ROutputBuffer *, const Us4RBuffer::Handle &, bool isTriggerSync) = 0;
+    virtual void unregisterOutputBuffer(Us4ROutputBuffer *, const Us4RBuffer::Handle &) = 0;
 };
 
 }
diff --git a/arrus/core/devices/us4r/Us4RImpl.cpp b/arrus/core/devices/us4r/Us4RImpl.cpp
index f5035ebc2..29b834c6d 100644
--- a/arrus/core/devices/us4r/Us4RImpl.cpp
+++ b/arrus/core/devices/us4r/Us4RImpl.cpp
@@ -111,11 +111,13 @@ Us4RImpl::upload(const ops::us4r::TxRxSequence &seq,
     auto dataType = element.getDataType();
     // If the output buffer already exists - remove it.
     if (this->buffer) {
+        // The buffer should be already unregistered (after stopping the device).
         this->buffer.reset();
     }
     // Create output buffer.
     this->buffer = std::make_shared<Us4ROutputBuffer>(us4oemComponentSize, shape, dataType, hostBufferNElements);
-    getProbeImpl()->registerOutputBuffer(this->buffer.get(), rxBuffer, isTriggerSync);
+    this->us4rBuffer = std::move(rxBuffer);
+    getProbeImpl()->registerOutputBuffer(this->buffer.get(), this->us4rBuffer, isTriggerSync);
     return {this->buffer, std::move(fcm)};
 }
 
@@ -153,6 +155,7 @@ void Us4RImpl::stopDevice() {
     if (this->buffer != nullptr) {
         this->buffer->shutdown();
         std::this_thread::sleep_for(std::chrono::milliseconds(1000));
+        getProbeImpl()->unregisterOutputBuffer(this->buffer.get(), this->us4rBuffer);
     }
     this->state = State::STOPPED;
 }
diff --git a/arrus/core/devices/us4r/Us4RImpl.h b/arrus/core/devices/us4r/Us4RImpl.h
index ea815b16d..abac6ad95 100644
--- a/arrus/core/devices/us4r/Us4RImpl.h
+++ b/arrus/core/devices/us4r/Us4RImpl.h
@@ -142,6 +142,7 @@ class Us4RImpl : public Us4R {
     std::optional<ProbeAdapterImplBase::Handle> probeAdapter;
     std::optional<ProbeImplBase::Handle> probe;
     std::optional<HighVoltageSupplier::Handle> hv;
+    std::unique_ptr<Us4RBuffer> us4rBuffer;
     std::shared_ptr<Us4ROutputBuffer> buffer;
     State state{State::STOPPED};
     // AFE parameters.
diff --git a/arrus/core/devices/us4r/probeadapter/ProbeAdapterImpl.cpp b/arrus/core/devices/us4r/probeadapter/ProbeAdapterImpl.cpp
index 8b6d7392c..51a300feb 100644
--- a/arrus/core/devices/us4r/probeadapter/ProbeAdapterImpl.cpp
+++ b/arrus/core/devices/us4r/probeadapter/ProbeAdapterImpl.cpp
@@ -234,9 +234,7 @@ ProbeAdapterImpl::setTxRxSequence(const std::vector<TxRxParameters> &seq,
                 // and has no assigned value.
                 ARRUS_REQUIRES_DATA_TYPE_E(
                     dstModuleChannel, int8,
-                    ::arrus::ArrusException(
-                        "Invalid dstModuleChannel data type, "
-                        "rx aperture is outise."));
+                    ::arrus::ArrusException("Invalid dstModuleChannel data type, rx aperture is outside."));
                 if (FrameChannelMapping::isChannelUnavailable((int8)dstModuleChannel)) {
                     outFcBuilder.setChannelMapping(
                         frameIdx, activeRxChIdx + op.getRxPadding()[0],
@@ -305,15 +303,7 @@ void ProbeAdapterImpl::registerOutputBuffer(Us4ROutputBuffer *buffer, const Us4R
 void ProbeAdapterImpl::registerOutputBuffer(Us4ROutputBuffer *outputBuffer, const Us4OEMBuffer &us4oemBuffer,
                                             Us4OEMImplBase::RawHandle us4oem, bool isTriggerSync) {
     // Each transfer should have the same size.
-    std::unordered_set<size_t> sizes;
-    for (auto &element: us4oemBuffer.getElements()) {
-        sizes.insert(element.getSize());
-    }
-    if (sizes.size() > 1) {
-        throw ::arrus::ArrusException("Each us4oem buffer element should have the same size.");
-    }
-    // This is the size of a single element produced by this us4oem.
-    const size_t elementSize = *std::begin(sizes);
+    size_t elementSize = getUniqueUs4OEMBufferElementSize(us4oemBuffer);
     if (elementSize == 0) {
         // This us4oem will not transfer any data, so the buffer registration has no sense here.
         return;
@@ -449,4 +439,43 @@ void ProbeAdapterImpl::registerOutputBuffer(Us4ROutputBuffer *outputBuffer, cons
     });
 }
 
+size_t ProbeAdapterImpl::getUniqueUs4OEMBufferElementSize(const Us4OEMBuffer &us4oemBuffer) const {
+    std::unordered_set<size_t> sizes;
+    for (auto &element: us4oemBuffer.getElements()) {
+        sizes.insert(element.getSize());
+    }
+    if (sizes.size() > 1) {
+        throw ArrusException("Each us4oem buffer element should have the same size.");
+    }
+    // This is the size of a single element produced by this us4oem.
+    const size_t elementSize = *std::begin(sizes);
+    return elementSize;
+}
+
+void ProbeAdapterImpl::unregisterOutputBuffer(Us4ROutputBuffer *hostBuffer, const Us4RBuffer::Handle &us4rBuffer) {
+    const size_t hostBufferNElements = hostBuffer->getNumberOfElements();
+
+    for (Ordinal i = 0; i < us4oems.size(); ++i) {
+        auto &us4oem = us4oems[i];
+        const Ordinal ordinal = us4oem->getDeviceId().getOrdinal();
+        auto ius4oem = us4oem->getIUs4oem();
+
+        auto us4oemBuffer = us4rBuffer->getUs4oemBuffer(i);
+        size_t elementSize = getUniqueUs4OEMBufferElementSize(us4oemBuffer);
+        const auto rxBufferNElements = ARRUS_SAFE_CAST(us4oemBuffer.getNumberOfElements(), uint16);
+        uint16 hostElement = 0, rxElement = 0;
+
+        while (hostElement < hostBufferNElements) {
+            auto dstAddress = hostBuffer->getAddress(hostElement, ordinal);
+            auto srcAddress = us4oemBuffer.getElement(rxElement).getAddress();
+            logger->log(LogSeverity::DEBUG,
+                        format("Unregistering transfer: to {} from {}, size {}",
+                                                            (size_t)dstAddress, (size_t)srcAddress, elementSize));
+            ius4oem->ReleaseTransferRxBufferToHost(dstAddress, elementSize, srcAddress);
+            ++hostElement;
+            rxElement = (rxElement + 1) % rxBufferNElements;
+        }
+    }
+}
+
 }
\ No newline at end of file
diff --git a/arrus/core/devices/us4r/probeadapter/ProbeAdapterImpl.h b/arrus/core/devices/us4r/probeadapter/ProbeAdapterImpl.h
index 42c32c4bc..0f13df505 100644
--- a/arrus/core/devices/us4r/probeadapter/ProbeAdapterImpl.h
+++ b/arrus/core/devices/us4r/probeadapter/ProbeAdapterImpl.h
@@ -50,6 +50,8 @@ class ProbeAdapterImpl : public ProbeAdapterImplBase {
                          const Us4RBuffer::Handle &us4rBuffer,
                          bool isTriggerSync);
 
+    void unregisterOutputBuffer(Us4ROutputBuffer *buffer, const Us4RBuffer::Handle &us4rBuffer);
+
 private:
     Logger::Handle logger;
     ProbeAdapterModelId modelId;
@@ -65,6 +67,8 @@ class ProbeAdapterImpl : public ProbeAdapterImplBase {
     Us4OEMImplBase::RawHandle getMasterUs4oem() const {
         return this->us4oems[0];
     }
+
+    size_t getUniqueUs4OEMBufferElementSize(const Us4OEMBuffer &us4oemBuffer) const;
 };
 
 }
diff --git a/arrus/core/devices/us4r/probeadapter/ProbeAdapterImplBase.h b/arrus/core/devices/us4r/probeadapter/ProbeAdapterImplBase.h
index 4ea6509cb..3f730a154 100644
--- a/arrus/core/devices/us4r/probeadapter/ProbeAdapterImplBase.h
+++ b/arrus/core/devices/us4r/probeadapter/ProbeAdapterImplBase.h
@@ -26,9 +26,9 @@ class ProbeAdapterImplBase : public ProbeAdapter {
                     bool triggerSync) = 0;
 
     virtual
-    void registerOutputBuffer(Us4ROutputBuffer *buffer,
-                              const Us4RBuffer::Handle &transfers,
-                              bool isTriggerSync) = 0;
+    void registerOutputBuffer(Us4ROutputBuffer *buffer, const Us4RBuffer::Handle &transfers, bool isTriggerSync) = 0;
+
+    virtual void unregisterOutputBuffer(Us4ROutputBuffer *buffer, const Us4RBuffer::Handle &handle) = 0;
 
     virtual Ordinal getNumberOfUs4OEMs() = 0;
 

From f029a3d88fa6717d893c43a4bf186279a428cf29 Mon Sep 17 00:00:00 2001
From: Piotr Jarosik <piotr.jarosik@us4us.eu>
Date: Fri, 17 Dec 2021 09:19:47 +0100
Subject: [PATCH 16/21] Us4R buffer unregistration will be performed only once.

---
 arrus/core/devices/us4r/Us4RImpl.cpp | 5 ++++-
 1 file changed, 4 insertions(+), 1 deletion(-)

diff --git a/arrus/core/devices/us4r/Us4RImpl.cpp b/arrus/core/devices/us4r/Us4RImpl.cpp
index 29b834c6d..a1ca547cb 100644
--- a/arrus/core/devices/us4r/Us4RImpl.cpp
+++ b/arrus/core/devices/us4r/Us4RImpl.cpp
@@ -155,7 +155,10 @@ void Us4RImpl::stopDevice() {
     if (this->buffer != nullptr) {
         this->buffer->shutdown();
         std::this_thread::sleep_for(std::chrono::milliseconds(1000));
-        getProbeImpl()->unregisterOutputBuffer(this->buffer.get(), this->us4rBuffer);
+        if(this->us4rBuffer) {
+            getProbeImpl()->unregisterOutputBuffer(this->buffer.get(), this->us4rBuffer);
+            this->us4rBuffer.reset();
+        }
     }
     this->state = State::STOPPED;
 }

From f1c9a5e164d26ad3fef10124c9c13325c65c6c5c Mon Sep 17 00:00:00 2001
From: Piotr Jarosik <piotr.jarosik@us4us.eu>
Date: Fri, 17 Dec 2021 09:25:55 +0100
Subject: [PATCH 17/21] Version 0.6.5.

---
 CMakeLists.txt | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/CMakeLists.txt b/CMakeLists.txt
index 3885dc468..25811e3fd 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -3,7 +3,7 @@ cmake_minimum_required(VERSION 3.17.0)
 project(arrus LANGUAGES CXX)
 set(CMAKE_CXX_STANDARD 17)
 # Version
-set(PROJECT_VERSION 0.6.4)
+set(PROJECT_VERSION 0.6.5)
 set(ARRUS_PROJECT_VERSION "${PROJECT_VERSION}")
 
 option(ARRUS_DEVELOP_VERSION "Build develop version." ON)

From ca03f79be85b0e01c35518e87c9223e761ea2724 Mon Sep 17 00:00:00 2001
From: Piotr Jarosik <piotr.jarosik@us4us.eu>
Date: Fri, 17 Dec 2021 09:29:59 +0100
Subject: [PATCH 18/21] Updated release notes.

---
 docs/content/misc/release_notes.rst | 9 +++++++++
 1 file changed, 9 insertions(+)

diff --git a/docs/content/misc/release_notes.rst b/docs/content/misc/release_notes.rst
index 1143f673a..b4d185c0d 100644
--- a/docs/content/misc/release_notes.rst
+++ b/docs/content/misc/release_notes.rst
@@ -5,6 +5,15 @@ Release notes
 -----
 
 
+0.6.5
+.....
+
+- core (C++ API):
+
+    - fixed memory leak on subsequent re-uploads
+    - some improvements in the us4R-lite driver compatibility with the us4R-lite system
+
+
 0.6.4
 .....
 

From 15a06d33d4859ec0fb294e9e7fff814ad2c228bd Mon Sep 17 00:00:00 2001
From: Piotr Jarosik <piotr.jarosik@us4us.eu>
Date: Fri, 17 Dec 2021 09:51:56 +0100
Subject: [PATCH 19/21] Increased us4r-api driver dependency to 0.6.9

---
 CMakeLists.txt | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/CMakeLists.txt b/CMakeLists.txt
index 25811e3fd..f0fb7dd6a 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -88,7 +88,7 @@ set(ARRUS_DOCS_INSTALL_DIR docs)
 ################################################################################
 # Common dependencies
 ################################################################################
-find_package(Us4 0.6.8 EXACT REQUIRED US4OEM HV256 DBARLite)
+find_package(Us4 0.6.9 EXACT REQUIRED US4OEM HV256 DBARLite)
 find_package(Boost REQUIRED)
 set(Boost_USE_STATIC_LIBS ON)
 find_package(Protobuf REQUIRED)

From 157e93a7374468538414a8ac2c877cf1581f3df3 Mon Sep 17 00:00:00 2001
From: Piotr Jarosik <piotr.jarosik@us4us.eu>
Date: Fri, 17 Dec 2021 09:55:46 +0100
Subject: [PATCH 20/21] Removed unreachable code.

---
 arrus/core/io/SettingsDictionary.h | 1 -
 1 file changed, 1 deletion(-)

diff --git a/arrus/core/io/SettingsDictionary.h b/arrus/core/io/SettingsDictionary.h
index 5aec62bc9..9a39cfda3 100644
--- a/arrus/core/io/SettingsDictionary.h
+++ b/arrus/core/io/SettingsDictionary.h
@@ -45,7 +45,6 @@ class SettingsDictionary {
                 probeModelId.getManufacturer(), probeModelId.getName(),
                 adapterModelId.getManufacturer(), adapterModelId.getName()));
         }
-        return probesMap.at(key);
     }
 
     void insertProbeSettings(ProbeSettings &&probe,

From 3a486887264aa81e84b6d2ed1ac365cb155d968c Mon Sep 17 00:00:00 2001
From: Piotr Jarosik <piotr.jarosik@us4us.eu>
Date: Fri, 17 Dec 2021 17:51:30 +0100
Subject: [PATCH 21/21] ARRUS-116: Exposed the number of us4oems and FPGA
 temperature in the us4r API. (#249)

---
 arrus/core/api/devices/us4r/Us4OEM.h                   |  2 ++
 arrus/core/api/devices/us4r/Us4R.h                     | 10 ++++++++++
 arrus/core/devices/us4r/Us4RImpl.cpp                   |  8 ++++++++
 arrus/core/devices/us4r/Us4RImpl.h                     |  2 ++
 .../devices/us4r/probeadapter/ProbeAdapterImplTest.cpp |  1 +
 arrus/core/devices/us4r/us4oem/Us4OEMImpl.cpp          |  4 +++-
 arrus/core/devices/us4r/us4oem/Us4OEMImpl.h            |  1 +
 7 files changed, 27 insertions(+), 1 deletion(-)

diff --git a/arrus/core/api/devices/us4r/Us4OEM.h b/arrus/core/api/devices/us4r/Us4OEM.h
index 66df7a319..2120a6498 100644
--- a/arrus/core/api/devices/us4r/Us4OEM.h
+++ b/arrus/core/api/devices/us4r/Us4OEM.h
@@ -17,6 +17,8 @@ class Us4OEM : public Device, public TriggerGenerator {
 
     virtual double getSamplingFrequency() = 0;
 
+    virtual float getFPGATemperature() = 0;
+
     Us4OEM(Us4OEM const&) = delete;
     Us4OEM(Us4OEM const&&) = delete;
     void operator=(Us4OEM const&) = delete;
diff --git a/arrus/core/api/devices/us4r/Us4R.h b/arrus/core/api/devices/us4r/Us4R.h
index 869133638..116f535c8 100644
--- a/arrus/core/api/devices/us4r/Us4R.h
+++ b/arrus/core/api/devices/us4r/Us4R.h
@@ -135,6 +135,16 @@ class Us4R : public DeviceWithComponents {
      */
     virtual void setRxSettings(const RxSettings &settings) = 0;
 
+    /**
+     * Returns the number of us4OEM modules that are used in this us4R system.
+     */
+    virtual uint8_t getNumberOfUs4OEMs() = 0;
+
+    /**
+     * Returns us4R device sampling frequency.
+     */
+    virtual float getSamplingFrequency() const = 0;
+
     virtual void start() = 0;
     virtual void stop() = 0;
 
diff --git a/arrus/core/devices/us4r/Us4RImpl.cpp b/arrus/core/devices/us4r/Us4RImpl.cpp
index a1ca547cb..6464a5757 100644
--- a/arrus/core/devices/us4r/Us4RImpl.cpp
+++ b/arrus/core/devices/us4r/Us4RImpl.cpp
@@ -276,4 +276,12 @@ void Us4RImpl::setActiveTermination(std::optional<uint16> value) {
     setRxSettings(newRxSettings);
 }
 
+uint8_t Us4RImpl::getNumberOfUs4OEMs() {
+    return us4oems.size();
+}
+
+float Us4RImpl::getSamplingFrequency() const {
+    return us4oems[0]->getSamplingFrequency();
+}
+
 }
\ No newline at end of file
diff --git a/arrus/core/devices/us4r/Us4RImpl.h b/arrus/core/devices/us4r/Us4RImpl.h
index abac6ad95..da1964b29 100644
--- a/arrus/core/devices/us4r/Us4RImpl.h
+++ b/arrus/core/devices/us4r/Us4RImpl.h
@@ -120,6 +120,8 @@ class Us4RImpl : public Us4R {
     void setLpfCutoff(uint32 value) override;
     void setDtgcAttenuation(std::optional<uint16> value) override;
     void setActiveTermination(std::optional<uint16> value) override;
+    uint8_t getNumberOfUs4OEMs() override;
+    float getSamplingFrequency() const override;
 
 private:
     UltrasoundDevice *getDefaultComponent();
diff --git a/arrus/core/devices/us4r/probeadapter/ProbeAdapterImplTest.cpp b/arrus/core/devices/us4r/probeadapter/ProbeAdapterImplTest.cpp
index e7f94d599..054aa4873 100644
--- a/arrus/core/devices/us4r/probeadapter/ProbeAdapterImplTest.cpp
+++ b/arrus/core/devices/us4r/probeadapter/ProbeAdapterImplTest.cpp
@@ -94,6 +94,7 @@ class MockUs4OEM : public Us4OEMImplBase {
         (override));
     MOCK_METHOD(Interval<Voltage>, getAcceptedVoltageRange, (), (override));
     MOCK_METHOD(double, getSamplingFrequency, (), (override));
+    MOCK_METHOD(float, getFPGATemperature, (), (override));
     MOCK_METHOD(void, startTrigger, (), (override));
     MOCK_METHOD(void, stopTrigger, (), (override));
     MOCK_METHOD(void, start, (), (override));
diff --git a/arrus/core/devices/us4r/us4oem/Us4OEMImpl.cpp b/arrus/core/devices/us4r/us4oem/Us4OEMImpl.cpp
index c8e364889..e12ecc4ad 100644
--- a/arrus/core/devices/us4r/us4oem/Us4OEMImpl.cpp
+++ b/arrus/core/devices/us4r/us4oem/Us4OEMImpl.cpp
@@ -639,6 +639,8 @@ inline void Us4OEMImpl::setActiveTerminationAfe(std::optional<uint16> param, boo
     }
 }
 
-
+float Us4OEMImpl::getFPGATemperature() {
+    return ius4oem->GetFPGATemp();
+}
 
 }
diff --git a/arrus/core/devices/us4r/us4oem/Us4OEMImpl.h b/arrus/core/devices/us4r/us4oem/Us4OEMImpl.h
index 1aed34bfa..39eadfd56 100644
--- a/arrus/core/devices/us4r/us4oem/Us4OEMImpl.h
+++ b/arrus/core/devices/us4r/us4oem/Us4OEMImpl.h
@@ -117,6 +117,7 @@ class Us4OEMImpl : public Us4OEMImplBase {
 
     std::vector<uint8_t> getChannelMapping() override;
     void setRxSettings(const RxSettings &newSettings) override;
+    float getFPGATemperature() override;
 
 private:
     using Us4OEMBitMask = std::bitset<Us4OEMImpl::N_ADDR_CHANNELS>;