Added support for cl_ext_float_atomics in CBasicTestFetchAddSpecialFloats with atomic_half

[shajder]

Related to #2142, according to the work plan, extending CBasicTestFetchAddSpecialFloats with support for atomic_half.

[lakshmih]

Can this be rebased, thanks.

[shajder]

Can this be rebased, thanks.

Done

[bashbaug]

This test is "failing" on our device because the incoming values in "finalDest" are all NaN. I have a very time following these tests, so I'm not exactly sure where things are going wrong. Can you please take a look and see what might be happening?

Found by modifying the kernel source to print the initial values:

__kernel void test_atomic_kernel(uint threadCount, uint numDestItems,
                                 __global half *finalDest,
                                 __global half *oldValues,
                                 volatile __local atomic_half *destMemory) {
  uint tid = get_global_id(0);

  // initialize atomics not reachable from host (first thread
  // is doing this, other threads are waiting on barrier)
  if (get_local_id(0) == 0)
    for (uint dstItemIdx = 0; dstItemIdx < numDestItems; dstItemIdx++) {
      printf("At index %u: destMemory is %f\n", dstItemIdx, (float)finalDest[dstItemIdx]);
      atomic_store_explicit(destMemory + dstItemIdx, finalDest[dstItemIdx],
                            memory_order_relaxed, memory_scope_work_group);
    }
  barrier(CLK_LOCAL_MEM_FENCE);
...

This prints:

At index 0: destMemory is -nan
At index 1: destMemory is -nan
At index 2: destMemory is -nan
...

[shajder]

Can you please take a look and see what might be happening?

I can't find any implementation that supports either the CL_DEVICE_LOCAL_FP_ATOMIC_ADD_EXT or CL_DEVICE_GLOBAL_FP_ATOMIC_ADD_EXT capabilities, so it's difficult to test this feature. I’ll try it on an Arc B570 machine tomorrow, but if you know of any simpler way to access a device that exposes these capabilities, I would appreciate the information. Thanks!

[shajder]

This test is "failing" on our device because the incoming values in "finalDest" are all NaN. I have a very time following these tests, so I'm not exactly sure where things are going wrong. Can you please take a look and see what might be happening?

I can reproduce the problem now and I am working on it, thanks!

[shajder]

Can you please take a look and see what might be happening?

Corrected, please review.

[bashbaug]

I still think these tests are still quite hard to follow and perhaps overly clever, but I think they're correct now, and they execute fairly promptly.

[bashbaug]

Merging as discussed in the January 27th teleconference.