直方圖內核內存問題

Question

我想實現一個算法來處理超過256個分檔的圖像。

在這種情況下處理直方圖的主要問題來自於不可能將超過32 Kb作爲GPU中的本地選項卡分配。

我發現的每像素8位圖像的所有算法在本地使用固定大小的選項卡。 直方圖是該選項卡中的第一個進程，然後屏障啓動並最終與輸出向量進行相加。

我正在使用具有超過32K動態垃圾箱的紅外圖像。 所以我不能在GPU內分配一個固定大小的選項卡。

我的算法使用atomic_add爲了直接創建輸出直方圖。

我與OpenCV接口，所以爲了管理可能的飽和情況，我的箱子使用浮點。取決於GPU管理單精度或雙精度的能力。 OpenCV不管理數據類型爲unsigned int,long和unsigned long作爲矩陣類型。

我有一個錯誤...我認爲這個錯誤是一種分段錯誤。 幾天後，我仍然不知道什麼是錯的。

這裏是我的代碼：

histogram.cl：

#pragma OPENCL EXTENSION cl_khr_fp64: enable 
#pragma OPENCL EXTENSION cl_khr_int64_base_atomics: enable 


static void Atomic_Add_f64(__global double *val, double delta) 
{ 
     union { 

    double f; 
     ulong i; 
     } old; 
     union { 
     double f; 
     ulong i; 
     } new; 
     do { 
     old.f = *val; 
     new.f = old.f + delta; 
     } 
     while (atom_cmpxchg ((volatile __global ulong *)val, old.i, new.i) != old.i); 
    } 

static void Atomic_Add_f32(__global float *val, double delta) 
{ 
     union 
     { 
     float f; 
     uint i; 
     } old; 

     union 
     { 
     float f; 
     uint i; 
     } new; 

     do 
     { 
     old.f = *val; 
     new.f = old.f + delta; 
     } 
     while (atom_cmpxchg ((volatile __global ulong *)val, old.i, new.i) != old.i); 
    } 


__kernel void khist(
        __global const uchar* _src, 
        const int src_steps, 
        const int src_offset, 
        const int rows, 
        const int cols, 
        __global uchar* _dst, 
        const int dst_steps, 
        const int dst_offset) 
{ 


        const int gid = get_global_id(0); 

//     printf("This message has been printed from the OpenCL kernel %d \n",gid); 


        if(gid < rows) 
        { 
         __global const _Sty* src = (__global const _Sty*)_src; 
         __global _Dty* dst = (__global _Dty*) _dst; 

         const int src_step1 = src_steps/sizeof(_Sty); 
         const int dst_step1 = dst_steps/sizeof(_Dty); 

         src += mad24(gid,src_step1,src_offset); 
         dst += mad24(gid,dst_step1,dst_offset); 

         _Dty one = (_Dty)1; 

         for(int c=0;c<cols;c++) 
         { 
          const _Rty idx = (_Rty)(*(src+c+src_offset)); 


           ATOMIC_FUN(dst+idx+dst_offset,one); 
         } 

         } 
} 

功能Atomic_Add_f64直接來自here和there

的main.cpp

#include <opencv2/core.hpp> 
#include <opencv2/core/ocl.hpp> 

#include <fstream> 
#include <sstream> 

#include <chrono> 

    int main() 
    { 
     cv::Mat_<unsigned short> a(480,640); 

     cv::RNG rng(std::time(nullptr)); 

     std::for_each(a.begin(),a.end(),[&](unsigned short& v){ v = rng.uniform(0,100);}); 

     bool ret = false; 

     cv::String file_content; 

     { 
      std::ifstream file_stream("../test/histogram.cl"); 
      std::ostringstream file_buf; 

      file_buf<<file_stream.rdbuf(); 

      file_content = file_buf.str(); 
     } 

     int output_flag = cv::ocl::Device::getDefault().doubleFPConfig() == 0 ? CV_32F : CV_64F; 

     cv::String atomic_fun = output_flag == CV_32F ? "Atomic_Add_f32" : "Atomic_Add_f64"; 

     cv::ocl::ProgramSource source(file_content); 

    // std::cout<<source.source()<<std::endl; 

     cv::ocl::Kernel k; 

     cv::UMat src; 
     cv::UMat dst = cv::UMat::zeros(1,65536,output_flag); 

     a.copyTo(src); 

     atomic_fun = cv::format("-D _Sty=%s -D _Rty=%s -D _Dty=%s -D ATOMIC_FUN=%s", 
           cv::ocl::typeToStr(src.depth()), 
           cv::ocl::typeToStr(src.depth()), // this to manage case like a matrix of usigned short stored as a matrix of float. 
           cv::ocl::typeToStr(output_flag), 
           atomic_fun.c_str()); 


     ret = k.create("khist",source,atomic_fun); 

     std::cout<<"check create : "<<ret<<std::endl; 

     k.args(cv::ocl::KernelArg::ReadOnly(src),cv::ocl::KernelArg::WriteOnlyNoSize(dst)); 

     std::size_t sz = a.rows; 

     ret = k.run(1,&sz,nullptr,false); 

     std::cout<<"check "<<ret<<std::endl; 

     cv::Mat b; 

     dst.copyTo(b); 

std::copy_n(b.ptr<double>(0),101,std::ostream_iterator<double>(std::cout," ")); 
std::cout<<std::endl; 


     return EXIT_SUCCESS; 
    } 

Answer 1

你好我到達要解決這個問題。 我真的不知道問題來自哪裏。 但是，如果我想輸出作爲指針而不是矩陣它的工作。

我所做的更改是這些：

histogram.cl：

__kernel void khist(
        __global const uchar* _src, 
        const int src_steps, 
        const int src_offset, 
        const int rows, 
        const int cols, 
        __global _Dty* _dst) 
{ 


        const int gid = get_global_id(0); 

        if(gid < rows) 
        { 
         __global const _Sty* src = (__global const _Sty*)_src; 
         __global _Dty* dst = _dst; 

         const int src_step1 = src_steps/sizeof(_Sty); 


         src += mad24(gid,src_step1,src_offset); 


         ulong one = 1; 

         for(int c=0;c<cols;c++) 
         { 
          const _Rty idx = (_Rty)(*(src+c+src_offset)); 


           ATOMIC_FUN(dst+idx,one); 

         } 

         } 
} 

的main.cpp

k.args(cv::ocl::KernelArg::ReadOnly(src),cv::ocl::KernelArg::PtrWriteOnly(dst)); 

代碼的其餘部分是在這兩個文件中的相同。 對我來說它工作正常。

如果有人知道爲什麼它的輸出被聲明爲一個指針，而不是一個矢量（一行矩陣），我很感興趣。

不過，我的問題是修復:)。