gistfile1.txt

shader: MESA_SHADER_KERNEL
local-size: 1, 1, 1 (variable)
shared-size: 0
inputs: 5
outputs: 0
uniforms: 0
shared: 0
decl_var uniform INTERP_MODE_NONE readonly isampler1D @0 (0, 0, 0)
decl_var uniform INTERP_MODE_NONE readonly usampler1D @1 (0, 0, 1)
decl_var uniform INTERP_MODE_NONE sampler @2 (0, 0, 0)
decl_var ubo INTERP_MODE_NONE uint[10] kernel_inputs (0, 0, 0)
decl_var ubo INTERP_MODE_NONE uint[11] kernel_work_properies (0, 0, 1)
decl_var shader_in INTERP_MODE_NONE uint64_t @3 (2.x, 16, 0)
decl_var shader_in INTERP_MODE_NONE uint64_t @4 (3.x, 24, 0)
decl_var shader_in INTERP_MODE_NONE uint64_t @5 (4.x, 32, 0)
decl_function __wrapped_sample_kernel (0 params) (entrypoint)

impl __wrapped_sample_kernel {
        block block_0:
        /* preds: */
        vec1 32 ssa_191 = load_const (0x00000000 /* 0.000000 */)
        vec1 32 ssa_385 = load_const (0x00000001 /* 0.000000 */)
        vec4 32 ssa_218 = intrinsic load_ubo_dxil (ssa_191, ssa_385) ()
        vec1 32 ssa_408 = mov ssa_218.y
        vec4 32 ssa_238 = intrinsic load_ubo_dxil (ssa_191, ssa_385) ()
        vec1 32 ssa_406 = mov ssa_238.w
        vec1 32 ssa_393 = load_const (0x00000002 /* 0.000000 */)
        vec4 32 ssa_258 = intrinsic load_ubo_dxil (ssa_191, ssa_393) ()
        vec1 32 ssa_404 = mov ssa_258.y
        vec1 32 ssa_13 = load_const (0x00000003 /* 0.000000 */)
        vec3 32 ssa_185 = intrinsic load_global_invocation_id () ()
        vec1 32 ssa_403 = mov ssa_185.x
        vec4 32 ssa_340 = vec4 ssa_185.x, ssa_185.x, ssa_185.x, ssa_185.x
        vec4 32 ssa_81 = txf ssa_340 (coord), ssa_191 (lod), 0 (texture), 0 (sampler)
        vec4 32 ssa_95 = txf ssa_191 (lod), ssa_403 (coord), 1 (texture), 0 (sampler)
        vec1 32 ssa_33 = ishl.nsw ssa_185.x, ssa_393
        vec1 32 ssa_401 = load_const (0x00000004 /* 0.000000 */)
        vec1 32 ssa_372 = ishl ssa_185.x, ssa_401
        vec1 32 ssa_176 = iadd ssa_218.x, ssa_372
        vec1 32 ssa_274 = load_const (0xfffffffc /* -nan */)
        vec1 32 ssa_275 = iand ssa_176, ssa_274
        vec1 32 ssa_278 = mov ssa_81.x
        intrinsic store_ssbo (ssa_278, ssa_408, ssa_275) (0, 0, 4, 0) /* wrmask= */ /* access=0 */ /* align_mul=4 */ /* align_offset=0 */
        vec1 32 ssa_167 = iadd ssa_238.z, ssa_372
        vec1 32 ssa_280 = iand ssa_167, ssa_274
        vec1 32 ssa_283 = mov ssa_95.x
        intrinsic store_ssbo (ssa_283, ssa_406, ssa_280) (0, 0, 4, 0) /* wrmask= */ /* access=0 */ /* align_mul=4 */ /* align_offset=0 */
        vec1 32 ssa_42 = ior ssa_33, ssa_385
        vec1 32 ssa_155 = ishl ssa_42, ssa_393
        vec1 32 ssa_158 = iadd ssa_218.x, ssa_155
        vec1 32 ssa_285 = iand ssa_158, ssa_274
        vec1 32 ssa_288 = mov ssa_81.y
        intrinsic store_ssbo (ssa_288, ssa_408, ssa_285) (0, 0, 4, 0) /* wrmask= */ /* access=0 */ /* align_mul=4 */ /* align_offset=0 */
        vec1 32 ssa_149 = iadd ssa_238.z, ssa_155
        vec1 32 ssa_290 = iand ssa_149, ssa_274
        vec1 32 ssa_293 = mov ssa_95.y
        intrinsic store_ssbo (ssa_293, ssa_406, ssa_290) (0, 0, 4, 0) /* wrmask= */ /* access=0 */ /* align_mul=4 */ /* align_offset=0 */
        vec1 32 ssa_50 = ior ssa_33, ssa_393
        vec1 32 ssa_137 = ishl ssa_50, ssa_393
        vec1 32 ssa_140 = iadd ssa_218.x, ssa_137
        vec1 32 ssa_295 = iand ssa_140, ssa_274
        vec1 32 ssa_298 = mov ssa_81.z
        intrinsic store_ssbo (ssa_298, ssa_408, ssa_295) (0, 0, 4, 0) /* wrmask= */ /* access=0 */ /* align_mul=4 */ /* align_offset=0 */
        vec1 32 ssa_131 = iadd ssa_238.z, ssa_137
        vec1 32 ssa_300 = iand ssa_131, ssa_274
        vec1 32 ssa_303 = mov ssa_95.z
        intrinsic store_ssbo (ssa_303, ssa_406, ssa_300) (0, 0, 4, 0) /* wrmask= */ /* access=0 */ /* align_mul=4 */ /* align_offset=0 */
        vec1 32 ssa_58 = ior ssa_33, ssa_13
        vec1 32 ssa_119 = ishl ssa_58, ssa_393
        vec1 32 ssa_122 = iadd ssa_218.x, ssa_119
        vec1 32 ssa_305 = iand ssa_122, ssa_274
        vec1 32 ssa_308 = mov ssa_81.w
        intrinsic store_ssbo (ssa_308, ssa_408, ssa_305) (0, 0, 4, 0) /* wrmask= */ /* access=0 */ /* align_mul=4 */ /* align_offset=0 */
        vec1 32 ssa_113 = iadd ssa_238.z, ssa_119
        vec1 32 ssa_310 = iand ssa_113, ssa_274
        vec1 32 ssa_313 = mov ssa_95.w
        intrinsic store_ssbo (ssa_313, ssa_406, ssa_310) (0, 0, 4, 0) /* wrmask= */ /* access=0 */ /* align_mul=4 */ /* align_offset=0 */
        vec1 1 ssa_341 = ine ssa_81.x, ssa_95.x
        vec1 1 ssa_342 = ine ssa_81.y, ssa_95.y
        vec1 1 ssa_343 = ine ssa_81.z, ssa_95.z
        vec1 1 ssa_344 = ine ssa_81.w, ssa_95.w
        vec1 1 ssa_357 = ior ssa_341, ssa_342
        vec1 1 ssa_361 = ior ssa_357, ssa_343
        vec1 1 ssa_375 = ior ssa_361, ssa_344
        vec1 32 ssa_101 = ishl ssa_185.x, ssa_393
        vec1 32 ssa_104 = iadd ssa_258.x, ssa_101
        vec1 32 ssa_319 = b2i32 ssa_375
        vec1 32 ssa_410 = load_const (0x00000000 /* 0.000000 */)
        vec1 32 ssa_411 = isub ssa_410, ssa_319
        vec1 32 ssa_315 = iand ssa_104, ssa_274
        intrinsic store_ssbo (ssa_411, ssa_404, ssa_315) (0, 0, 4, 0) /* wrmask= */ /* access=0 */ /* align_mul=4 */ /* align_offset=0 */
        /* succs: block_1 */
        block block_1:
}


const char *read1DKernelSourcePattern =
"__kernel void sample_kernel( read_only image1d_t input, sampler_t sampler, __global int *a, __global int *b, __global int *neq )\n"
"{\n"
"   int tidX = get_global_id(0);\n"
"   int offset = tidX;\n"
"   %s clr = read_image%s( input, tidX );\n"
"   %s clr2 = read_image%s( input, sampler, tidX );\n"
"   offset *= 4;\n"
"   a[offset] = clr.x;\n"
"   b[offset] = clr2.x;\n"
"   a[offset + 1] = clr.y;\n"
"   b[offset + 1] = clr2.y;\n"
"   a[offset + 2] = clr.z;\n"
"   b[offset + 2] = clr2.z;\n"
"   a[offset + 3] = clr.w;\n"
"   b[offset + 3] = clr2.w;\n"
"   int4 test = (clr != clr2);\n"
"   if ( test.x || test.y || test.z || test.w )\n"
"      neq[tidX] = -1;\n"
"   else\n"
"      neq[tidX] = 0;\n"
"}";

__kernel void sample_kernel( read_only image1d_t input, sampler_t sampler, __global int *a, __global int *b, __global int *neq )
{
int tidX = get_global_id(0);
int offset = tidX;
int4 clr = read_imagei( input, tidX );
int4 clr2 = read_imagei( input, sampler, tidX );
offset *= 4;
a[offset] = clr.x;
b[offset] = clr2.x;
a[offset + 1] = clr.y;
b[offset + 1] = clr2.y;
a[offset + 2] = clr.z;
b[offset + 2] = clr2.z;
a[offset + 3] = clr.w;
b[offset + 3] = clr2.w;
int4 test = (clr != clr2);
if ( test.x || test.y || test.z || test.w )
neq[tidX] = -1;
else
neq[tidX] = 0;
}