[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