maximmoroz

KernelAnalyzer crashes with Catalyst 11.6

Discussion created by maximmoroz on Jun 22, 2011
Latest reply on Jun 23, 2011 by maximmoroz

R6950, Catalyst 11.6, AMD APP 2.4, Windows 7 Ultimate

I am attaching the code it crashes with (I picked the most simple one). If I specify this kernel in clBuildProgram my application just crashes.

If I supply analyzer with simpler code it says "OpenCL Compile Error: Can't find the IL for Cayman".

The kernels (the ones which are compiled) went wild with 11.6 and produce wrong results.

I didn't expirience all this with Catalyst 11.5.

The tragic is that I am unable to rollback to 11.5. Well, CCC shows 11.5, but all the problems remain. And I am unable to rollback to 11.4 at all - I install 11.4 and the system shows that I have 11.5.

#define width 96 #define height 96 #define window_width 13 #define window_height 13 #define feature_map_count 2 #define input_entry_size 18432 #define intermediate_entry_size 9216 __kernel void SubtractiveLocalContrastAverage( const __global float * restrict intermediate2, __global float * restrict intermediate, const __global float * restrict weights, const __global float * restrict adjustments, const uint dimension_size ) { const uint dd = get_global_id(0); if (dd < dimension_size) { const uint entry_id = dd / (width * height); const uint average_id = dd % (width * height); const int initial_dest_y = average_id / width; const int initial_dest_x = average_id % width; const int dest_y = (int)initial_dest_y - (int)(window_height >> 1); const int dest_x = (int)initial_dest_x - (int)(window_width >> 1); const uint input_offset = intermediate_entry_size * entry_id; float sum = 0.0F; for(uint window_y = 0; window_y < window_height; window_y++) { const int y = dest_y + (int)window_y; if ((y >= 0) && (y < height)) { const uint input_offset2 = y * width + input_offset; const uint weights_offset = window_y * window_width; for(uint window_x = 0; window_x < window_width; window_x++) { int x = dest_x + (int)window_x; float w = weights[weights_offset + window_x]; const int a1 = (x >= 0); const int a2 = (x < width); w = (a1 & a2) ? w : 0.0F; x += a1 ? 0 : width; x -= a2 ? 0 : width; sum += intermediate2[input_offset2 + x] * w; } } } intermediate[intermediate_entry_size * entry_id + average_id] = sum; } }

Outcomes