Spaces:
Runtime error
Runtime error
File size: 3,030 Bytes
be11144 |
1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 28 29 30 31 32 33 34 35 36 37 38 39 40 41 42 43 44 45 46 47 48 49 50 51 52 53 54 55 56 57 58 59 60 61 62 63 64 65 66 67 68 69 70 71 72 73 74 75 76 77 78 79 80 81 82 83 84 85 86 87 88 89 90 91 92 93 94 95 96 97 98 99 100 101 102 103 104 105 106 107 108 109 110 111 |
#include <thrust/device_vector.h>
#include <thrust/functional.h>
#include <thrust/gather.h>
#include <thrust/iterator/transform_input_output_iterator.h>
#include <thrust/sequence.h>
#include <iostream>
// Base 2 fixed point
class ScaledInteger
{
int value_;
int scale_;
public:
__host__ __device__
ScaledInteger(int value, int scale): value_{value}, scale_{scale} {}
__host__ __device__
int value() const { return value_; }
__host__ __device__
ScaledInteger rescale(int scale) const
{
int shift = scale - scale_;
int result = shift < 0 ? value_ << (-shift) : value_ >> shift;
return ScaledInteger{result, scale};
}
__host__ __device__
friend ScaledInteger operator+(ScaledInteger a, ScaledInteger b)
{
// Rescale inputs to the lesser of the two scales
if (b.scale_ < a.scale_)
a = a.rescale(b.scale_);
else if (a.scale_ < b.scale_)
b = b.rescale(a.scale_);
return ScaledInteger{a.value_ + b.value_, a.scale_};
}
};
struct ValueToScaledInteger
{
int scale;
__host__ __device__
ScaledInteger operator()(const int& value) const
{
return ScaledInteger{value, scale};
}
};
struct ScaledIntegerToValue
{
int scale;
__host__ __device__
int operator()(const ScaledInteger& scaled) const
{
return scaled.rescale(scale).value();
}
};
int main(void)
{
const size_t size = 4;
thrust::device_vector<int> A(size);
thrust::device_vector<int> B(size);
thrust::device_vector<int> C(size);
thrust::sequence(A.begin(), A.end(), 1);
thrust::sequence(B.begin(), B.end(), 5);
const int A_scale = 16; // Values in A are left shifted by 16
const int B_scale = 8; // Values in B are left shifted by 8
const int C_scale = 4; // Values in C are left shifted by 4
auto A_begin = thrust::make_transform_input_output_iterator(A.begin(),
ValueToScaledInteger{A_scale}, ScaledIntegerToValue{A_scale});
auto A_end = thrust::make_transform_input_output_iterator(A.end(),
ValueToScaledInteger{A_scale}, ScaledIntegerToValue{A_scale});
auto B_begin = thrust::make_transform_input_output_iterator(B.begin(),
ValueToScaledInteger{B_scale}, ScaledIntegerToValue{B_scale});
auto C_begin = thrust::make_transform_input_output_iterator(C.begin(),
ValueToScaledInteger{C_scale}, ScaledIntegerToValue{C_scale});
// Sum A and B as ScaledIntegers, storing the scaled result in C
thrust::transform(A_begin, A_end, B_begin, C_begin, thrust::plus<ScaledInteger>{});
thrust::host_vector<int> A_h(A);
thrust::host_vector<int> B_h(B);
thrust::host_vector<int> C_h(C);
std::cout << std::hex;
std::cout << "Expected [ ";
for (size_t i = 0; i < size; i++) {
const int expected = ((A_h[i] << A_scale) + (B_h[i] << B_scale)) >> C_scale;
std::cout << expected << " ";
}
std::cout << "] \n";
std::cout << "Result [ ";
for (size_t i = 0; i < size; i++) {
std::cout << C_h[i] << " ";
}
std::cout << "] \n";
return 0;
}
|