[Critical bug]Cat 10.12+SDK v2.3 hangs/resets computer

Discussion created by bubu on Dec 29, 2010
Latest reply on Jan 11, 2011 by jeff_golds

I've found a critical bug/hang. I simply write to an int2 R/W buffer the values of a const struct and I add "one" to the buffer.


I've packed the source code and executables for you at



It happens randomly so, please, run the executable several times until you get it ( usually 5-10 times from the MS-DOS console ). Usually the x86 version hangs faster than the x64.

I use Win7 x64, i7 920, Radeon 5750, Cat 10.12+SDK 2.3 normally.

Runs ok with Intel / NVIDIA OpenCL and Cat 10.9, so the problem should be a routine in the new Cat.10.10+ GPU code ( because the CPU one apparently runs ok ). Catalyst 10.10, 10.11 and 10.12 makes the hang to appear using GPU devices.


I'll put here the kernel too:


I've noticed in the SKA too a strange "LSHR" command. strange because I see several LSHL which I think are correct ( because the block/work sizes are pow2 ). That LSHR is probably computing bad the buffer offset and causing something to mess badly...


typedef struct tParams { int4 c1; int c21; int c22; int c23; int c24; }Params; __kernel __attribute__((reqd_work_group_size(WORK_WIDTH,WORK_HEIGHT,1))) void MyKernel ( __constant Params *params, __global int2 *outb ) { const size_t g1D = get_global_id(1U)*BLOCK_WIDTH + get_global_id(0U); size_t i; for ( i=0; i<ITERATIONS; ++i ) { const int val = params->c1.x + params->c21 + params->c22 + params->c23 + params->c24; int2 i2v; i2v.x = val; i2v.y = val; outb[g1D+i*BLOCK_SIZE] = i2v; outb[g1D+i*BLOCK_SIZE] += (int2)(1); } }