[Beignet] assertion failed in llvm_gen_backend.cpp

[Matt Harvey]

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/20131012/444ce3a0/attachment.html>