OpenCL 内核函数崩溃

Question

我在 OpenCL 中编写了一个代码，其中我没有使用本地（共享）内存。我的代码在执行过程中崩溃并给出错误-5。当我用一些常量值替换对 cvt_img 缓冲区（在代码中间）的全局内存访问时，错误消失了。

我不明白为什么会发生这种情况，因为我使用 if 语句阻止访问超出范围的内存位置。

此代码是 3D 管道的一部分，但现在，我已将其与我的主应用程序分开，并将其放在一个单独的项目中，其中所有缓冲区都被随机初始化。

网格的大小（以线程数计）与图像的大小（img_size.x, img_size.y）相同，块的大小为（16, 16）。该应用程序正在运行 15 个图像。

void compute_cost_volume( 全局浮点3 *cvt_img， 全局 float8 *spixl_map,
全局浮动 *disp_level, 全局 int *view_subset, 全局 int *subset_num, int array_width, int2 map_size, int2 img_size，浮动 bl_ratio， int sp_size, int num_disp, float2 step, int x, int y, int z, int view_count ) { 屏障（CLK_GLOBAL_MEM_FENCE）；

int idx = map_size.x * map_size.y * z + map_size.x * y + x;

float8 spixl = spixl_map[idx];
float2 center = spixl.s12;
int2 camIdx  = (int2)(z % array_width, z / array_width);
float cost_est = 1000000.0, disp_est = 0.0;


for (int dl = 0 ; dl < num_disp ; dl++)
{
    float d = disp_level[dl];
    float min_val = 1000000.0;

    for (int n = 0 ; n < subset_num[z] ; n++)
    {
        int view = view_subset[n];
        int2 viewIdx = (int2)(view % array_width, view / array_width);
        float val = 0.0;

        for (int i = -2 ; i <= 2 ; i++) for (int j = -2 ; j <= 2 ; j++)
        {
            //int2 xy_ref = (int2)(center.x - 2*step.x + i*step.x, center.y - 2*step.y + j*step.y);
            int2 xy_ref = (int2)(center.x + i*step.x, center.y + j*step.y);
            int2 xy_proj = (int2)((int)(xy_ref.x - d*(viewIdx.x - camIdx.x)), (int)(xy_ref.y - bl_ratio*d*(viewIdx.y - camIdx.y) ) );                   

            if (xy_ref.x >= 0 && xy_ref.y >= 0 && xy_proj.x >= 0 && xy_proj.y >= 0 && xy_ref.x < img_size.x && xy_ref.y < img_size.y && xy_proj.x < img_size.x  && xy_proj.y < img_size.y)
            {
                float3 color_ref  = cvt_img[img_size.x*img_size.y*z     + img_size.x*xy_ref.y  + xy_ref.x];
                float3 color_proj = cvt_img[img_size.x*img_size.y*view  + img_size.x*xy_proj.y + xy_proj.x];
                val += fabs(color_ref.x - color_proj.x) + fabs(color_ref.y - color_proj.y) + fabs(color_ref.z - color_proj.z);
            }
            else 
                val += 30;
        }
        if (val < min_val)
            min_val = val;
    }
    if (min_val < cost_est)
    {
        cost_est = min_val;
        disp_est = d;
    }
}

spixl_map[idx].s7 = disp_est;

}

内核无效初始深度估计（ 全局浮点3 *cvt_img， 全局 float8 *spixl_map, 全局浮动 *disp_level, int array_width, int2 map_size, int2 img_size，浮动 bl_ratio， int sp_size, int disp_num, 全局 int *view_subset，全局 int *subset_num ) {

int x = get_global_id(0);
int y = get_global_id(1);

if (x >= map_size.x || y >= map_size.y)
    return;

//float2 step = (float2)(1, 1);
for (int z = 0 ; z < 15 ; z++){

    int idx = map_size.x*map_size.y*z + map_size.x*y + x;

    // Set The Bounding Box

    float2 step = (float2)(1.0, 1.0);

    compute_cost_volume(cvt_img, spixl_map, disp_level, view_subset, subset_num, 
                                    array_width, map_size, img_size, bl_ratio, sp_size, disp_num, step, x, y, z, 15);

    barrier(CLK_LOCAL_MEM_FENCE);
}

}

Answer 1

来自文档

https://www.khronos.org/registry/OpenCL/sdk/1.0/docs/man/xhtml/vectorDataTypes.html

" 向量数据类型使用类型名称定义，即 char、uchar、short、ushort、int、uint、float、long 和 ulong，后跟定义向量中元素数量的字面值 n。支持n 的值为 2、4、8 和 16。"

所以没有float3，或许你可以尝试使用float4，让最后一个元素为零？

另外，假设float3存在，这行代码

float3 color_proj = cvt_img[img_size.x*img_size.y*view  + img_size.x*xy_proj.y + xy_proj.x];

没有做你想做的事，这将产生一个不能分配给向量的值，你应该使用类似的东西

float3 color_proj = (float3) cvt_img[img_size.x*img_size.y*view  + img_size.x*xy_proj.y + xy_proj.x];

这会将 cvt_img[...] 返回的一个值复制到 3 个向量元素。