arrayfire · 9prady9 · Mar 14, 2020 · Mar 4, 2020 · Mar 4, 2020
diff --git a/src/backend/cuda/CMakeLists.txt b/src/backend/cuda/CMakeLists.txt
@@ -375,6 +375,7 @@ cuda_add_library(afcuda
 
     Array.cpp
     Array.hpp
+    LookupTable1D.hpp
     Param.hpp
     anisotropic_diffusion.hpp
     approx.hpp

diff --git a/src/backend/cuda/LookupTable1D.hpp b/src/backend/cuda/LookupTable1D.hpp
@@ -0,0 +1,66 @@
+/*******************************************************
+ * Copyright (c) 2020, ArrayFire
+ * All rights reserved.
+ *
+ * This file is distributed under 3-clause BSD license.
+ * The complete license agreement can be obtained at:
+ * http://arrayfire.com/licenses/BSD-3-Clause
+ ********************************************************/
+
+#pragma once
+
+#include <Array.hpp>
+#include <err_cuda.hpp>
+
+#include <type_traits>
+
+namespace cuda {
+
+template<typename T>
+class LookupTable1D {
+   public:
+    LookupTable1D()                          = delete;
+    LookupTable1D(const LookupTable1D& arg)  = delete;
+    LookupTable1D(const LookupTable1D&& arg) = delete;
+    LookupTable1D& operator=(const LookupTable1D& arg) = delete;
+    LookupTable1D& operator=(const LookupTable1D&& arg) = delete;
+
+    LookupTable1D(const Array<T>& lutArray) : mTexture(0), mData(lutArray) {
+        cudaResourceDesc resDesc;
+        memset(&resDesc, 0, sizeof(resDesc));
+
+        cudaTextureDesc texDesc;
+        memset(&texDesc, 0, sizeof(texDesc));
+
+        resDesc.resType                = cudaResourceTypeLinear;
+        resDesc.res.linear.devPtr      = mData.get();
+        resDesc.res.linear.desc.x      = sizeof(T) * 8;
+        resDesc.res.linear.sizeInBytes = mData.elements() * sizeof(T);
+
+        if (std::is_signed<T>::value)
+            resDesc.res.linear.desc.f = cudaChannelFormatKindSigned;
+        else if (std::is_unsigned<T>::value)
+            resDesc.res.linear.desc.f = cudaChannelFormatKindUnsigned;
+        else
+            resDesc.res.linear.desc.f = cudaChannelFormatKindFloat;
+
+        texDesc.readMode = cudaReadModeElementType;
+
+        CUDA_CHECK(
+            cudaCreateTextureObject(&mTexture, &resDesc, &texDesc, NULL));
+    }
+
+    ~LookupTable1D() {
+        if (mTexture) { cudaDestroyTextureObject(mTexture); }
+    }
+
+    cudaTextureObject_t get() const noexcept { return mTexture; }
+
+   private:
+    // Keep a copy so that ref count doesn't go down to zero when
+    // original Array<T> goes out of scope before LookupTable1D object does.
+    Array<T> mData;
+    cudaTextureObject_t mTexture;
+};
+
+}  // namespace cuda
diff --git a/src/backend/cuda/fast.cu b/src/backend/cuda/fast.cu
@@ -7,11 +7,14 @@
  * http://arrayfire.com/licenses/BSD-3-Clause
  ********************************************************/
 
-#include <Array.hpp>
-#include <err_cuda.hpp>
+#include <fast.hpp>
+
+#include <LookupTable1D.hpp>
 #include <kernel/fast.hpp>
+#include <kernel/fast_lut.hpp>
 #include <af/dim4.hpp>
-#include <af/features.h>
+
+#include <mutex>
 
 using af::dim4;
 using af::features;
@@ -28,8 +31,14 @@ unsigned fast(Array<float> &x_out, Array<float> &y_out, Array<float> &score_out,
     float *d_y_out;
     float *d_score_out;
 
+    // TODO(pradeep) Figure out a better way to create lut Array only once
+    const Array<unsigned char> lut = createHostDataArray(
+        af::dim4(sizeof(FAST_LUT) / sizeof(unsigned char)), FAST_LUT);
+
+    LookupTable1D<unsigned char> fastLUT(lut);
+
     kernel::fast<T>(&nfeat, &d_x_out, &d_y_out, &d_score_out, in, thr,
-                    arc_length, non_max, feature_ratio, edge);
+                    arc_length, non_max, feature_ratio, edge, fastLUT);
 
     if (nfeat > 0) {
         const dim4 out_dims(nfeat);
@@ -38,7 +47,6 @@ unsigned fast(Array<float> &x_out, Array<float> &y_out, Array<float> &score_out,
         y_out     = createDeviceDataArray<float>(out_dims, d_y_out);
         score_out = createDeviceDataArray<float>(out_dims, d_score_out);
     }
-
     return nfeat;
 }
 

diff --git a/src/backend/cuda/kernel/fast.hpp b/src/backend/cuda/kernel/fast.hpp
@@ -9,14 +9,13 @@
 
 #pragma once
 
+#include <LookupTable1D.hpp>
 #include <common/dispatch.hpp>
 #include <debug_cuda.hpp>
-#include <err_cuda.hpp>
-#include <kernel/fast_lut.hpp>
+#include <kernel/shared.hpp>
 #include <math.hpp>
 #include <memory.hpp>
 #include <cub/block/block_reduce.cuh>
-#include "shared.hpp"
 
 namespace cuda {
 namespace kernel {
@@ -102,11 +101,16 @@ inline __device__ double abs_diff(const double x, const double y) {
     return fabs(x - y);
 }
 
+inline __device__ int lookup(const int n, cudaTextureObject_t tex) {
+    return (int)tex1Dfetch<unsigned char>(tex, n);
+}
+
 template<typename T, int arc_length>
 __device__ void locate_features_core(T *local_image, float *score,
                                      const unsigned idim0, const unsigned idim1,
                                      const float thr, int x, int y,
-                                     const unsigned edge) {
+                                     const unsigned edge,
+                                     cudaTextureObject_t luTable) {
     if (x >= idim0 - edge || y >= idim1 - edge) return;
 
     score[y * idim0 + x] = 0.f;
@@ -159,8 +163,8 @@ __device__ void locate_features_core(T *local_image, float *score,
 
     // Checks LUT to verify if there is a segment for which all pixels are much
     // brighter or much darker than central pixel p.
-    if ((int)FAST_LUT[bright] >= arc_length ||
-        (int)FAST_LUT[dark] >= arc_length)
+    if (lookup(bright, luTable) >= arc_length ||
+        lookup(dark, luTable) >= arc_length)
         score[x + idim0 * y] = max_val(s_bright, s_dark);
 }
 
@@ -187,7 +191,8 @@ __device__ void load_shared_image(CParam<T> in, T *local_image, unsigned ix,
 
 template<typename T, int arc_length>
 __global__ void locate_features(CParam<T> in, float *score, const float thr,
-                                const unsigned edge) {
+                                const unsigned edge,
+                                cudaTextureObject_t luTable) {
     unsigned ix = threadIdx.x;
     unsigned iy = threadIdx.y;
     unsigned bx = blockDim.x;
@@ -202,7 +207,7 @@ __global__ void locate_features(CParam<T> in, float *score, const float thr,
     load_shared_image(in, local_image_curr, ix, iy, bx, by, x, y, lx, ly, edge);
     __syncthreads();
     locate_features_core<T, arc_length>(local_image_curr, score, in.dims[0],
-                                        in.dims[1], thr, x, y, edge);
+                                        in.dims[1], thr, x, y, edge, luTable);
 }
 
 template<bool nonmax>
@@ -316,8 +321,8 @@ __global__ void get_features(float *x_out, float *y_out, float *score_out,
 template<typename T>
 void fast(unsigned *out_feat, float **x_out, float **y_out, float **score_out,
           const Array<T> &in, const float thr, const unsigned arc_length,
-          const unsigned nonmax, const float feature_ratio,
-          const unsigned edge) {
+          const unsigned nonmax, const float feature_ratio, const unsigned edge,
+          const LookupTable1D<unsigned char> &luTable) {
     dim4 indims             = in.dims();
     const unsigned max_feat = ceil(indims[0] * indims[1] * feature_ratio);
 
@@ -342,35 +347,43 @@ void fast(unsigned *out_feat, float **x_out, float **y_out, float **score_out,
     switch (arc_length) {
         case 9:
             CUDA_LAUNCH_SMEM((locate_features<T, 9>), blocks, threads,
-                             shared_size, in, d_score.get(), thr, edge);
+                             shared_size, in, d_score.get(), thr, edge,
+                             luTable.get());
             break;
         case 10:
             CUDA_LAUNCH_SMEM((locate_features<T, 10>), blocks, threads,
-                             shared_size, in, d_score.get(), thr, edge);
+                             shared_size, in, d_score.get(), thr, edge,
+                             luTable.get());
             break;
         case 11:
             CUDA_LAUNCH_SMEM((locate_features<T, 11>), blocks, threads,
-                             shared_size, in, d_score.get(), thr, edge);
+                             shared_size, in, d_score.get(), thr, edge,
+                             luTable.get());
             break;
         case 12:
             CUDA_LAUNCH_SMEM((locate_features<T, 12>), blocks, threads,
-                             shared_size, in, d_score.get(), thr, edge);
+                             shared_size, in, d_score.get(), thr, edge,
+                             luTable.get());
             break;
         case 13:
             CUDA_LAUNCH_SMEM((locate_features<T, 13>), blocks, threads,
-                             shared_size, in, d_score.get(), thr, edge);
+                             shared_size, in, d_score.get(), thr, edge,
+                             luTable.get());
             break;
         case 14:
             CUDA_LAUNCH_SMEM((locate_features<T, 14>), blocks, threads,
-                             shared_size, in, d_score.get(), thr, edge);
+                             shared_size, in, d_score.get(), thr, edge,
+                             luTable.get());
             break;
         case 15:
             CUDA_LAUNCH_SMEM((locate_features<T, 15>), blocks, threads,
-                             shared_size, in, d_score.get(), thr, edge);
+                             shared_size, in, d_score.get(), thr, edge,
+                             luTable.get());
             break;
         case 16:
             CUDA_LAUNCH_SMEM((locate_features<T, 16>), blocks, threads,
-                             shared_size, in, d_score.get(), thr, edge);
+                             shared_size, in, d_score.get(), thr, edge,
+                             luTable.get());
             break;
     }
 

diff --git a/src/backend/cuda/kernel/fast_lut.hpp b/src/backend/cuda/kernel/fast_lut.hpp
@@ -1,5 +1,5 @@
 /*******************************************************
- * Copyright (c) 2014, ArrayFire
+ * Copyright (c) 2020, ArrayFire
  * All rights reserved.
  *
  * This file is distributed under 3-clause BSD license.
@@ -9,7 +9,7 @@
 
 #pragma once
 
-__constant__ unsigned char FAST_LUT[] = {
+unsigned char FAST_LUT[] = {
     0,  0,  0,  0,  0,  0,  0,  0,  0,  0,  0,  0,  0,  0,  0,  0,  0,  0,  0,
     0,  0,  0,  0,  0,  0,  0,  0,  0,  0,  0,  0,  0,  0,  0,  0,  0,  0,  0,
     0,  0,  0,  0,  0,  0,  0,  0,  0,  0,  0,  0,  0,  0,  0,  0,  0,  0,  0,

diff --git a/src/backend/cuda/kernel/orb.hpp b/src/backend/cuda/kernel/orb.hpp
@@ -9,28 +9,26 @@
 
 #pragma once
 
+#include <LookupTable1D.hpp>
 #include <common/dispatch.hpp>
 #include <debug_cuda.hpp>
-#include <err_cuda.hpp>
+#include <kernel/convolve.hpp>
+#include <kernel/orb_patch.hpp>
+#include <kernel/range.hpp>
+#include <kernel/sort_by_key.hpp>
 #include <memory.hpp>
 
-#include "convolve.hpp"
-#include "orb_patch.hpp"
-#include "range.hpp"
-#include "sort_by_key.hpp"
-
 using std::unique_ptr;
 using std::vector;
 
 namespace cuda {
-
 namespace kernel {
 
-static const int THREADS   = 256;
-static const int THREADS_X = 16;
-static const int THREADS_Y = 16;
+constexpr int THREADS   = 256;
+constexpr int THREADS_X = 16;
+constexpr int THREADS_Y = 16;
 
-static const float PI_VAL = 3.14159265358979323846f;
+constexpr float PI_VAL = 3.14159265358979323846f;
 
 template<typename T>
 void gaussian1D(T* out, const int dim, double sigma = 0.0) {
@@ -213,12 +211,17 @@ inline __device__ T get_pixel(unsigned x, unsigned y, const float ori,
     return image.ptr[x * image.dims[0] + y];
 }
 
+inline __device__ int lookup(const int n, cudaTextureObject_t tex) {
+    return tex1Dfetch<int>(tex, n);
+}
+
 template<typename T>
 __global__ void extract_orb(unsigned* desc_out, const unsigned n_feat,
                             float* x_in_out, float* y_in_out,
                             const float* ori_in, float* size_out,
                             CParam<T> image, const float scl,
-                            const unsigned patch_size) {
+                            const unsigned patch_size,
+                            cudaTextureObject_t luTable) {
     unsigned f = blockDim.x * blockIdx.x + threadIdx.x;
 
     if (f < n_feat) {
@@ -240,13 +243,13 @@ __global__ void extract_orb(unsigned* desc_out, const unsigned n_feat,
             for (unsigned j = 0; j < 16; j++) {
                 // Get position from distribution pattern and values of points
                 // p1 and p2
-                int dist_x = d_ref_pat[i * 16 * 4 + j * 4];
-                int dist_y = d_ref_pat[i * 16 * 4 + j * 4 + 1];
+                int dist_x = lookup(i * 16 * 4 + j * 4, luTable);
+                int dist_y = lookup(i * 16 * 4 + j * 4 + 1, luTable);
                 T p1       = get_pixel(x, y, ori, size, dist_x, dist_y, image,
                                  patch_size);
 
-                dist_x = d_ref_pat[i * 16 * 4 + j * 4 + 2];
-                dist_y = d_ref_pat[i * 16 * 4 + j * 4 + 3];
+                dist_x = lookup(i * 16 * 4 + j * 4 + 2, luTable);
+                dist_y = lookup(i * 16 * 4 + j * 4 + 3, luTable);
                 T p2   = get_pixel(x, y, ori, size, dist_x, dist_y, image,
                                  patch_size);
 
@@ -274,7 +277,8 @@ void orb(unsigned* out_feat, float** d_x, float** d_y, float** d_score,
          vector<float*>& d_y_pyr, vector<unsigned>& lvl_best,
          vector<float>& lvl_scl, vector<Array<T>>& img_pyr,
          const float fast_thr, const unsigned max_feat, const float scl_fctr,
-         const unsigned levels, const bool blur_img) {
+         const unsigned levels, const bool blur_img,
+         const LookupTable1D<int>& luTable) {
     UNUSED(fast_thr);
     UNUSED(max_feat);
     UNUSED(scl_fctr);
@@ -381,7 +385,7 @@ void orb(unsigned* out_feat, float** d_x, float** d_y, float** d_score,
         blocks  = dim3(divup(feat_pyr[i], threads.x), 1);
         CUDA_LAUNCH((extract_orb<T>), blocks, threads, d_desc_lvl, feat_pyr[i],
                     d_x_lvl, d_y_lvl, d_ori_lvl, d_size_lvl, img_pyr[i],
-                    lvl_scl[i], patch_size);
+                    lvl_scl[i], patch_size, luTable.get());
         POST_LAUNCH_CHECK();
 
         // Store results to pyramids
@@ -446,5 +450,4 @@ void orb(unsigned* out_feat, float** d_x, float** d_y, float** d_score,
 }
 
 }  // namespace kernel
-
 }  // namespace cuda
diff --git a/src/backend/cuda/kernel/orb_patch.hpp b/src/backend/cuda/kernel/orb_patch.hpp
@@ -10,19 +10,18 @@
 #pragma once
 
 namespace cuda {
-namespace kernel {
 
 // Reference pattern, generated for a patch size of 31x31, as suggested by
 // original ORB paper
-#define REF_PAT_SIZE 31
-#define REF_PAT_SAMPLES 256
-#define REF_PAT_COORDS 4
-#define REF_PAT_LENGTH (REF_PAT_SAMPLES * REF_PAT_COORDS)
+constexpr unsigned REF_PAT_SIZE    = 31;
+constexpr unsigned REF_PAT_SAMPLES = 256;
+constexpr unsigned REF_PAT_COORDS  = 4;
+constexpr unsigned REF_PAT_LENGTH  = (REF_PAT_SAMPLES * REF_PAT_COORDS);
 
 // Current reference pattern was borrowed from OpenCV, a randomly generated
 // pattern will not achieve same quality as it must be trained like described
 // in sections 4.2 and 4.3 of the original ORB paper.
-__constant__ int d_ref_pat[REF_PAT_LENGTH] = {
+int d_ref_pat[REF_PAT_LENGTH] = {
     8,   -3,  9,   5,   4,   2,   7,   -12, -11, 9,   -8,  2,   7,   -12, 12,
     -13, 2,   -13, 2,   12,  1,   -7,  1,   6,   -2,  -10, -2,  -4,  -13, -13,
     -11, -8,  -13, -3,  -12, -9,  10,  4,   11,  9,   -13, -8,  -8,  -9,  -11,
@@ -94,6 +93,4 @@ __constant__ int d_ref_pat[REF_PAT_LENGTH] = {
     -1,  -6,  0,   -11,
 };
 
-}  // namespace kernel
-
 }  // namespace cuda