Commit 4adcb11e authored by Tobias Winchen's avatar Tobias Winchen
Browse files

Loop over output blocks

parent 723a3828
......@@ -131,6 +131,7 @@ GatedSpectrometer<HandlerType, InputType, OutputType>::GatedSpectrometer(
BOOST_LOG_TRIVIAL(debug) << "Integrating " << nsamps_per_output_spectra <<
" samples from " << _nBlocks << "blocks into one output spectrum.";
// plan the FFT
size_t nsamps_per_buffer = _dadaBufferLayout.sizeOfData() * 8 / nbits;
int batch = nsamps_per_pol / _fft_length;
......@@ -286,8 +287,8 @@ void GatedSpectrometer<HandlerType, InputType, OutputType>::gated_fft(
CUFFT_ERROR_CHECK(cufftExecR2C(_fft_plan, (cufftReal *)_unpacked_voltage_ptr,
(cufftComplex *)_channelised_voltage_ptr));
CUDA_ERROR_CHECK(cudaStreamSynchronize(_proc_stream));
BOOST_LOG_TRIVIAL(debug) << "Exit processing";
// CUDA_ERROR_CHECK(cudaStreamSynchronize(_proc_stream));
// BOOST_LOG_TRIVIAL(debug) << "Exit processing";
} // process
......@@ -370,6 +371,8 @@ void GatedSpectrometer<HandlerType, InputType, OutputType>::process(SinglePolari
{
gated_fft(*inputDataStream, outputDataStream->G0._noOfBitSets.a(), outputDataStream->G1._noOfBitSets.a());
kernels::detect_and_accumulate<IntegratedPowerType> <<<1024, 1024, 0, _proc_stream>>>(
thrust::raw_pointer_cast(inputDataStream->_channelised_voltage_G0.data()),
thrust::raw_pointer_cast(outputDataStream->G0.data.a().data()),
......@@ -398,26 +401,31 @@ void GatedSpectrometer<HandlerType, InputType, OutputType>::process(DualPolariza
gated_fft(inputDataStream->polarization0, outputDataStream->G0._noOfBitSets.a(), outputDataStream->G1._noOfBitSets.a());
gated_fft(inputDataStream->polarization1, outputDataStream->G0._noOfBitSets.a(), outputDataStream->G1._noOfBitSets.a());
stokes_accumulate<<<1024, 1024, 0, _proc_stream>>>(
thrust::raw_pointer_cast(inputDataStream->polarization0._channelised_voltage_G0.data()),
thrust::raw_pointer_cast(inputDataStream->polarization1._channelised_voltage_G0.data()),
thrust::raw_pointer_cast(outputDataStream->G0.I.a().data()),
thrust::raw_pointer_cast(outputDataStream->G0.Q.a().data()),
thrust::raw_pointer_cast(outputDataStream->G0.U.a().data()),
thrust::raw_pointer_cast(outputDataStream->G0.V.a().data()),
_nchans, _naccumulate / _nBlocks
);
stokes_accumulate<<<1024, 1024, 0, _proc_stream>>>(
thrust::raw_pointer_cast(inputDataStream->polarization0._channelised_voltage_G1.data()),
thrust::raw_pointer_cast(inputDataStream->polarization1._channelised_voltage_G1.data()),
thrust::raw_pointer_cast(outputDataStream->G1.I.a().data()),
thrust::raw_pointer_cast(outputDataStream->G1.Q.a().data()),
thrust::raw_pointer_cast(outputDataStream->G1.U.a().data()),
thrust::raw_pointer_cast(outputDataStream->G1.V.a().data()),
_nchans, _naccumulate / _nBlocks
);
for(int output_block_number = 0; output_block_number < outputDataStream->G0._noOfBitSets.size(); output_block_number++)
{
size_t input_offset = output_block_number * inputDataStream->polarization0._channelised_voltage_G0.size() / outputDataStream->G0._noOfBitSets.size();
size_t output_offset = output_block_number * outputDataStream->G0.I.a().size() / outputDataStream->G0._noOfBitSets.size();
BOOST_LOG_TRIVIAL(debug) << "Accumulating data for output block " << output_block_number << " with input offset " << input_offset << " and output_offset " << output_offset;
stokes_accumulate<<<1024, 1024, 0, _proc_stream>>>(
thrust::raw_pointer_cast(inputDataStream->polarization0._channelised_voltage_G0.data() + input_offset),
thrust::raw_pointer_cast(inputDataStream->polarization1._channelised_voltage_G0.data() + input_offset),
thrust::raw_pointer_cast(outputDataStream->G0.I.a().data() + output_offset),
thrust::raw_pointer_cast(outputDataStream->G0.Q.a().data() + output_offset),
thrust::raw_pointer_cast(outputDataStream->G0.U.a().data() + output_offset),
thrust::raw_pointer_cast(outputDataStream->G0.V.a().data() + output_offset),
_nchans, _naccumulate / _nBlocks
);
stokes_accumulate<<<1024, 1024, 0, _proc_stream>>>(
thrust::raw_pointer_cast(inputDataStream->polarization0._channelised_voltage_G1.data() + input_offset),
thrust::raw_pointer_cast(inputDataStream->polarization1._channelised_voltage_G1.data() + input_offset),
thrust::raw_pointer_cast(outputDataStream->G1.I.a().data() + output_offset),
thrust::raw_pointer_cast(outputDataStream->G1.Q.a().data() + output_offset),
thrust::raw_pointer_cast(outputDataStream->G1.U.a().data() + output_offset),
thrust::raw_pointer_cast(outputDataStream->G1.V.a().data() + output_offset),
_nchans, _naccumulate / _nBlocks
);
}
//thrust::fill(thrust::cuda::par.on(_proc_stream),outputDataStream->G0.I.a().begin(), outputDataStream->G0.I.a().end(), _call_count);
//thrust::fill(thrust::cuda::par.on(_proc_stream),outputDataStream->G0.Q.a().begin(), outputDataStream->G0.Q.a().end(), _call_count);
//thrust::fill(thrust::cuda::par.on(_proc_stream),outputDataStream->G0.U.a().begin(), outputDataStream->G0.U.a().end(), _call_count);
......@@ -429,7 +437,7 @@ void GatedSpectrometer<HandlerType, InputType, OutputType>::process(DualPolariza
//thrust::fill(thrust::cuda::par.on(_proc_stream),outputDataStream->G1.U.a().begin(), outputDataStream->G1.U.a().end(), _call_count);
//thrust::fill(thrust::cuda::par.on(_proc_stream),outputDataStream->G1.V.a().begin(), outputDataStream->G1.V.a().end(), _call_count);
cudaDeviceSynchronize();
// cudaDeviceSynchronize();
}
} // edd
......
......@@ -80,12 +80,15 @@ with open(args.filename[0], 'rb') as inputfile:
for i in range(n_rows):
start = i * size_of_line
D0[i] = getSingleSpectrum(rawData[start: start + nChannels * bitDepth / 8], bitDepth)
D0N[i] = numpy.fromstring(rawData[start + nChannels * bitDepth / 8:][:8], dtype='uint64')
end = start + nChannels * bitDepth / 8
D0[i] = getSingleSpectrum(rawData[start:end], bitDepth)
start += size_of_line/2
D1[i] = getSingleSpectrum(rawData[start: start + nChannels * bitDepth / 8], bitDepth)
D1N[i] = numpy.fromstring(rawData[start + nChannels * bitDepth / 8:][:8], dtype='uint64')
start = end
end = start + nChannels * bitDepth / 8
D1[i] = getSingleSpectrum(rawData[start:end], bitDepth)
D0N[i] = numpy.fromstring(rawData[end:][:8], dtype='uint64')
D1N[i] = numpy.fromstring(rawData[end:][:8], dtype='uint64')
fig = plt.figure()
s1 = fig.add_subplot(121)
......@@ -119,28 +122,25 @@ with open(args.filename[0], 'rb') as inputfile:
N_off = numpy.zeros(n_rows)
size_of_spectrum = nChannels * bitDepth / 8
start_of_spectrum = 0
end_of_spectrum = start_of_spectrum + size_of_spectrum
for i in range(n_rows):
log.debug("Line {}:".format(i))
for j in range(4):
start_of_spectrum = i * (8 * (size_of_spectrum + 8)) + 2 * j * size_of_spectrum
end_of_spectrum = start_of_spectrum + size_of_spectrum
log.debug(" - j = {} off: [{}:{}]".format(j, start_of_spectrum, end_of_spectrum))
D_off[j,i] = getSingleSpectrum(rawData[start_of_spectrum:end_of_spectrum], bitDepth)
N_off[i] = numpy.fromstring(rawData[end_of_spectrum:end_of_spectrum+8], dtype='uint64')
log.debug(" P = {}, N = {}".format(sum(D_off[j,i]), N_off[i]))
start_of_spectrum = end_of_spectrum + 64/8
start_of_spectrum = end_of_spectrum
end_of_spectrum = start_of_spectrum + size_of_spectrum
log.debug(" - j = {}, on: [{}:{}]".format(j, start_of_spectrum, end_of_spectrum))
D_on[j,i] = getSingleSpectrum(rawData[start_of_spectrum:end_of_spectrum], bitDepth)
N_on[i] = numpy.fromstring(rawData[end_of_spectrum:end_of_spectrum+8], dtype='uint64')
log.debug(" P = {}, N = {}".format(sum(D_on[j,i]), N_on[i]))
start_of_spectrum = end_of_spectrum + 64/8
end_of_spectrum = start_of_spectrum + size_of_spectrum
start_of_data= i * (8 * (size_of_spectrum + 8)) + 8 * size_of_spectrum
end_of_data = start_of_data + 8*8
N = numpy.fromstring(rawData[start_of_data:end_of_data], dtype='uint64')
N_off[i] = N[0]
N_on[i] = N[1]
log.debug(" j = {} :: ON: P = {}, N = {} OFF: P = {}, N = {}".format(j, sum(D_on[j,i]), N_on[i], sum(D_off[j,i]), N_off[i]))
fig = plt.figure()
s1 = fig.add_subplot(241)
......
Markdown is supported
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment