Commit f14d1933 authored by Tobias Winchen's avatar Tobias Winchen
Browse files

Switched to appropriate unsigned integer types

parent 30450808
......@@ -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
......
......@@ -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;
......
......@@ -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;
......
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