src/cuda_functions.cu

   1 #include "cuda_functions.h"
   2
   3 __global__ void gaussian_correlation_kernel(float *data_in, float *data_out, float *xf_sqr_norm, float *yf_sqr_norm,
   4                                             int rows, int cols, int channels_per_scale, double sigma)
   5 {
   6     extern __shared__ float sdata[];
   7     int blockId = blockIdx.y * gridDim.x + blockIdx.x;
   8     int threadId = blockId * (blockDim.x + channels_per_scale / 2) + threadIdx.x;
   9
  10     sdata[threadIdx.x] = 0;
  11     sdata[threadIdx.x] = data_in[threadId] + data_in[threadId + blockDim.x];
  12     __syncthreads();
  13
  14     for (unsigned int s = (channels_per_scale / 2 + 1) / 2, old_s = channels_per_scale / 2; s > 0; s >>= 1) {
  15
  16         if (old_s & 1) s += 1;
  17
  18         if (threadIdx.x < s && threadIdx.x + s < old_s) {
  19             sdata[threadIdx.x] += sdata[threadIdx.x + s];
  20         }
  21         old_s = s;
  22         __syncthreads();
  23     }
  24
  25     if (threadIdx.x == 0) {
  26         float accumulate_res = sdata[0];
  27
  28         float numel_xf_inv = 1.f / ((cols / 2 + 1) * rows * (channels_per_scale));
  29
  30         float tmp = (xf_sqr_norm[blockIdx.x] + yf_sqr_norm[0] - 2 * accumulate_res) * numel_xf_inv;
  31
  32         if (tmp > 0) {
  33             data_out[blockIdx.x * rows * cols + blockIdx.y] = expf(-1.f / (sigma * sigma) * tmp);
  34         } else {
  35             data_out[blockIdx.x * rows * cols + blockIdx.y] = expf(0);
  36         }
  37     }
  38 }
  39
  40 void cuda_gaussian_correlation(float *data_in, float *data_out, float *xf_sqr_norm, float *yf_sqr_norm, double sigma,
  41                                int n_channels, int n_scales, int rows, int cols)
  42 {
  43     dim3 threadsPerBlock((n_channels / n_scales) / 2);
  44     dim3 numBlocks(n_scales, rows * cols);
  45
  46     gaussian_correlation_kernel<<<numBlocks, threadsPerBlock, ((n_channels / n_scales) / 2) * sizeof(float)>>>(
  47         data_in, data_out, xf_sqr_norm, yf_sqr_norm, rows, cols, n_channels / n_scales, sigma);
  48     CudaCheckError();
  49
  50     //    float *data_cpu = (float*) malloc(rows*cols*n_scales*sizeof(float));
  51     //    CudaSafeCall(cudaMemcpy(data_cpu, data_out, rows*cols*n_scales*sizeof(float), cudaMemcpyDeviceToHost));
  52     //    for (int j = 0; j < rows*n_scales; ++j) {
  53     //                for (int k = 0; k < cols-1; ++k)
  54     //                   std::cout  << data_cpu[j*cols  + k]  << ", ";
  55     //                std::cout << data_cpu[j*cols + cols-1] <<  std::endl;
  56     //            }
  57     //    std::cout << std::endl << std::endl;
  58     //    free(data_cpu);
  59     return;
  60 }