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

Reduce floating point precision error

Using floats here instead of doubles causes a change by 1 in the 8-bit modem in
the test.
parent 90bda5db
......@@ -16,18 +16,18 @@ __global__
void detect_and_accumulate(float2 const* __restrict__ in, int8_t* __restrict__ out,
int nchans, int nsamps, int naccumulate, float scale, float offset)
{
// grid stride loop over output array, if input,output, nchans and naccumulate are all nice powers of 2
// grid stride loop over output array to keep
for (size_t i = blockIdx.x * blockDim.x + threadIdx.x; (i < nsamps * nchans / naccumulate); i += blockDim.x * gridDim.x)
{
float sum = 0.0f;
double sum = 0.0f;
size_t currentOutputSpectra = i / nchans;
size_t currentChannel = i % nchans;
for (size_t j = 0; j < naccumulate; j++)
{
float2 tmp = in[ j * nchans + currentOutputSpectra * nchans * naccumulate + currentChannel];
float x = tmp.x * tmp.x;
float y = tmp.y * tmp.y;
double x = tmp.x * tmp.x;
double y = tmp.y * tmp.y;
sum += x + y;
}
out[i] = (int8_t) ((sum - offset)/scale);
......@@ -50,8 +50,8 @@ void detect_and_accumulate(float2 const* __restrict__ in, float* __restrict__ ou
for (size_t j = 0; j < naccumulate; j++)
{
float2 tmp = in[ j * nchans + currentOutputSpectra * nchans * naccumulate + currentChannel];
float x = tmp.x * tmp.x;
float y = tmp.y * tmp.y;
double x = tmp.x * tmp.x;
double y = tmp.y * tmp.y;
sum += x + y;
}
out[i] = sum;
......
......@@ -47,7 +47,7 @@ void DetectorAccumulatorTester::detect_c_reference(
output_sample_idx < nsamples_out;
++output_sample_idx)
{
float value = 0.0f;
double value = 0.0f;
for (int input_sample_offset=0;
input_sample_offset < tscrunch;
++input_sample_offset)
......
Supports Markdown
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