[cuda] Added the Mapping type

cuda/CMakeLists.txt

@@ -12,6 +12,7 @@ list(APPEND rtac_cuda_headers
     include/rtac_base/cuda/DeviceMesh.h
     include/rtac_base/cuda/DeviceObject.h
     include/rtac_base/cuda/Texture2D.h
+    include/rtac_base/cuda/Mapping.h
     include/rtac_base/cuda/reductions.hcu
     include/rtac_base/cuda/operators.h
 )

cuda/include/rtac_base/cuda/Mapping.h

+#ifndef _DEF_RTAC_BASE_CUDA_MAPPING_H_
+#define _DEF_RTAC_BASE_CUDA_MAPPING_H_
+
+#include <rtac_base/cuda/Texture2D.h>
+
+namespace rtac { namespace cuda {
+
+/**
+ * This struct encodes a generic mapping from an arbitrary input space.  The
+ * mapping is encoded into the texture.
+ *
+ * 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.
+ */
+template <class FunctorT, typename T>
+struct DeviceMapping
+{
+    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);
+    }
+    #endif //RTAC_CUDACC
+};
+
+}; //namespace cuda
+}; //namespace rtac
+
+#ifndef RTAC_CUDACC
+
+namespace rtac { namespace cuda {
+
+/**
+ * Host side class with manage device side Device mapping data.
+ */
+template <class FunctorT, typename T>
+class Mapping
+{
+    public:
+
+    using Ptr       = std::shared_ptr<Mapping>;
+    using ConstPtr  = std::shared_ptr<const Mapping>;
+    using Texture   = Texture2D<T>;
+    using DeviceMap = DeviceMapping<FunctorT,T>;
+
+    protected:
+    
+    Texture  data_;
+    FunctorT f_;
+    
+    Mapping(const FunctorT& f, Texture&& data);
+
+    public:
+
+    static Ptr Create(const FunctorT& f, Texture&& data);
+    static Ptr Create(const FunctorT& f, const Texture& data);
+
+    void set_texture(const Texture& texture);
+    void set_texture(Texture&& texture);
+    void set_FunctorT(const FunctorT& a);
+
+    const Texture& texture() const;
+    const FunctorT& functor() const;
+
+    DeviceMap device_map() const;
+};
+
+template <class FunctorT, typename T>
+Mapping<FunctorT,T>::Mapping(const FunctorT& f, Texture&& data) :
+    data_(data),
+    f_(f)
+{}
+
+template <class FunctorT, typename T>
+typename Mapping<FunctorT,T>::Ptr 
+Mapping<FunctorT,T>::Create(const FunctorT& f, const Texture& data)
+{
+    return Ptr(new Mapping<FunctorT,T>(f, data));
+}
+
+template <class FunctorT, typename T>
+typename Mapping<FunctorT,T>::Ptr 
+Mapping<FunctorT,T>::Create(const FunctorT& f, Texture&& data)
+{
+    return Ptr(new Mapping<FunctorT,T>(f, std::move(Texture(data))));
+}
+
+template <class FunctorT, typename T>
+void Mapping<FunctorT,T>::set_texture(const Texture& texture)
+{
+    data_ = texture;
+}
+
+template <class FunctorT, typename T>
+void Mapping<FunctorT,T>::set_texture(Texture&& texture)
+{
+    data_ = texture;
+}
+
+template <class FunctorT, typename T>
+void Mapping<FunctorT,T>::set_FunctorT(const FunctorT& f)
+{
+    f_ = f;
+}
+
+template <class FunctorT, typename T>
+const Texture2D<T>& Mapping<FunctorT,T>::texture() const
+{
+    return data_;
+}
+
+template <class FunctorT, typename T>
+const FunctorT& Mapping<FunctorT,T>::functor() const
+{
+    return f_;
+}
+
+template <class FunctorT, typename T>
+typename Mapping<FunctorT,T>::DeviceMap Mapping<FunctorT,T>::device_map() const
+{
+    return DeviceMapping<FunctorT,T>({data_.texture(), f_});
+}
+
+}; //namespace cuda
+}; //namespace rtac
+
+#endif //RTAC_CUDACC
+
+#endif //_DEF_RTAC_BASE_CUDA_MAPPING_H_

cuda/tests/CMakeLists.txt

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

cuda/tests/mapping/CMakeLists.txt

+
+
+set(target_name ${PROJECT_NAME}_mapping_test)
+add_executable(${target_name}
+    src/mapping_test.cpp
+    src/mapping_test.cu
+)
+target_link_libraries(${target_name} PRIVATE
+    rtac_cuda
+)
+

cuda/tests/mapping/src/mapping_test.cpp

+#include <iostream>
+using namespace std;
+
+#include "mapping_test.h"
+
+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;
+    texture.set_wrap_mode(Texture2D<float>::WrapClamp);
+    texture.set_image(data.size(), 1, data.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;
+
+    return 0;
+}

cuda/tests/mapping/src/mapping_test.cu

+#include "mapping_test.h"
+
+
+__global__ void do_map(DeviceMapping<Affine1D,float> mapping, 
+                       const float* x, float* out, unsigned int size)
+{
+    for(int i = 0; i < size; i++) {
+        out[i] = mapping(x[i]);
+    }
+}
+
+DeviceVector<float> map(const DeviceMapping<Affine1D,float>& mapping,
+                        const DeviceVector<float>& x)
+{
+    DeviceVector<float> out(x.size());
+    
+    do_map<<<1,1>>>(mapping, x.data(), out.data(), out.size());
+    cudaDeviceSynchronize();
+
+    return out;
+}

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>
+using namespace rtac::cuda;
+
+struct Affine1D
+{
+    using Input = float;
+
+    float a;
+    float b;
+
+    #ifdef RTAC_CUDACC
+    __device__ float2 operator()(float x) const {
+        return float2({a*x+b, 0.0});
+    }
+    #endif //RTAC_CUDACC
+};
+
+DeviceVector<float> map(const DeviceMapping<Affine1D,float>& mapping,
+                        const DeviceVector<float>& x);
+
+#endif //_DEF_RTAC_BASE_CUDA_TESTS_MAPPING_TEST_H_