[Beignet] assertion failed in llvm_gen_backend.cpp

Yang, Rong R rong.r.yang at intel.com
Wed Oct 16 22:24:34 PDT 2013


I just send a patch “Remove newValueProxy from scalarize pass to genWriter pass.” to fix this bug, can you apply and try it? Thanks.

From: beignet-bounces+rong.r.yang=intel.com at lists.freedesktop.org [mailto:beignet-bounces+rong.r.yang=intel.com at lists.freedesktop.org] On Behalf Of Matt Harvey
Sent: Sunday, October 13, 2013 6:48 AM
To: beignet at lists.freedesktop.org
Subject: [Beignet] assertion failed in llvm_gen_backend.cpp

Hi list,

When I run the kernel copied below, I find that it compiles just fine, but upon running emits the following error:

ASSERTION FAILED: scalarMap.find(key) != scalarMap.end()
  at file /media/Programming/src/beignet/backend/src/llvm/llvm_gen_backend.cpp, function gbe::ir::Register gbe::RegisterTranslator::getScalar(llvm::Value*, uint32_t), line 373
Stack dump:
0.  Running pass 'Function Pass Manager' on module '/tmp/fileLxplMp.ll'.
1.  Running pass 'Gen Back-End' on function '@edge_detect'

This looks like a bug to me.
Thanks,
Matt

Here is the kernel that caused this error:

__constant sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;

float calculate_intensity(float4 rgba)
{
    //equation taken from http://en.wikipedia.org/wiki/YUV#Converting_between_Y.27UV_and_RGB
    return (0.299f * rgba[0]) +
            (0.587f * rgba[1]) +
            (0.114f * rgba[2]);
}

__kernel void edge_detect(__read_only image2d_t image_in,
                        __write_only image2d_t image_out,
                        __read_only ulong image_width,
                        __read_only ulong image_height)
{
    const int2 pos = { get_global_id(0), get_global_id(1) };
    float4 output = (float4)(0.0f, 0.0f, 0.0f, 1.0f); //black
    if ((pos[0] != 0) && (pos[1] != 0) && (pos[0] != image_width) && (pos[1] != image_height))
    {
        float A = calculate_intensity(read_imagef(image_in, sampler, (int2)(pos[0] - 1, pos[1] - 1)));
        float F = calculate_intensity(read_imagef(image_in, sampler, (int2)(pos[0], pos[1] + 1)));
        float G = calculate_intensity(read_imagef(image_in, sampler, (int2)(pos[0] - 1, pos[1] + 1)));
        float H = calculate_intensity(read_imagef(image_in, sampler, (int2)(pos[0] - 1, pos[1])));
        float addFG = F + G;
        float addAH = A + H;
        float subtraction = addFG - addAH;
        if (subtraction > 383.0f)
        {
            output = (float4)(1.0f, 1.0f, 1.0f, 1.0f);
        }
    }
    write_imagef(image_out, pos, output);
}

-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.freedesktop.org/archives/beignet/attachments/20131017/b445333e/attachment-0001.html>


More information about the Beignet mailing list