Hi, sorry for the slow reply.
I have added the following code inside an APP SDK sample (SimpleImage), and it seemed to run very well with the current build of CodeXL (which will soon be released as CodeXL 1.2) and the Catalyst driver (Catalyst 13.4).
int i, j;
int k = 0;
__local float b[64];
for (i=0; i<8; i++)
for (j=0; j<8; j++)
b[k++] = read_imagef(input, imageSampler, (int2)(j, i)).x;
a. If the work items are indeed reading from the same image and writing to the same buffer, the performance I saw is fairly good. See 2, 3 below why this might be. Suggest making the kernel more effective by making it work in parallel, e.g:
int id = get_global_id(0)
j = (id / 8);
i = (id % 8);
b[id] = read_imagef(input, imageSampler, (int2)(j, i)).x;
b. If you prevent the work items from writing to the same location somehow (e.g initializing k to a different starting offset for each work item, such as 64 * get_global_id(0) and having b be a larger buffer), the performance I see is the same.
c. If the prevention is by doing something like using barriers or making k volatile - that in itself might cause the slowdown in stepping. That also would create a slowdown in the kernel when you're not debugging it, so I'd suggest moving away from that algorithm, as you're effectively serializing the operation and losing the benefits of OpenCL. Try using CodeXL's GPU profiling capabilities to see how to improve this.
2. If you were debugging through the CodeXL Visual Studio extension, specifically on Visual Studio 2012, there is a known issue causing slowdown on that version in CodeXL 1.1. Try using the standalone CodeXL application (CodeXL.exe) until CodeXL 1.2 (which fixes this) comes out.
3. Otherwise, the way CodeXL currently performs kernel debugging requires it to recompile the kernel as you step through it. This is a CPU intensive operation. Have a look at the application you're debugging in Windows' task manager during the 2-5 minute wait you're experiencing.
a. If the application is working and set to a high CPU percentage (usually this will be 100 / the number of cores your CPU has, e.g. 50%, 25% or 16%), the wait is probably in the compiler. Try upgrading to the latest Catalyst version or freeing up the CPU and see if that improves anything. Otherwise, if your CPU is not very new, it might simply be that it does not match up to the GPU's capabilities. You can also try using CodeXL's CPU profiler to see if your application itself is what's taking up the resources - if most profiling samples are inside the AMD drivers (C:\Windows\[System32|SysWOW64]\amd*.dll ), it's probably the compiler thing.
b. If the application is mostly sitting at 0% CPU, and only occasionally jumps up (or never jumps up at all), this is more likely a wait in the GPU. Either you are running many graphic / compute applications at the same time, or your kernel is too complex or too serialized (such as in case 1c above). Try making the kernel more parallelized or running it through CodeXL's GPU profiler to see what can be improved on this front.
Hope this information is useful for you, and sorry again for the slow reply.
Regards,
