Does the Atomic Flag C11 Atomic Test Need Improved Coverage

#954 Skipped the `atomic_flag` test in the case that it required `atomic_scope_device` on a 3.0 device which optionally doesn't support this scope.

The purpose of this ticket is to clarify whether we now need to update this test to improve coverage.

For an implementation not supporting the `memory_scope_device` scope and supporting the minimum required atomic features, the following four kernels are still getting generated and run as part of this test:
1.
```
__kernel void test_atomic_kernel(uint threadCount, uint numDestItems,
__global int *finalDest,
__global int *oldValues,
volatile __local atomic_flag *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++) {
if (finalDest[dstItemIdx])
atomic_flag_test_and_set_explicit(destMemory + dstItemIdx,
memory_order_relaxed,
memory_scope_work_group);
else
atomic_flag_clear_explicit(destMemory + dstItemIdx,
memory_order_relaxed,
memory_scope_work_group);
}
barrier(CLK_LOCAL_MEM_FENCE);

uint cnt, stop = 0;
for (cnt = 0; !stop && cnt < threadCount;
cnt++) // each thread must find critical section where it is the first
// visitor
{
bool set = atomic_flag_test_and_set_explicit(
&destMemory[cnt], memory_order_relaxed, memory_scope_work_group);
atomic_work_item_fence(CLK_LOCAL_MEM_FENCE, memory_order_acquire,
memory_scope_work_group);
if (!set) {
uint csIndex = get_enqueued_local_size(0) * get_group_id(0) + cnt;
// verify that thread is the first visitor
if (oldValues[csIndex] == 1000000000) {
oldValues[csIndex] = tid; // set the winner id for this critical section
stop = 1;
}
atomic_work_item_fence(CLK_LOCAL_MEM_FENCE, memory_order_release,
memory_scope_work_group);
atomic_flag_clear_explicit(&destMemory[cnt], memory_order_relaxed,
memory_scope_work_group);
}
}

// Copy final values to host reachable buffer
barrier(CLK_LOCAL_MEM_FENCE);
if (get_local_id(0) == 0) // first thread in workgroup
for (uint dstItemIdx = 0; dstItemIdx < numDestItems; dstItemIdx++)

finalDest[dstItemIdx] = atomic_flag_test_and_set_explicit(
destMemory + dstItemIdx, memory_order_relaxed,
memory_scope_work_group);
}

```
2.
```
void test_atomic_function(uint tid, uint threadCount, uint numDestItems,
volatile __local atomic_flag *destMemory,
__global int *oldValues) {
uint cnt, stop = 0;
for (cnt = 0; !stop && cnt < threadCount;
cnt++) // each thread must find critical section where it is the first
// visitor
{
bool set = atomic_flag_test_and_set_explicit(
&destMemory[cnt], memory_order_relaxed, memory_scope_work_group);
atomic_work_item_fence(CLK_LOCAL_MEM_FENCE, memory_order_acquire,
memory_scope_work_group);
if (!set) {
uint csIndex = get_enqueued_local_size(0) * get_group_id(0) + cnt;
// verify that thread is the first visitor
if (oldValues[csIndex] == 1000000000) {
oldValues[csIndex] = tid; // set the winner id for this critical section
stop = 1;
}
atomic_work_item_fence(CLK_LOCAL_MEM_FENCE, memory_order_release,
memory_scope_work_group);
atomic_flag_clear_explicit(&destMemory[cnt], memory_order_relaxed,
memory_scope_work_group);
}
}
}

__kernel void test_atomic_kernel(uint threadCount, uint numDestItems,
__global int *finalDest,
__global int *oldValues,
volatile __local atomic_flag *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++) {
if (finalDest[dstItemIdx])
atomic_flag_test_and_set_explicit(destMemory + dstItemIdx,
memory_order_relaxed,
memory_scope_work_group);
else
atomic_flag_clear_explicit(destMemory + dstItemIdx,
memory_order_relaxed,
memory_scope_work_group);
}
barrier(CLK_LOCAL_MEM_FENCE);

test_atomic_function(tid, threadCount, numDestItems, destMemory, oldValues);

// Copy final values to host reachable buffer
barrier(CLK_LOCAL_MEM_FENCE);
if (get_local_id(0) == 0) // first thread in workgroup
for (uint dstItemIdx = 0; dstItemIdx < numDestItems; dstItemIdx++)

finalDest[dstItemIdx] = atomic_flag_test_and_set_explicit(
destMemory + dstItemIdx, memory_order_relaxed,
memory_scope_work_group);
}
```
3.
```
__kernel void
test_atomic_kernel(uint threadCount, uint numDestItems, __global int *finalDest,
__global int *oldValues,
volatile __local atomic_flag *notUseddestMemory) {
__local volatile atomic_flag destMemory[1024];
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++) {
if (finalDest[dstItemIdx])
atomic_flag_test_and_set_explicit(destMemory + dstItemIdx,
memory_order_relaxed,
memory_scope_work_group);
else
atomic_flag_clear_explicit(destMemory + dstItemIdx,
memory_order_relaxed,
memory_scope_work_group);
}
barrier(CLK_LOCAL_MEM_FENCE);

uint cnt, stop = 0;
for (cnt = 0; !stop && cnt < threadCount;
cnt++) // each thread must find critical section where it is the first
// visitor
{
bool set = atomic_flag_test_and_set_explicit(
&destMemory[cnt], memory_order_relaxed, memory_scope_work_group);
atomic_work_item_fence(CLK_LOCAL_MEM_FENCE, memory_order_acquire,
memory_scope_work_group);
if (!set) {
uint csIndex = get_enqueued_local_size(0) * get_group_id(0) + cnt;
// verify that thread is the first visitor
if (oldValues[csIndex] == 1000000000) {
oldValues[csIndex] = tid; // set the winner id for this critical section
stop = 1;
}
atomic_work_item_fence(CLK_LOCAL_MEM_FENCE, memory_order_release,
memory_scope_work_group);
atomic_flag_clear_explicit(&destMemory[cnt], memory_order_relaxed,
memory_scope_work_group);
}
}

// Copy final values to host reachable buffer
barrier(CLK_LOCAL_MEM_FENCE);
if (get_local_id(0) == 0) // first thread in workgroup
for (uint dstItemIdx = 0; dstItemIdx < numDestItems; dstItemIdx++)

finalDest[dstItemIdx] = atomic_flag_test_and_set_explicit(
destMemory + dstItemIdx, memory_order_relaxed,
memory_scope_work_group);
}
```
4.
```
void test_atomic_function(uint tid, uint threadCount, uint numDestItems,
volatile __local atomic_flag *destMemory,
__global int *oldValues) {
uint cnt, stop = 0;
for (cnt = 0; !stop && cnt < threadCount;
cnt++) // each thread must find critical section where it is the first
// visitor
{
bool set = atomic_flag_test_and_set_explicit(
&destMemory[cnt], memory_order_relaxed, memory_scope_work_group);
atomic_work_item_fence(CLK_LOCAL_MEM_FENCE, memory_order_acquire,
memory_scope_work_group);
if (!set) {
uint csIndex = get_enqueued_local_size(0) * get_group_id(0) + cnt;
// verify that thread is the first visitor
if (oldValues[csIndex] == 1000000000) {
oldValues[csIndex] = tid; // set the winner id for this critical section
stop = 1;
}
atomic_work_item_fence(CLK_LOCAL_MEM_FENCE, memory_order_release,
memory_scope_work_group);
atomic_flag_clear_explicit(&destMemory[cnt], memory_order_relaxed,
memory_scope_work_group);
}
}
}

__kernel void
test_atomic_kernel(uint threadCount, uint numDestItems, __global int *finalDest,
__global int *oldValues,
volatile __local atomic_flag *notUseddestMemory) {
__local volatile atomic_flag destMemory[1024];
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++) {
if (finalDest[dstItemIdx])
atomic_flag_test_and_set_explicit(destMemory + dstItemIdx,
memory_order_relaxed,
memory_scope_work_group);
else
atomic_flag_clear_explicit(destMemory + dstItemIdx,
memory_order_relaxed,
memory_scope_work_group);
}
barrier(CLK_LOCAL_MEM_FENCE);

test_atomic_function(tid, threadCount, numDestItems, destMemory, oldValues);

// Copy final values to host reachable buffer
barrier(CLK_LOCAL_MEM_FENCE);
if (get_local_id(0) == 0) // first thread in workgroup
for (uint dstItemIdx = 0; dstItemIdx < numDestItems; dstItemIdx++)

finalDest[dstItemIdx] = atomic_flag_test_and_set_explicit(
destMemory + dstItemIdx, memory_order_relaxed,
memory_scope_work_group);
}
```
The question is whether this is enough coverage.

Start with the atomic_flag test discussed here and issue #954; compare the four generated kernels and the unsupported memory_scope_device case. Determine whether they cover the intended minimum atomic features, then record the coverage decision or update the test accordingly and run the relevant CTS test.

Written by the indexing model from the issue text.