From f14d193343c729703e2cd2e8fc9f83c90b8b0c7b Mon Sep 17 00:00:00 2001
From: Tobias Winchen <tobias.winchen@rwth-aachen.de>
Date: Tue, 12 Mar 2019 09:36:36 +0000
Subject: [PATCH] Switched to appropriate unsigned integer types

---
 psrdada_cpp/effelsberg/edd/GatedSpectrometer.cuh     |  8 ++++----
 .../effelsberg/edd/detail/GatedSpectrometer.cu       | 12 +++++++-----
 .../effelsberg/edd/test/src/GatedSpectrometerTest.cu |  2 +-
 3 files changed, 12 insertions(+), 10 deletions(-)

diff --git a/psrdada_cpp/effelsberg/edd/GatedSpectrometer.cuh b/psrdada_cpp/effelsberg/edd/GatedSpectrometer.cuh
index f460e826..3e585f0b 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 02918c81..59e52c00 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 bbdb08d3..15b6ccf8 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;
-- 
GitLab