加权乘以 OpenCL，发现全范围？

Question

我正在编写内核以处理索引的加权乘法。我的意思是 i*vec[i]。我想尽可能高效地使用我的 计算单元 (CL_DEVICE_MAX_COMPUTE_UNITS) 而不会超出我的 工作组大小的最大值 (CL_KERNEL_WORK_GROUP_SIZE).

为此，我编写了以下内核：

wght_mul.cl:

#ifdef MU_ND

__kernel void kmu(__global const int* _hist, const ulong _total, __global int* _tmp)
{
    int id = get_global_id(0)*kercn;
    const int lid = get_local_id(0);
    const int gid = get_group_id(0);

    const int gsz = get_global_size(0);
    const int lsz = get_local_size(0);

    int sum = 0;

    #if kercn == 4
    int4 sum4 = (int4)(0.f);
    #elif kercn == 8
    int8 sum8 = (int8)(0.f);
    #endif

    for(int grain = CU*WGS*kercn;id<_total;id+=grain)
    {
        #if kercn == 1
        sum += *(_hist+id) * id;
        #elif kercn == 4
        int4 idx = (int4)(id,id+1,id+2,id+3);
        int4 val = vload4(0,_hist+id);

        val*=idx;

        sum4 += val;
        #elif kercn == 8
        int8 idx = (int8)(id,id+1,id+2,id+3,id+4,id+5,id+6,id+7);
        int8 val = vload8(0,_hist+id);

        val*=idx;

        sum8 += val;
        #else
        #pragma unroll
        for(int i =0;i<kercn;i++)
            sum += *(_hist+id+i) * (id+i);

        #endif
    }
    barrier(CLK_LOCAL_MEM_FENCE);

    #if kercn == 4
    sum = sum4.s0 + sum4.s1 + sum4.s2 + sum4.s3;
    #elif kercn == 8
    sum = sum8.s0 + sum8.s1 + sum8.s2 + sum8.s3 + sum8.s4 + sum8.s5 + sum8.s6 + sum8.s7;
    #endif

    *(_tmp+gid) += sum;

    printf("kercn %d\n",kercn);

    //printf("%d %d %d %d %d\n",id,lid,gid,gsz,lsz);
}

#endif

这是我想要的和我得到的之间的演示：

main.cpp :

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

#include <fstream>
#include <sstream>
#include <algorithm>

int main()
{
    cv::Mat vec = cv::Mat::zeros(1,65536,CV_32SC1);
    cv::Mat idx = cv::Mat::zeros(1,65536,CV_32SC1);

    std::iota(idx.begin<int>(),idx.end<int>(),0);

    std::for_each(vec.begin<int>(),vec.end<int>(),[&](int& v){ v = cv::theRNG().uniform(0,1000);});


    // MY CODE

    cv::UMat uvec;
    cv::UMat utmp;

    vec.copyTo(uvec);

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

    file_buf<<file_stream.rdbuf();

    cv::ocl::ProgramSource source(file_buf.str());

    const cv::ocl::Device& dev = cv::ocl::Device::getDefault();

    std::size_t wgs = dev.maxWorkGroupSize();
    std::size_t compute_units = dev.maxComputeUnits();

    std::size_t kercn = 8;

    std::size_t globalsize = wgs*compute_units;
    std::size_t localsize = wgs;


    cv::ocl::Kernel k("kmu",source,cv::format("-D MU_ND -D CU=%ld -D WGS=%ld -D kercn=%ld ",compute_units,wgs,kercn));

    k.run(1,&globalsize,&localsize,false);


    utmp = cv::UMat::zeros(compute_units,1,CV_32SC1);

    k.args(cv::ocl::KernelArg::PtrReadOnly(uvec),uvec.total(),cv::ocl::KernelArg::PtrWriteOnly(utmp));

    k.run(1,&globalsize,&localsize,false);

    std::cout<<"utmp : "<<utmp.getMat(cv::ACCESS_RW)<<std::endl;

    // RESULT EXPECTED

    cv::multiply(vec,idx,vec);

    std::cout<<"result expected : "<<cv::sum(vec)(0)<<" results using my kernel : "<<cv::sum(utmp)(0)<<std::endl;

    return EXIT_SUCCESS;
}

执行成功，但结果不是我所期望的。

所以我的问题是：

我的代码有什么问题？
我观察如果我将变量kercn设置为4我所有的计算单元都被使用，否则只有一半被使用。使用所有计算单元是否更好？

感谢您的帮助。

Answer 1

要知道内核函数中的局部变量使用的寄存器文件在space中是非常有限的。每个局部变量使用 8 个字节而不是 4 个字节会导致每个工作项需要更多的寄存器。 GPU 核心之间的共享寄存器文件不足以同时激活所有核心。

您可以使用 GPU Profiling 工具找出每种情况下所需的寄存器文件并对其进行优化。

加权乘以 OpenCL，发现全范围？

weighted multiply with OpenCL, uncovered full range?

c++

gpgpu

opencl

c++11

opencv3.0