Commit 3cce1e9a authored by Pierre NARVOR's avatar Pierre NARVOR
Browse files

[cuda] Added and tested a simple mapping type

parent 8d7bd981
...@@ -15,6 +15,7 @@ list(APPEND rtac_cuda_headers ...@@ -15,6 +15,7 @@ list(APPEND rtac_cuda_headers
include/rtac_base/cuda/functors.h include/rtac_base/cuda/functors.h
include/rtac_base/cuda/FunctorCompound.h include/rtac_base/cuda/FunctorCompound.h
include/rtac_base/cuda/Mapping.h include/rtac_base/cuda/Mapping.h
include/rtac_base/cuda/mappings.h
include/rtac_base/cuda/reductions.hcu include/rtac_base/cuda/reductions.hcu
include/rtac_base/cuda/operators.h include/rtac_base/cuda/operators.h
) )
......
...@@ -63,12 +63,12 @@ struct device_map_type { ...@@ -63,12 +63,12 @@ struct device_map_type {
// Deducing DeviceMap type to use (1D or 2D), depending on the // Deducing DeviceMap type to use (1D or 2D), depending on the
// FunctorT::OutputT type. // FunctorT::OutputT type.
using DeviceMapT = typename std::conditional<std::is_same<float, FoutT>::value, using BaseDeviceMapT = typename std::conditional<std::is_same<float, FoutT>::value,
DeviceMapping1D<T>, DeviceMapping1D<T>,
DeviceMapping2D<T>>::type; DeviceMapping2D<T>>::type;
// Building final type. // Building final type.
using type = functors::FunctorCompound<DeviceMapT, FunctorT>; using type = functors::FunctorCompound<BaseDeviceMapT, FunctorT>;
}; };
/** /**
...@@ -85,7 +85,8 @@ class Mapping ...@@ -85,7 +85,8 @@ class Mapping
using Ptr = std::shared_ptr<Mapping>; using Ptr = std::shared_ptr<Mapping>;
using ConstPtr = std::shared_ptr<const Mapping>; using ConstPtr = std::shared_ptr<const Mapping>;
using Texture = Texture2D<T>; using Texture = Texture2D<T>;
using DeviceMap = typename device_map_type<T, FunctorT>::type; using BaseDeviceMapT = typename device_map_type<T, FunctorT>::BaseDeviceMapT;
using DeviceMap = typename device_map_type<T, FunctorT>::type;
protected: protected:
...@@ -185,7 +186,7 @@ const FunctorT& Mapping<T,FunctorT>::functor() const ...@@ -185,7 +186,7 @@ const FunctorT& Mapping<T,FunctorT>::functor() const
template <typename T, class FunctorT> template <typename T, class FunctorT>
typename Mapping<T,FunctorT>::DeviceMap Mapping<T,FunctorT>::device_map() const typename Mapping<T,FunctorT>::DeviceMap Mapping<T,FunctorT>::device_map() const
{ {
return DeviceMap(DeviceMapping2D<T>({data_.texture()}), f_); return DeviceMap(BaseDeviceMapT({data_.texture()}), f_);
} }
}; //namespace cuda }; //namespace cuda
......
...@@ -51,6 +51,19 @@ struct Offset { ...@@ -51,6 +51,19 @@ struct Offset {
} }
}; };
template <typename Tout, typename Tin = Tout, typename Tscaling = Tout, typename Toff = Tout>
struct AffineTransform {
using InputT = Tin;
using OutputT = Tout;
Tscaling scaling;
Toff offset;
RTAC_HOSTDEVICE Tout operator()(const Tin& input) const {
return scaling*input + offset;
}
};
}; //namespace functors }; //namespace functors
}; //namespace cuda }; //namespace cuda
}; //namespace rtac }; //namespace rtac
......
#ifndef _DEF_RTAC_BASE_CUDA_MAPPINGS_H_
#define _DEF_RTAC_BASE_CUDA_MAPPINGS_H_
#include <rtac_base/cuda/functors.h>
#include <rtac_base/cuda/Mapping.h>
namespace rtac { namespace cuda {
template <typename T>
struct Mapping1D : public Mapping<T, functors::AffineTransform<float>>
{
using MappingT = Mapping<T, functors::AffineTransform<float>>;
using Ptr = typename MappingT::Ptr;
using ConstPtr = typename MappingT::ConstPtr;
using DeviceMap = typename MappingT::DeviceMap;
static Ptr Create(float domainMin, float domainMax,
const T* domainData, unsigned int dataSize)
{
Texture2D<T> tex;
tex.set_image(dataSize, 1, domainData);
auto t = functors::AffineTransform<float>(
{1.0f / (domainMax - domainMin),
-domainMin / (domainMax - domainMin)});
return MappingT::Create(std::move(tex), t);
}
};
}; //namespace cuda
}; //namespace rtac
#endif //_DEF_RTAC_BASE_CUDA_MAPPINGS_H_
#include <iostream> #include <iostream>
#include <vector>
using namespace std; using namespace std;
#include <rtac_base/files.h> #include <rtac_base/files.h>
...@@ -7,9 +8,35 @@ using namespace rtac::files; ...@@ -7,9 +8,35 @@ using namespace rtac::files;
#include "mapping_test.h" #include "mapping_test.h"
using namespace rtac::cuda; using namespace rtac::cuda;
std::vector<float> mapping_data(int N, float k = 2.f)
{
std::vector<float> output(N);
//for(int n = 0; n < N; n++) {
// float x = 0.9*M_PI*(n - 0.5f*(N-1)) / (N - 1);
// output[n] = tan(0.5*x);
//}
//float m = output[0], M = output[output.size() - 1];
//for(int n = 0; n < N; n++) {
// m = min(m, output[n]);
// M = max(M, output[n]);
//}
//for(int n = 0; n < N; n++) {
// output[n] = (output[n] - m) / (M - m);
//}
for(int n = 0; n < N; n++) {
float x = 4*(n - 0.5f*(N-1)) / (N - 1);
output[n] = 1.0 / (1 + exp(-k*x));
}
return output;
}
int main() int main()
{ {
Texture2D<float> texture(std::move(Texture2D<float>::checkerboard(4,4,0.0f,1.0f)));
Texture2D<float> texture(std::move(Texture2D<float>::checkerboard(64,64,0.0f,1.0f)));
int W = 1024, H = 720; int W = 1024, H = 720;
...@@ -24,5 +51,11 @@ int main() ...@@ -24,5 +51,11 @@ int main()
HostVector<float> rendered2 = render_mapping2(W,H,map2->device_map()); HostVector<float> rendered2 = render_mapping2(W,H,map2->device_map());
write_pgm("output2.pgm", W, H, rendered2.data()); write_pgm("output2.pgm", W, H, rendered2.data());
auto v = mapping_data(256);
auto map3 = Mapping3::Create(0,1,v.data(), v.size());
HostVector<float> rendered3 = render_mapping3(W,H,map1->device_map(), map3->device_map());
write_pgm("output3.pgm", W, H, rendered3.data());
return 0; return 0;
} }
...@@ -30,3 +30,16 @@ DeviceVector<float> render_mapping2(int W, int H, const Mapping2::DeviceMap& map ...@@ -30,3 +30,16 @@ DeviceVector<float> render_mapping2(int W, int H, const Mapping2::DeviceMap& map
return output; return output;
} }
DeviceVector<float> render_mapping3(int W, int H,
const Mapping1::DeviceMap& map1,
const Mapping3::DeviceMap& map3)
{
DeviceVector<float> output(W*H);
do_render_mapping3<<<1,1>>>(output.data(), W, H, map1, map3);
cudaDeviceSynchronize();
return output;
}
...@@ -5,6 +5,7 @@ ...@@ -5,6 +5,7 @@
#include <rtac_base/cuda/HostVector.h> #include <rtac_base/cuda/HostVector.h>
#include <rtac_base/cuda/Mapping.h> #include <rtac_base/cuda/Mapping.h>
#include <rtac_base/cuda/functors.h> #include <rtac_base/cuda/functors.h>
#include <rtac_base/cuda/mappings.h>
using namespace rtac::cuda; using namespace rtac::cuda;
// This should fail successfully. // This should fail successfully.
...@@ -31,10 +32,16 @@ struct NormalizerUV { ...@@ -31,10 +32,16 @@ struct NormalizerUV {
using Mapping1 = Mapping<float>; using Mapping1 = Mapping<float>;
using Mapping2 = Mapping<float, NormalizerUV>; using Mapping2 = Mapping<float, NormalizerUV>;
using Mapping3 = Mapping1D<float>;
DeviceVector<float> render_texture(int W, int H, const Texture2D<float>& texData); DeviceVector<float> render_texture(int W, int H, const Texture2D<float>& texData);
DeviceVector<float> render_mapping1(int W, int H, const Mapping1::DeviceMap& map); DeviceVector<float> render_mapping1(int W, int H, const Mapping1::DeviceMap& map);
DeviceVector<float> render_mapping2(int W, int H, const Mapping2::DeviceMap& map); DeviceVector<float> render_mapping2(int W, int H, const Mapping2::DeviceMap& map);
DeviceVector<float> render_mapping3(int W, int H,
const Mapping1::DeviceMap& map1,
const Mapping3::DeviceMap& map3);
#endif //_DEF_RTAC_BASE_CUDA_TESTS_MAPPING_TEST_H_ #endif //_DEF_RTAC_BASE_CUDA_TESTS_MAPPING_TEST_H_
...@@ -41,6 +41,19 @@ __global__ void do_render_mapping2(T* output, int W, int H, MappingT map) ...@@ -41,6 +41,19 @@ __global__ void do_render_mapping2(T* output, int W, int H, MappingT map)
} }
} }
template <typename T, class MappingT1, class MappingT2>
__global__ void do_render_mapping3(T* output, int W, int H,
MappingT1 map1, MappingT2 map2)
{
float2 p;
for(int h = 0; h < H; h++) {
p.y = map2(((float)h) / (H - 1));
for(int w = 0; w < W; w++) {
p.x = map2(1.0f - ((float)w) / (W - 1));
output[W*h + w] = map1(p);
}
}
}
#endif //_DEF_RTAC_BASE_CUDA_TESTS_MAPPING_TEST_HCU_ #endif //_DEF_RTAC_BASE_CUDA_TESTS_MAPPING_TEST_HCU_
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