error: calling a __host__ function from a __global__ function is not allowed

Question

I have written cuda function for dense sampling of feature points but i am getting error. My cuda code is given below. I am using cuda 7.5 toolkit.

#include <stdlib.h>
#include <stdio.h>
#include <math.h>
#include <opencv2/core/core.hpp>
#include <opencv2/highgui/highgui.hpp>
#include "opencv2/imgproc/imgproc.hpp"
#include <opencv2/gpu/gpu.hpp>
#include <opencv2/opencv.hpp>


using namespace cv::gpu;
using namespace cv;
using namespace std;

__global__ void densefun(std::vector<int>* d_counters,std::vector<Point2f>* d_points,int d_x_max,int d_y_max,int width, int min_distance)
{
  int i = blockDim.x * blockIdx.x + threadIdx.x;
  Point2f point = (*d_points)[i];
  int x = cvFloor(point.x);
  int y = cvFloor(point.y);
  //if(x >= d_x_max || y >= d_y_max)
      //continue;
  x /= min_distance;
  y /= min_distance;
  (*d_counters)[y*width+x]++;
}


void dense(std::vector<int>& counters,std::vector<Point2f>& points,int x_max,int y_max,int width)
{
  std::vector<int>* d_counters;
  std::vector<Point2f>* d_points;
  int min_distance=5; 
  cudaMalloc(&d_counters,counters.size());
  cudaMalloc(&d_points,points.size());
  cudaMemcpy(d_points, &points, points.size(), cudaMemcpyHostToDevice);
  densefun<<<1,points.size()>>>(d_counters,d_points,x_max,y_max,width,min_distance);
  cudaMemcpy(&counters, d_counters, counters.size(), cudaMemcpyDeviceToHost);
  cudaFree(d_counters);
  cudaFree(d_points);
}

Output:

/home/supriya/Desktop/5Dec/CalculateFV_merged_gpu_old/build/denseCuda.cu(28): error: calling a host function("cv::Point_ ::Point_") from a global function("densefun") is not allowed

/home/supriya/Desktop/5Dec/CalculateFV_merged_gpu_old/build/denseCuda.cu(28): error: calling a host function("std::vector , std::allocator > > ::operator []") from a global function("densefun") is not allowed

/home/supriya/Desktop/5Dec/CalculateFV_merged_gpu_old/build/denseCuda.cu(29) (col. 7): error: calling a host function("cvFloor") from a global function("densefun") is not allowed

/home/supriya/Desktop/5Dec/CalculateFV_merged_gpu_old/build/denseCuda.cu(30) (col. 7): error: calling a host function("cvFloor") from a global function("densefun") is not allowed

/home/supriya/Desktop/5Dec/CalculateFV_merged_gpu_old/build/denseCuda.cu(35): error: calling a host function("std::vector > ::operator []") from a global function("densefun") is not allowed

5 errors detected in the compilation of "/tmp/tmpxft_00000c85_00000000-7_denseCuda.cpp1.ii". CMake Error at testVideo_generated_denseCuda.cu.o.cmake:260 (message): Error generating file
/home/supriya/Desktop/5Dec/CalculateFV_merged_gpu_old/build/CMakeFiles/testVideo.dir//./testVideo_generated_denseCuda.cu.o

CMakeFiles/testVideo.dir/build.make:392: recipe for target 'CMakeFiles/testVideo.dir/./testVideo_generated_denseCuda.cu.o' failed make[2]: * [CMakeFiles/testVideo.dir/./testVideo_generated_denseCuda.cu.o] Error 1 CMakeFiles/Makefile2:130: recipe for target 'CMakeFiles/testVideo.dir/all' failed make[1]: * [CMakeFiles/testVideo.dir/all] Error 2 Makefile:76: recipe for target 'all' failed make: *** [all] Error 2

Answer 1

You cannot use C++ standard library, OpenCV or any other non-CUDA specific library inside a CUDA kernel.

Instead of std::vector you need to use raw pointers to arrays allocated on the device, instead of Point2f you need to use CUDA specific vector type float2, instead of cvFloor you need to use __device__ ​ floorf() and so on.