直方图内核内存问题

Question

我正在尝试实现一种算法来处理超过 256 个 bin 的图像。

在这种情况下处理直方图的主要问题是无法在 GPU 中分配超过 32 Kb 作为本地选项卡。

我为每像素 8 位图像找到的所有算法都在本地使用固定大小的选项卡。 直方图是该选项卡中的第一个过程，然后是障碍物，最后与输出向量进行了相加。

我正在处理具有超过 32K 动态箱的 IR 图像。 所以我无法在 GPU 内分配固定大小的选项卡。

我的算法使用atomic_add 来直接创建输出直方图。

我正在与 OpenCV 交互，因此，为了管理可能出现的饱和情况，我的 bin 使用浮点数。取决于 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));

两个文件中的其余代码相同。 对我来说它工作正常。

如果有人知道为什么当输出被声明为指针而不是向量（一行矩阵）时它会起作用，我很感兴趣。

不过我的问题已经解决了:)。