[cuda/mappings] Completed mapping types, using functor types

cuda/include/rtac_base/cuda/Mapping.h

@@ -2,33 +2,34 @@
 #define _DEF_RTAC_BASE_CUDA_MAPPING_H_
 
 #include <rtac_base/cuda/Texture2D.h>
+#include <rtac_base/cuda/FunctorCompound.h>
+#include <rtac_base/cuda/functors.h>
 
 namespace rtac { namespace cuda {
 
 /**
- * This struct encodes a generic mapping from an arbitrary input space.  The
- * mapping is encoded into the texture.
+ * Mappings are special classes of functors which output value is fetched from
+ * a texture. The main purpose of this type is to leverage the fast local
+ * caching and interpolation of GPU texture units. This is especially relevant
+ * for operations such as correlations or convolutions.
  *
- * A Functor type is used to transform the input space input local texture
- * coordinates float2.
- *
- * This can be used to implement fast linear or nearest interpolation on a GPU,
- * colormapping a gray scale image, or simply reading in an image using
- * coordinates in an arbitrary unit.
- *
- * Caution : the result of the operation ax + b must be defined and be
- * implicitly castable to a float2.
+ * Mapping are handled differently than regular functors because the texture
+ * memory must be managed from host side (creation, memory allocation and
+ * destruction...). On the device side however, a mapping behaves like a
+ * regular functor.
  */
-template <class FunctorT, typename T>
-struct DeviceMapping
+template <typename T>
+struct DeviceMapping2D
 {
+    // Declaring input and output types to be compatible with other functors.
+    using InputT  = float2;
+    using OutputT = T;
+
     cudaTextureObject_t data;
-    FunctorT f;
 
-    #ifdef RTAC_CUDACC
-    __device__ T operator()(const typename FunctorT::Input& x) const {
-        float2 p = f(x);
-        return tex2D<T>(data, p.x, p.y);
+    #ifdef RTAC_CUDACC // this cannot be used from host side because of the texture fetch
+    __device__ T operator()(const float2& uv) const {
+        return tex2D<T>(data, uv.x, uv.y);
     }
     #endif //RTAC_CUDACC
 };
@@ -36,14 +37,15 @@ struct DeviceMapping
 }; //namespace cuda
 }; //namespace rtac
 
-#ifndef RTAC_CUDACC
-
 namespace rtac { namespace cuda {
 
 /**
- * Host side class with manage device side Device mapping data.
+ * Host side class with manage device side Device mapping data. This also
+ * handle an additional functor which allows to perform an operation on the
+ * input coordinates (namely a nomalization between 0 and 1) before effectively
+ * performing a texture fetch.
  */
-template <class FunctorT, typename T>
+template <typename T, class FunctorT = functors::IdentityFunctor<float2>>
 class Mapping
 {
     public:
@@ -51,23 +53,27 @@ class Mapping
     using Ptr       = std::shared_ptr<Mapping>;
     using ConstPtr  = std::shared_ptr<const Mapping>;
     using Texture   = Texture2D<T>;
-    using DeviceMap = DeviceMapping<FunctorT,T>;
+    using DeviceMap = functors::FunctorCompound<DeviceMapping2D<T>, FunctorT>;
 
     protected:
     
     Texture  data_;
     FunctorT f_;
     
-    Mapping(const FunctorT& f, Texture&& data);
+    Mapping(Texture&& data);
+    Mapping(Texture&& data, const FunctorT& f);
 
     public:
 
-    static Ptr Create(const FunctorT& f, Texture&& data);
-    static Ptr Create(const FunctorT& f, const Texture& data);
+    static Ptr Create(Texture&& data);
+    static Ptr Create(const Texture& data);
+
+    static Ptr Create(Texture&& data, const FunctorT& f);
+    static Ptr Create(const Texture& data, const FunctorT& f);
 
     void set_texture(const Texture& texture);
     void set_texture(Texture&& texture);
-    void set_FunctorT(const FunctorT& a);
+    void set_FunctorT(const FunctorT& f);
 
     const Texture& texture() const;
     const FunctorT& functor() const;
@@ -75,65 +81,82 @@ class Mapping
     DeviceMap device_map() const;
 };
 
-template <class FunctorT, typename T>
-Mapping<FunctorT,T>::Mapping(const FunctorT& f, Texture&& data) :
+template <typename T, class FunctorT>
+Mapping<T,FunctorT>::Mapping(Texture&& data) :
+    data_(data)
+{}
+
+template <typename T, class FunctorT>
+Mapping<T,FunctorT>::Mapping(Texture&& data, const FunctorT& f) :
     data_(data),
     f_(f)
 {}
 
-template <class FunctorT, typename T>
-typename Mapping<FunctorT,T>::Ptr 
-Mapping<FunctorT,T>::Create(const FunctorT& f, const Texture& data)
+template <typename T, class FunctorT>
+typename Mapping<T,FunctorT>::Ptr 
+Mapping<T,FunctorT>::Create(Texture&& data)
 {
-    return Ptr(new Mapping<FunctorT,T>(f, data));
+    return Ptr(new Mapping<T,FunctorT>(std::move(Texture(data))));
 }
 
-template <class FunctorT, typename T>
-typename Mapping<FunctorT,T>::Ptr 
-Mapping<FunctorT,T>::Create(const FunctorT& f, Texture&& data)
+template <typename T, class FunctorT>
+typename Mapping<T,FunctorT>::Ptr 
+Mapping<T,FunctorT>::Create(const Texture& data)
 {
-    return Ptr(new Mapping<FunctorT,T>(f, std::move(Texture(data))));
+    return Ptr(new Mapping<T,FunctorT>(std::move(Texture(data))));
 }
 
-template <class FunctorT, typename T>
-void Mapping<FunctorT,T>::set_texture(const Texture& texture)
+template <typename T, class FunctorT>
+typename Mapping<T,FunctorT>::Ptr 
+Mapping<T,FunctorT>::Create(Texture&& data, const FunctorT& f)
+{
+    return Ptr(new Mapping<T,FunctorT>(std::move(Texture(data)), f));
+}
+
+template <typename T, class FunctorT>
+typename Mapping<T,FunctorT>::Ptr 
+Mapping<T,FunctorT>::Create(const Texture& data, const FunctorT& f)
+{
+    return Ptr(new Mapping<T,FunctorT>(std::move(Texture(data)), f));
+}
+
+template <typename T, class FunctorT>
+void Mapping<T,FunctorT>::set_texture(const Texture& texture)
 {
     data_ = texture;
 }
 
-template <class FunctorT, typename T>
-void Mapping<FunctorT,T>::set_texture(Texture&& texture)
+template <typename T, class FunctorT>
+void Mapping<T,FunctorT>::set_texture(Texture&& texture)
 {
     data_ = texture;
 }
 
-template <class FunctorT, typename T>
-void Mapping<FunctorT,T>::set_FunctorT(const FunctorT& f)
+template <typename T, class FunctorT>
+void Mapping<T,FunctorT>::set_FunctorT(const FunctorT& f)
 {
     f_ = f;
 }
 
-template <class FunctorT, typename T>
-const Texture2D<T>& Mapping<FunctorT,T>::texture() const
+template <typename T, class FunctorT>
+const Texture2D<T>& Mapping<T,FunctorT>::texture() const
 {
     return data_;
 }
 
-template <class FunctorT, typename T>
-const FunctorT& Mapping<FunctorT,T>::functor() const
+template <typename T, class FunctorT>
+const FunctorT& Mapping<T,FunctorT>::functor() const
 {
     return f_;
 }
 
-template <class FunctorT, typename T>
-typename Mapping<FunctorT,T>::DeviceMap Mapping<FunctorT,T>::device_map() const
+template <typename T, class FunctorT>
+typename Mapping<T,FunctorT>::DeviceMap Mapping<T,FunctorT>::device_map() const
 {
-    return DeviceMapping<FunctorT,T>({data_.texture(), f_});
+    return DeviceMap(DeviceMapping2D<T>({data_.texture()}), f_);
 }
 
 }; //namespace cuda
 }; //namespace rtac
 
-#endif //RTAC_CUDACC
-
 #endif //_DEF_RTAC_BASE_CUDA_MAPPING_H_

cuda/include/rtac_base/cuda/functors.h

@@ -10,6 +10,22 @@
 
 namespace rtac { namespace cuda { namespace functors {
 
+/**
+ * This functor is usefull to define default template argument which have no
+ * effects. This shouldn't have any impact on performance after an optimized
+ * compilation.
+ */
+template <typename T>
+struct IdentityFunctor
+{
+    using InputT  = T;
+    using OutputT = T;
+
+    RTAC_HOSTDEVICE const T& operator()(const T& input) const {
+        return input;
+    }
+};
+
 template <typename Tout, typename Tin = Tout, typename Tscale = Tin>
 struct Scaling {
     using InputT  = Tin;

cuda/tests/CMakeLists.txt

@@ -32,8 +32,8 @@ foreach(name ${cuda_test_names})
 endforeach(name)
 
 add_subdirectory(texture)
-add_subdirectory(mapping)
 add_subdirectory(functors)
+add_subdirectory(mapping)
 
 
 add_executable(reductions_test

cuda/tests/mapping/src/mapping_test.cpp

 #include <iostream>
 using namespace std;
 
+#include <rtac_base/files.h>
+using namespace rtac::files;
+
 #include "mapping_test.h"
+using namespace rtac::cuda;
 
 int main()
 {
-    std::vector<float> data(16);
-    for(int i = 0; i < data.size(); i++) {
-        float x = (2.0*i) / (data.size() - 1) - 1.0f;
-        data[i] = 1.0 - x*x;
-        cout << " " << data[i];
-    }
-    cout << endl;
+    Texture2D<float> texture(std::move(Texture2D<float>::checkerboard(4,4,0.0f,1.0f)));
+
+    int W = 1024, H = 720;
     
-    Texture2D<float> texture;
-    texture.set_wrap_mode(Texture2D<float>::WrapClamp);
-    texture.set_image(data.size(), 1, data.data());
+    HostVector<float> rendered0 = render_texture(W,H,texture);
+    write_pgm("output0.pgm", W, H, rendered0.data());
+
+    auto map1 = Mapping1::Create(texture);
+    HostVector<float> rendered1 = render_mapping1(W,H,map1->device_map());
+    write_pgm("output1.pgm", W, H, rendered1.data());
 
-    auto mapping = Mapping<Affine1D,float>::Create(Affine1D({1,0}),
-                                                   std::move(texture));
-    HostVector<float> x(8);
-    for(int i = 0; i < x.size(); i++) {
-        x[i] = ((float)i) / (x.size() - 1);
-    }
-    HostVector<float> out = map(mapping->device_map(), x);
-    for(auto v : out) {
-        cout << " " << v;
-    }
-    cout << endl;
+    auto map2 = Mapping2::Create(texture, NormalizerUV({uint2({W,H})}));
+    HostVector<float> rendered2 = render_mapping2(W,H,map2->device_map());
+    write_pgm("output2.pgm", W, H, rendered2.data());
 
     return 0;
 }

cuda/tests/mapping/src/mapping_test.cu

 #include "mapping_test.h"
+#include "mapping_test.hcu"
 
+DeviceVector<float> render_texture(int W, int H, const Texture2D<float>& texData)
+{
+    DeviceVector<float> output(W*H);
+    
+    do_render_texture<<<1,1>>>(output.data(), W, H, texData.texture());
+    cudaDeviceSynchronize();
+
+    return output;
+}
 
-__global__ void do_map(DeviceMapping<Affine1D,float> mapping, 
-                       const float* x, float* out, unsigned int size)
+DeviceVector<float> render_mapping1(int W, int H, const Mapping1::DeviceMap& map)
 {
-    for(int i = 0; i < size; i++) {
-        out[i] = mapping(x[i]);
-    }
+    DeviceVector<float> output(W*H);
+    
+    do_render_mapping1<<<1,1>>>(output.data(), W, H, map);
+    cudaDeviceSynchronize();
+
+    return output;
 }
 
-DeviceVector<float> map(const DeviceMapping<Affine1D,float>& mapping,
-                        const DeviceVector<float>& x)
+DeviceVector<float> render_mapping2(int W, int H, const Mapping2::DeviceMap& map)
 {
-    DeviceVector<float> out(x.size());
+    DeviceVector<float> output(W*H);
     
-    do_map<<<1,1>>>(mapping, x.data(), out.data(), out.size());
+    do_render_mapping2<<<1,1>>>(output.data(), W, H, map);
     cudaDeviceSynchronize();
 
-    return out;
+    return output;
 }

cuda/tests/mapping/src/mapping_test.h

 #ifndef _DEF_RTAC_BASE_CUDA_TESTS_MAPPING_TEST_H_
 #define _DEF_RTAC_BASE_CUDA_TESTS_MAPPING_TEST_H_
 
-#include <rtac_base/cuda/Mapping.h>
 #include <rtac_base/cuda/DeviceVector.h>
 #include <rtac_base/cuda/HostVector.h>
+#include <rtac_base/cuda/Mapping.h>
 using namespace rtac::cuda;
 
-struct Affine1D
-{
-    using Input = float;
+// This Functor transforms pixel coordinates in normalized texture coordinates,
+// given a size (width, height).
+struct NormalizerUV {
+    using InputT  = uint2;
+    using OutputT = float2;
 
-    float a;
-    float b;
+    uint2 shape;
 
-    #ifdef RTAC_CUDACC
-    __device__ float2 operator()(float x) const {
-        return float2({a*x+b, 0.0});
+    RTAC_HOSTDEVICE float2 operator()(uint2 pixCoords) const {
+        return float2({(2.0f* pixCoords.x) / (shape.x - 1),
+                       (2.0f* pixCoords.y) / (shape.y - 1)});
     }
-    #endif //RTAC_CUDACC
 };
 
-DeviceVector<float> map(const DeviceMapping<Affine1D,float>& mapping,
-                        const DeviceVector<float>& x);
+using Mapping1 = Mapping<float>;
+using Mapping2 = Mapping<float, NormalizerUV>;
+
+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_mapping2(int W, int H, const Mapping2::DeviceMap& map);
 
 #endif //_DEF_RTAC_BASE_CUDA_TESTS_MAPPING_TEST_H_

cuda/tests/mapping/src/mapping_test.hcu

+#ifndef _DEF_RTAC_BASE_CUDA_TESTS_MAPPING_TEST_HCU_
+#define _DEF_RTAC_BASE_CUDA_TESTS_MAPPING_TEST_HCU_
+
+#include <rtac_base/cuda/utils.h>
+
+template <typename T>
+__global__ void do_render_texture(T* output, int W, int H, cudaTextureObject_t texData)
+{
+    for(int h = 0; h < H; h++) {
+        float y = ((float)h) / (H - 1);
+        for(int w = 0; w < W; w++) {
+            float x = ((float)w) / (W - 1);
+            output[W*h + w] = tex2D<T>(texData, x, y);
+        }
+    }
+}
+
+template <typename T, class MappingT>
+__global__ void do_render_mapping1(T* output, int W, int H, MappingT map)
+{
+    float2 p;
+    for(int h = 0; h < H; h++) {
+        p.y = ((float)h) / (H - 1);
+        for(int w = 0; w < W; w++) {
+            p.x = 1.0f - ((float)w) / (W - 1);
+            output[W*h + w] = map(p);
+        }
+    }
+}
+
+template <typename T, class MappingT>
+__global__ void do_render_mapping2(T* output, int W, int H, MappingT map)
+{
+    uint2 p;
+    for(int h = 0; h < H; h++) {
+        p.y = h;
+        for(int w = 0; w < W; w++) {
+            p.x = w;
+            output[W*h + w] = map(p);
+        }
+    }
+}
+
+
+
+#endif //_DEF_RTAC_BASE_CUDA_TESTS_MAPPING_TEST_HCU_