diff --git a/RTCP/Cobalt/GPUProc/share/gpu/kernels/IntToFloat.cuh b/RTCP/Cobalt/GPUProc/share/gpu/kernels/IntToFloat.cuh index fa0485b7a5e572241f3d34f5f0e3a0b31ca17255..45311aaa73fc1cc6c2bc0e00d884639602a85da3 100644 --- a/RTCP/Cobalt/GPUProc/share/gpu/kernels/IntToFloat.cuh +++ b/RTCP/Cobalt/GPUProc/share/gpu/kernels/IntToFloat.cuh @@ -52,7 +52,9 @@ inline __device__ float convertIntToFloat(signed char x) // Extract the 4-bit real or imaginary part of an 8-bit input sample inline __device__ signed char extractRI(signed char x, bool imag) { - return imag ? x >> 4 : (x << 4) >> 4; // preserve sign + // Note 1: Imaginary part is in the top 4 bits. See also RSP::decode4bit() in InputProc/Station/RSP.h. + // Note 2: Preserve the sign, so use sign-extending shifts to extract the right bits. + return imag ? x >> 4 : (x << 4) >> 4; } // WARNING: Caller is responsible for extracting 4-bit real or imaginary part from sample byte @@ -63,9 +65,7 @@ inline __device__ float convertIntToFloat(signed char x) // TODO: Is this the right scaling for 4-bit mode? // Keep output scale the same as 16 bit mode. - // Gains (input and complex voltages) end up x16, - // power (visibilities and Stokes) end up x16^2. - return 16 * i; + return 64 * i; } #else #error unsupported NR_BITS_PER_SAMPLE