hazeman

Does uav_cached works ?

Discussion created by hazeman on Apr 27, 2011
Latest reply on Oct 5, 2011 by MicahVillmow
uav_raw_load_id(n)_cached gives random values

I've created small IL program to test how "uav_cached" works. On linux ( ubuntu maverick 64 bit ) with driver 11.3 and 11.4 it returns almost random* values.

Example has uav0 as input. Changing to uav11 doesn't improve anything. With driver 11.3 uav11 works without _cached flag, with 11.4 it gives always random* values.

* Driver 11.3 returns (int)1, Driver 11.4 returns a lot of (int)1 and some random values

CAL++ kernel void kernel_A1() { uav_raw<uint4> in(0,UAV_CACHED); uav_raw<uint4> out(1); uint1 gid = get_global_id(0); uint4 tmp; tmp = in[16*gid]; out[16*gid] = tmp+1; } generated IL il_cs dcl_num_thread_per_group 128 dcl_raw_uav_id(0) dcl_raw_uav_id(1) dcl_literal l0, 0x0, 0x0, 0x0, 0x0 dcl_literal l3, 0x1, 0x1, 0x1, 0x1 dcl_literal l2, 0x10, 0x0, 0x0, 0x0 dcl_literal l1, 0xffffffff, 0xffffffff, 0xffffffff, 0xffffffff mov r1.x,vAbsTid.x umul r4.x,l2.x,r1.x uav_raw_load_id(0)_cached r7.xyzw,r4.x mov r8.xyzw,r7.xyzw mov r2,r8 iadd r11,r2,l3 umul r15.x,l2.x,r1.x uav_raw_store_id(1) mem.xyzw,r15.x,r11 end

Outcomes