...
 
Commits (197)
......@@ -129,6 +129,7 @@ unit_test_TMSS:
unit_test_RAServices:
stage: unit_test
image: ci_sas:latest
allow_failure: true
script:
- PACKAGE=RAServices
- echo "Testing $PACKAGE..."
......@@ -151,6 +152,7 @@ unit_test_RAServices:
unit_test_LTAIngest:
stage: unit_test
image: ci_lta:latest
allow_failure: true
script:
- PACKAGE=LTAIngest
- echo "Testing $PACKAGE..."
......@@ -243,6 +245,7 @@ integration_test_TMSS:
integration_test_RAServices:
stage: integration_test
image: ci_sas:latest
allow_failure: true
services:
- rabbitmq:latest
variables:
......@@ -265,6 +268,7 @@ integration_test_RAServices:
integration_test_LTAIngest:
stage: integration_test
image: ci_lta:latest
allow_failure: true
script:
- PACKAGE=LTAIngest
- echo "Integration Testing $PACKAGE..."
......
......@@ -647,6 +647,7 @@ namespace LOFAR
8, // nrBits
5, // scaleMax
-5, // scaleMin
false // sIpositive
};
ObservationSettings::BeamFormer::StokesSettings
......@@ -720,12 +721,14 @@ namespace LOFAR
struct ObservationSettings::BeamFormer::StokesSettings::QuantizerSettings *qtSettings
= &stSettings->quantizerSettings;
qtSettings->enabled=doQuantize;
qtSettings->nrBits=getUint32(prefix + ".quantizeBits");
qtSettings->nrBits=getUint32(prefix + ".quantizeBits",stSettings->quantizerSettings.nrBits);
ASSERT(qtSettings->nrBits == 8 || qtSettings->nrBits == 16);
qtSettings->scaleMax=getDouble(prefix + ".quantizeScaleMax");
qtSettings->scaleMin=getDouble(prefix + ".quantizeScaleMin");
qtSettings->scaleMax=getDouble(prefix + ".quantizeScaleMax",stSettings->quantizerSettings.scaleMax);
qtSettings->scaleMin=getDouble(prefix + ".quantizeScaleMin",stSettings->quantizerSettings.scaleMin);
ASSERT(qtSettings->scaleMax>qtSettings->scaleMin);
qtSettings->sIpositive= getBool(prefix+".quantizeIpositive",stSettings->quantizerSettings.sIpositive);
}
}
......
......@@ -708,6 +708,9 @@ namespace LOFAR
// key: Cobalt.BeamFormer.{Coherent|Incoherent}Stokes.quantizeScaleMin
double scaleMax;
double scaleMin;
// Keeping range of Stokes I positive if true
// key: Cobalt.BeamFormer.{Coherent|Incoherent}Stokes.quantizeIpositive
bool sIpositive;
};
struct QuantizerSettings quantizerSettings;
......
......@@ -2,11 +2,15 @@
Cobalt in DFDs
-----------------------------------
The data flow through Cobalt is described in DFDs in
The data flows through Cobalt are summarized in the following PDF files
cobalt-data-flow.eap (Enterprise Architect Project file)
Component Model.pdf
OutputProc.pdf
Station Input Pipeline.pdf
CPU Processing Pipeline.pdf
Signal Processing Pipeline.pdf
Please update the above diagram along with this document. The data through Cobalt flows though Queues between processes. Both are briefly described before the full pipeline is described in more detail.
The data through Cobalt flows though Queues between processes. Both are briefly described before the full pipeline is described in more detail.
---------
Processes
......
This document describes the inner workings of the CUDA kernel used in quantizing the 32 bit input data into 8 bit data.
The Parset keys used in quantization of beamformed data are
For coherent Stokes:
* `Cobalt.BeamFormer.CoherentStokes.quantize=false`
* `Cobalt.BeamFormer.CoherentStokes.quantizeBits=8`
* `Cobalt.BeamFormer.CoherentStokes.quantizeScaleMax=5`
* `Cobalt.BeamFormer.CoherentStokes.quantizeScaleMin=-5`
* `Cobalt.BeamFormer.CoherentStokes.quantizeIpositive=false`
For incoherent Stokes:
* `Cobalt.BeamFormer.IncoherentStokes.quantize=false`
* `Cobalt.BeamFormer.IncoherentStokes.quantizeBits=8`
* `Cobalt.BeamFormer.IncoherentStokes.quantizeScaleMax=5`
* `Cobalt.BeamFormer.IncoherentStokes.quantizeScaleMin=-5`
* `Cobalt.BeamFormer.IncoherentStokes.quantizeIpositive=false`
The values for each key shown above are the default values (if that particular key is not defined in the parset). The description of the keys:
* `.quantize=true|false`: If true, the output will be quantized (instead of using 32 bit float as output datatype, a reduced number of bits will be used).
* `.quantizeBits=8`: Currently, 8 bits will be used to store each quantized data point. This implies the output data type will be signed char (int8_t) or unsigned char (uint8_t). In addtion, scale and offset of each data block will also be produced as 32 bit float values.
* `.quantizeScaleMax` and .`quantizeScaleMin`: These two keys will be used to cut off extreme data points above or below a threshold, prior to quantization.
* `.quantizeIpositive=true|false`: If `quantizeScaleMin` is negative, the usable range for Stokes I will also include some values below zero. However, by definition, Stokes I is always positive. By setting this key to true, we can override this to only consider the positive range of Stokes I for quantization. In this way, the available number of bits in the quantizer are not wasted by representing values that do not exist in the orignal data.
Let us call the values defined by `quantizeScaleMax` and `quantizeScaleMin` as `Smax` and `Smin`, respectively.
<img src="quantization.png" alt="How quantization works" width="700"/>
We can explain the workings of the quantization kernel by looking at the probability density function (PDF) from the input 32 bit data to output 8 bit data as shown in the above figure.
* The input PDF will have data in the range -inf...inf with one exception, for Stokes I, it will be 0...inf. (The notation inf represents infinity). This is shown on the top plot.
* Using the user defined *Smax* and *Smin*. the range of the data to be quantized is selected. In most cases it is the range between *Smin*&times;&sigma; and *Smax*&times;&sigma;. (&sigma; is the standard deviation). The user can specify the best values for *Smax* and *Smin* depending on how the input data are distributed (i.e., by looking at the PDF). The input data that fall outside this range are cut-off as show in in the middle plot. In addition the mean (&mu;) of the original data will be subtracted in the cases of Stokes Q,U,V.
* Finally, the original range is mapped to values in the range 0 to 255 (for unsigned data, i.e., Stokes I) or in the range -128 to 127 (for signed data) as shown in the bottom plot. The output PDF will be the area-sampled version of the input PDF (after cutoff) as shown in this plot.
* A special case is quantization of Stokes I, because the original data have values in 0...inf. If the user-specified *Smin* value is negative, some useful number of quantized levels will be used to represent negative data that does not exist. Therefore, in this case, the user supplied *Smin* value will be over-ridden to use 0 instead. This behaviour is enabled by setting `.quantizeIpositive=true`.
In order to (approximately) recover the original data from the quantized data, the scale and offset that are used to transform the data to the ranges 0 to 255 or -128 to 127 are also produced as output. The scale and offset are determined per each block of data (duration approximately 1 sec) and per each channel and polarization. The number of data samples per each block is dependent on the integration factor and the sampling clock frequency.
......@@ -245,6 +245,7 @@ __device__ int clamp(int x, int a, int b)
* NR_SAMPLES_PER_CHANNEL | >= 4 | number of input samples per channel
* NR_TABS | >= 1 | number of tabs to create
* NR_QUANTIZE_BITS | 8, 16 | number of bits used for the quantization
* STOKES_I_POSITIVE | 1 or 0 | if 1, and if smin<0, override and set smin=0 for stokes I
*
* Every block of data will require sizeof(INT) + 2*sizeof(float),
* therefore NR_SAMPLES_PER_CHANNEL should be at least 5 (in case of 16-bit ouput) to
......@@ -306,17 +307,24 @@ __global__ void quantizeOutput(
// Make sure that the full range is used, if smin<0 therefore:
// - for Stokes I, map the values to (0:std*smax)
// - for Stokes QUV use the range (std*smin:std:smax)
if (!COMPLEX_VOLTAGES && !stokes_quv && smin<0.0f)
#if (STOKES_I_POSITIVE == 1)
bool unsigned_stokes_i=(!COMPLEX_VOLTAGES && !stokes_quv && smin<0.0f);
if (unsigned_stokes_i)
{
smin=0.0f;
}
#endif
for (unsigned sample_idx = lane; sample_idx < NR_SAMPLES_PER_CHANNEL; sample_idx += block.size())
{
// Compute scale and offset
int xmin = 0;
int xmax = (1 << NR_QUANTIZE_BITS) - 1;
#if (STOKES_I_POSITIVE == 1)
float xoffset = (unsigned_stokes_i ? 0.0f: mean + smin * std);
#else
float xoffset = mean + smin * std;
#endif
float xscale = ((smax - smin) * std) / (xmax + 1);
if (sample_idx < NR_SAMPLES_PER_CHANNEL)
......
......@@ -63,6 +63,7 @@ namespace LOFAR
nrQuantizeBits = stokes.quantizerSettings.nrBits;
quantizeScaleMax = stokes.quantizerSettings.scaleMax;
quantizeScaleMin = stokes.quantizerSettings.scaleMin;
sIpositive = stokes.quantizerSettings.sIpositive;
nrThreadsPerBlock = 512;
// The QuantizeOutputKernel reuses the buffers of CoherentStokesKernel in reverse order:
......@@ -229,6 +230,8 @@ namespace LOFAR
lexical_cast<string>(itsParameters.nrTABs);
defs["COMPLEX_VOLTAGES"] =
itsParameters.outputComplexVoltages ? "1" : "0";
defs["STOKES_I_POSITIVE"] =
itsParameters.sIpositive ? "1" : "0";
defs["NR_COHERENT_STOKES"] =
lexical_cast<string>(itsParameters.nrStokes);
......
......@@ -58,6 +58,7 @@ namespace LOFAR
unsigned nrQuantizeBits;
float quantizeScaleMax;
float quantizeScaleMin;
bool sIpositive;
unsigned int nrThreadsPerBlock;
......