bubu

compiler(and stream kernel analyzer) error

Discussion created by bubu on Jun 2, 2010
Latest reply on Jun 2, 2010 by MicahVillmow

The following code:

 

__constant sampler_t imgSampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;

typedef struct tParamBlock
{
    float4 a, b, c;
    float4 d;
    float4 e;
    float2 f;
    int g;

}ParamBlock;

__kernel void myShader
(
    __read_only image2d_t img0,
    __read_only image2d_t img1,
    __read_only image2d_t img2,
   
    __constant ParamBlock *PB,

    const int2 other,
   
    __global float4 *outImg
)
{
    const int x = get_global_id(0);
    const int y = get_global_id(1);
    const int width = get_global_size(0);

    float4 data = read_imagef ( img0, imgSampler, (int2)(x,y) );
    data.w = 0.0f;
    float4 res = normalize(data)*0.5f + 0.5f;
           
    outImg[y*width + x] = res;
}

crashes the compiler on clCreateProgramWithSource/clBuildProgram and also the Stream Kernel Analyzer with this strage error:

 

Stack dump:
0.      Program arguments: G:\ATIStream\bin\x86\llc -mcpu=gummy -mattr=mwgs-3-25
6-1-1,+byte_addressable_store,+images -regalloc=linearscan -mtriple=amdil-pc-amd
opencl g:\Temp\OCLDA3.tmp.bc -f -o g:\Temp\OCLDA3.tmp.il
1.      Running pass 'AMDIL Assembly Printer' on function '@__OpenCL_myShader_ke
rnel'
005AC188 (0x014A0000 0x014B1AE4 0x00000000 0x014C46E0)
01002801 (0x014B1AE4 0x00000000 0x014C46E0 0x00000000)
014A0000 (0x00000000 0x014C46E0 0x00000000 0x014C2AB8)
014B1AE4 (0x014C46E0 0x00000000 0x014C2AB8 0x00000000)

 

 

curiously, if I multiply the data by * (PB->a-PB->a+1.0f), which is equivalent to multiply to 1.0f, the error is solved:

float4 data = read_imagef ( img0, imgSampler, (int2)(x,y) )* (PB->a-PB->a+1.0f);

 

so it seems you've a problem with unused variables.

Outcomes