ahmed/OpenCL-CTS
1
0
Fork 0
mirror of https://github.com/KhronosGroup/OpenCL-CTS.git synced 2026-03-19 06:09:01 +00:00
Code Issues Packages Projects Releases Wiki Activity

cl20: Khronos Bug 15619 Race condition in device execution enqueue_flags test

Browse Source
This commit is contained in:
Bartosz Sochacki
2016-03-02 16:29:29 +01:00
committed by Kévin Petit
parent 4c31b69980
commit b99c6004ce
1 changed files with 20 additions and 18 deletions
Download Patch File Download Diff File Expand all files Collapse all files

38
test_conformance/device_execution/enqueue_flags.cpp
View File

@@ -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, ""