diff --git a/psrdada_cpp/effelsberg/edd/GatedSpectrometer.cuh b/psrdada_cpp/effelsberg/edd/GatedSpectrometer.cuh index f460e8262e09bf58478ce9ff6cc36cf20b587269..3e585f0bf32f475925ef3d0e2d07dac2a3d4bd05 100644 --- a/psrdada_cpp/effelsberg/edd/GatedSpectrometer.cuh +++ b/psrdada_cpp/effelsberg/edd/GatedSpectrometer.cuh @@ -88,7 +88,7 @@ private: thrust::device_vector<RawVoltageType> const &sideChannelData, thrust::device_vector<IntegratedPowerType> &detected_G0, thrust::device_vector<IntegratedPowerType> &detected_G1, - thrust::device_vector<int> &noOfBitSet); + thrust::device_vector<unsigned int> &noOfBitSet); private: std::size_t _buffer_bytes; @@ -116,7 +116,7 @@ private: DoubleDeviceBuffer<IntegratedPowerType> _power_db_G0; DoubleDeviceBuffer<IntegratedPowerType> _power_db_G1; DoubleDeviceBuffer<RawVoltageType> _sideChannelData_db; - DoubleDeviceBuffer<int> _noOfBitSetsInSideChannel; + DoubleDeviceBuffer<unsigned int> _noOfBitSetsInSideChannel; thrust::device_vector<UnpackedVoltageType> _unpacked_voltage_G0; thrust::device_vector<UnpackedVoltageType> _unpacked_voltage_G1; @@ -154,8 +154,8 @@ private: * */ __global__ void gating(float *G0, float *G1, const int64_t *sideChannelData, - size_t N, size_t heapSize, int64_t bitpos, - int64_t noOfSideChannels, int64_t selectedSideChannel); + size_t N, size_t heapSize, size_t bitpos, + size_t noOfSideChannels, size_t selectedSideChannel); } // edd diff --git a/psrdada_cpp/effelsberg/edd/detail/GatedSpectrometer.cu b/psrdada_cpp/effelsberg/edd/detail/GatedSpectrometer.cu index 02918c81a031d5466825e508b927f7acd188ad05..59e52c00db987e930cdcfe301e244f7b3d95c603 100644 --- a/psrdada_cpp/effelsberg/edd/detail/GatedSpectrometer.cu +++ b/psrdada_cpp/effelsberg/edd/detail/GatedSpectrometer.cu @@ -13,8 +13,8 @@ namespace edd { __global__ void gating(float *G0, float *G1, const int64_t *sideChannelData, - size_t N, size_t heapSize, int64_t bitpos, - int64_t noOfSideChannels, int64_t selectedSideChannel) { + size_t N, size_t heapSize, size_t bitpos, + size_t noOfSideChannels, size_t selectedSideChannel) { for (int i = blockIdx.x * blockDim.x + threadIdx.x; (i < N); i += blockDim.x * gridDim.x) { const float w = G0[i]; @@ -32,7 +32,9 @@ __global__ void gating(float *G0, float *G1, const int64_t *sideChannelData, } -__global__ void countBitSet(const int64_t *sideChannelData, size_t N, int64_t bitpos, int64_t noOfSideChannels, int64_t selectedSideChannel, int *nBitsSet) +__global__ void countBitSet(const int64_t *sideChannelData, size_t N, size_t + bitpos, size_t noOfSideChannels, size_t selectedSideChannel, unsigned int + *nBitsSet) { // really not optimized reduction, but here only trivial array sizes. int i = blockIdx.x * blockDim.x + threadIdx.x; @@ -181,7 +183,7 @@ void GatedSpectrometer<HandlerType>::process( thrust::device_vector<RawVoltageType> const &digitiser_raw, thrust::device_vector<RawVoltageType> const &sideChannelData, thrust::device_vector<IntegratedPowerType> &detected_G0, - thrust::device_vector<IntegratedPowerType> &detected_G1, thrust::device_vector<int> &noOfBitSet) { + thrust::device_vector<IntegratedPowerType> &detected_G1, thrust::device_vector<unsigned int> &noOfBitSet) { BOOST_LOG_TRIVIAL(debug) << "Unpacking raw voltages"; switch (_nbits) { case 8: @@ -300,7 +302,7 @@ bool GatedSpectrometer<HandlerType>::operator()(RawBytes &block) { int R[1]; CUDA_ERROR_CHECK(cudaMemcpyAsync(static_cast<void *>(R), static_cast<void *>(_noOfBitSetsInSideChannel.b_ptr()), - 1 * sizeof(int),cudaMemcpyDeviceToHost, _d2h_stream)); + 1 * sizeof(unsigned int),cudaMemcpyDeviceToHost, _d2h_stream)); BOOST_LOG_TRIVIAL(info) << "NOOF BIT SET IN SIDE CHANNEL: " << R[0] << std::endl; diff --git a/psrdada_cpp/effelsberg/edd/test/src/GatedSpectrometerTest.cu b/psrdada_cpp/effelsberg/edd/test/src/GatedSpectrometerTest.cu index bbdb08d3415dce6b73d3b8db004c7f4516dff287..15b6ccf8b08f4f8bd6a89017d4cfbda4df19e772 100644 --- a/psrdada_cpp/effelsberg/edd/test/src/GatedSpectrometerTest.cu +++ b/psrdada_cpp/effelsberg/edd/test/src/GatedSpectrometerTest.cu @@ -115,7 +115,7 @@ TEST(GatedSpectrometer, countBitSet) { const int64_t *sideCD = (int64_t *)(thrust::raw_pointer_cast(_sideChannelData.data())); - thrust::device_vector<int> res(1); + thrust::device_vector<unsigned int> res(1); // test 1 side channel res[0] = 0;