[ad_1]
\n<\/p>\n
Following is a naive method to achieve the desired functionality:<\/p>\n
Following is a sample code for the above mentioned process.<\/p>\n
#include <cstdio>\n#include <vector>\n#include <cuda_runtime.h>\n#include <thrust\/count.h>\n#include <thrust\/host_vector.h>\n#include <thrust\/device_vector.h>\n#include <thrust\/execution_policy.h>\n#include <opencv2\/opencv.hpp>\n\n\nstatic void _check_err(cudaError_t err, const char* file, int line)\n{\n if(err)\n {\n const char* err_str = cudaGetErrorString(err);\n\n printf(\"CUDA Error: %s\\nFile: %s\\nLine: %d\\n\", err_str, file, line);\n exit(EXIT_FAILURE);\n }\n}\n\n#define CHECK_ERR(err) _check_err((err), __FILE__, __LINE__)\n\n\n\n__global__ void kernel_find_indices(const unsigned char* input, int width, int height, int step, int2* indices)\n{\n const int x = blockIdx.x * blockDim.x + threadIdx.x;\n const int y = blockIdx.y * blockDim.y + threadIdx.y;\n\n if(x < width && y < height)\n {\n const int tidPixel = y * step + x;\n const int tidIndex = y * width + x;\n\n unsigned char value = input[tidPixel];\n\n int2 index_to_write;\n\n\n if(value)\n {\n \/\/Write actual index to pixels with non-zero value\n index_to_write.x = x;\n index_to_write.y = y;\n }\n else\n {\n \/\/Write dummy index to pixels with zero value\n index_to_write.x = -1;\n index_to_write.y = -1;\n }\n\n indices[tidIndex] = index_to_write;\n }\n}\n\n\n\/\/Operator to check whether an index is of a non-zero pixel\nstruct isNonZeroIndex\n{\n __host__ __device__ bool operator()(const int2 &idx)\n {\n return (idx.x != -1) && (idx.y != -1);\n }\n};\n\n\nstd::vector<cv::Point> getIndicesOfNonZeroPixels(cv::Mat input)\n{\n std::vector<int2> output_int2;\n std::vector<cv::Point> output;\n\n int pixelCount = input.cols * input.rows;\n size_t imageBytes= input.step * input.rows;\n\n unsigned char* image_d;\n thrust::device_vector<int2> index_buffer_d(pixelCount);\n\n \/\/Allocate device memory for input image\n CHECK_ERR(cudaMalloc(&image_d, imageBytes));\n \/\/Copy input image to device\n CHECK_ERR(cudaMemcpy(image_d, input.ptr(), imageBytes, cudaMemcpyHostToDevice));\n\n dim3 block(16,16);\n dim3 grid;\n grid.x = (input.cols + block.x - 1) \/ block.x;\n grid.y = (input.rows + block.y - 1) \/ block.y;\n\n \/\/Generate an index mask with dummy values for indices with zero pixel value\n kernel_find_indices<<<grid, block>>>(image_d, input.cols, input.rows, input.step, thrust::raw_pointer_cast(index_buffer_d.data()));\n CHECK_ERR(cudaDeviceSynchronize());\n\n int nonZeroCount = thrust::count_if(index_buffer_d.begin(), index_buffer_d.end(), isNonZeroIndex());\n\n \/\/Keep only those indices whose pixel value is non-zero (stream compaction)\n thrust::device_vector<int2> compacted(nonZeroCount);\n thrust::copy_if(index_buffer_d.begin(), index_buffer_d.end(), compacted.begin(), isNonZeroIndex());\n\n \/\/Copy non-zero pixel indices to host\n output_int2.resize(nonZeroCount);\n thrust::copy(compacted.begin(), compacted.end(), output_int2.begin());\n\n CHECK_ERR(cudaFree(image_d));\n\n \/\/Convert vector<int2> to vector<cv::Point>\n output.resize(nonZeroCount);\n for(size_t i=0; i<nonZeroCount; i++)\n output[i] = cv::Point(output_int2[i].x, output_int2[i].y);\n\n return output;\n}\n\nvoid run_test()\n{\n \/\/Generate a sample test image\n cv::Mat test = cv::Mat::zeros(100,100, CV_8UC1);\n cv::rectangle(test, cv::Rect(5,5,20,20), cv::Scalar::all(255), CV_FILLED);\n\n \/\/Get pixel indices of non-zero pixels\n std::vector<cv::Point> indices = getIndicesOfNonZeroPixels(test);\n\n \/\/Display those indices\n for(size_t i=0; i<indices.size(); i++)\n {\n printf(\"%d, %d\\n\", indices[i].x, indices[i].y);\n }\n\n \/\/Show image\n cv::imshow(\"Sample\", test);\n cv::waitKey();\n}\n\nint main(int argc, char** argv)\n{\n run_test();\n return 0;\n}\n<\/code><\/pre>\nCompilation Command<\/h1>\n\nnvcc -o nz nz.cu -arch=sm_61 -L\/usr\/local\/lib -lopencv_core
\n -lopencv_highgui -lopencv_imgproc<\/p>\n<\/blockquote>\n
Please keep in mind that this code is for image of type 8UC1<\/code> (8 bit, single channel) only. You can easily extend it to other data-types as required.<\/p>\n<\/p><\/div>\n