From b99c6004ce8ea7bf1ac3ff56e1c0a1431c93581b Mon Sep 17 00:00:00 2001 From: Bartosz Sochacki Date: Wed, 2 Mar 2016 16:29:29 +0100 Subject: [PATCH] cl20: Khronos Bug 15619 Race condition in device execution enqueue_flags test --- .../device_execution/enqueue_flags.cpp | 38 ++++++++++--------- 1 file changed, 20 insertions(+), 18 deletions(-) diff --git a/test_conformance/device_execution/enqueue_flags.cpp b/test_conformance/device_execution/enqueue_flags.cpp index cad996c6..6dd4e122 100644 --- a/test_conformance/device_execution/enqueue_flags.cpp +++ b/test_conformance/device_execution/enqueue_flags.cpp @@ -39,8 +39,6 @@ static const char* enqueue_flags_wait_kernel_simple[] = NL, " size_t lid = get_local_id(0);" NL, " size_t tid = get_global_id(0);" NL, "" - NL, " array[index * gs + tid] = array[(index - 1) * gs + tid] + 1;" - NL, "" NL, " if(tid == 0)" NL, " {" NL, " if((index + 1) < BITS_DEPTH)" @@ -52,6 +50,8 @@ static const char* enqueue_flags_wait_kernel_simple[] = NL, " }" NL, " }" NL, "" + NL, " array[index * gs + tid] = array[(index - 1) * gs + tid] + 1;" + NL, "" NL, " if((index + 1) == BITS_DEPTH)" NL, " {" NL, " barrier(CLK_GLOBAL_MEM_FENCE);" @@ -106,8 +106,6 @@ static const char* enqueue_flags_wait_kernel_event[] = NL, " size_t lid = get_local_id(0);" NL, " size_t tid = get_global_id(0);" NL, "" - NL, " array[index * gs + tid] = array[(index - 1) * gs + tid] + 1;" - NL, "" NL, " if(tid == 0)" NL, " {" NL, " if((index + 1) < BITS_DEPTH)" @@ -124,6 +122,8 @@ static const char* enqueue_flags_wait_kernel_event[] = NL, " }" NL, " }" NL, "" + NL, " array[index * gs + tid] = array[(index - 1) * gs + tid] + 1;" + NL, "" NL, " if((index + 1) == BITS_DEPTH)" NL, " {" NL, " barrier(CLK_GLOBAL_MEM_FENCE);" @@ -193,7 +193,6 @@ static const char* enqueue_flags_wait_kernel_local[] = NL, " val += sub_array[i];" NL, " val -= (tid == id)? 0: (id + index - 1);" NL, " }" - NL, " array[index * gs + tid] = val + 1;" NL, "" NL, " if(tid == 0)" NL, " {" @@ -206,6 +205,8 @@ static const char* enqueue_flags_wait_kernel_local[] = NL, " }" NL, " }" NL, "" + NL, " array[index * gs + tid] = val + 1;" + NL, "" NL, " if((index + 1) == BITS_DEPTH)" NL, " {" NL, " barrier(CLK_GLOBAL_MEM_FENCE);" @@ -269,7 +270,6 @@ static const char* enqueue_flags_wait_kernel_event_local[] = NL, " val += sub_array[i];" NL, " val -= (tid == id)? 0: (id + index - 1);" NL, " }" - NL, " array[index * gs + tid] = val + 1;" NL, "" NL, " if(tid == 0)" NL, " {" @@ -287,6 +287,8 @@ static const char* enqueue_flags_wait_kernel_event_local[] = NL, " }" NL, " }" NL, "" + NL, " array[index * gs + tid] = val + 1;" + NL, "" NL, " if((index + 1) == BITS_DEPTH)" NL, " {" NL, " barrier(CLK_GLOBAL_MEM_FENCE);" @@ -348,8 +350,6 @@ static const char* enqueue_flags_wait_work_group_simple[] = NL, " " NL, " if(gid == group_id)" NL, " {" - NL, " array[index * gs + tid] = array[(index - 1) * gs + tid] + 1;" - NL, " " NL, " if((index + 1) < BITS_DEPTH && lid == 0)" NL, " {" NL, " enqueue_kernel(get_default_queue(), CLK_ENQUEUE_FLAGS_WAIT_WORK_GROUP, ndrange_1D(gs, ls), " @@ -357,13 +357,15 @@ static const char* enqueue_flags_wait_work_group_simple[] = NL, " block_fn(array, index + 1, ls, res, gid);" NL, " });" NL, " }" + NL, " " + NL, " array[index * gs + tid] = array[(index - 1) * gs + tid] + 1;" NL, " }" NL, "" NL, " if((index + 1) == BITS_DEPTH)" NL, " {" NL, " barrier(CLK_GLOBAL_MEM_FENCE);" NL, "" - NL, " if(lid == 0)" + NL, " if(lid == 0 && gid == group_id)" NL, " {" NL, " res[gid] = 1;" NL, "" @@ -417,8 +419,6 @@ static const char* enqueue_flags_wait_work_group_event[] = NL, " " NL, " if(gid == group_id)" NL, " {" - NL, " array[index * gs + tid] = array[(index - 1) * gs + tid] + 1;" - NL, " " NL, " if((index + 1) < BITS_DEPTH && lid == 0)" NL, " {" NL, " clk_event_t block_evt;" @@ -431,6 +431,8 @@ static const char* enqueue_flags_wait_work_group_event[] = NL, " release_event(user_evt);" NL, " release_event(block_evt);" NL, " }" + NL, " " + NL, " array[index * gs + tid] = array[(index - 1) * gs + tid] + 1;" NL, " }" NL, "" NL, "" @@ -438,7 +440,7 @@ static const char* enqueue_flags_wait_work_group_event[] = NL, " {" NL, " barrier(CLK_GLOBAL_MEM_FENCE);" NL, "" - NL, " if(lid == 0)" + NL, " if(lid == 0 && gid == group_id)" NL, " {" NL, " res[gid] = 1;" NL, "" @@ -508,8 +510,6 @@ static const char* enqueue_flags_wait_work_group_local[] = NL, " " NL, " if(gid == group_id)" NL, " {" - NL, " array[index * gs + tid] = val + 1;" - NL, " " NL, " if((index + 1) < BITS_DEPTH && lid == 0)" NL, " {" NL, " enqueue_kernel(get_default_queue(), CLK_ENQUEUE_FLAGS_WAIT_WORK_GROUP, ndrange_1D(gs, ls), " @@ -517,6 +517,8 @@ static const char* enqueue_flags_wait_work_group_local[] = NL, " block_fn(array, index + 1, ls, res, sub_array, gid);" NL, " }, ls * sizeof(int));" NL, " }" + NL, " " + NL, " array[index * gs + tid] = val + 1;" NL, " }" NL, "" NL, "" @@ -524,7 +526,7 @@ static const char* enqueue_flags_wait_work_group_local[] = NL, " {" NL, " barrier(CLK_GLOBAL_MEM_FENCE);" NL, "" - NL, " if(lid == 0)" + NL, " if(lid == 0 && gid == group_id)" NL, " {" NL, " res[gid] = 1;" NL, "" @@ -589,8 +591,6 @@ static const char* enqueue_flags_wait_work_group_event_local[] = NL, "" NL, " if(gid == group_id)" NL, " {" - NL, " array[index * gs + tid] = val + 1;" - NL, " " NL, " if((index + 1) < BITS_DEPTH && lid == 0)" NL, " {" NL, " clk_event_t block_evt;" @@ -603,13 +603,15 @@ static const char* enqueue_flags_wait_work_group_event_local[] = NL, " release_event(user_evt);" NL, " release_event(block_evt);" NL, " }" + NL, " " + NL, " array[index * gs + tid] = val + 1;" NL, " }" NL, "" NL, " if((index + 1) == BITS_DEPTH)" NL, " {" NL, " barrier(CLK_GLOBAL_MEM_FENCE);" NL, "" - NL, " if(lid == 0)" + NL, " if(lid == 0 && gid == group_id)" NL, " {" NL, " res[gid] = 1;" NL, ""