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.