//
// Copyright (c) 2017 The Khronos Group Inc.
// 
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
//    http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
//
#include <stdio.h>
#include <string.h>
#include "harness/testHarness.h"
#include "harness/typeWrappers.h"

#include <vector>

#include "procs.h"
#include "utils.h"
#include <time.h>


#ifdef CL_VERSION_2_0
extern int gWimpyMode;
#define BITS_DEPTH 28

static const char* enqueue_flags_wait_kernel_simple[] =
{
    NL, "#define BITS_DEPTH " STRINGIFY_VALUE(BITS_DEPTH)
    NL, ""
    NL, "void block_fn(__global int* array, int index, size_t ls, size_t gs, __global int* res)"
    NL, "{"
    NL, "  int val = 0;"
    NL, "  size_t lid = get_local_id(0);"
    NL, "  size_t tid = get_global_id(0);"
    NL, ""
    NL, "  if(tid == 0)"
    NL, "  {"
    NL, "    if((index + 1) < BITS_DEPTH)"
    NL, "    {"
    NL, "      enqueue_kernel(get_default_queue(), CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange_1D(gs, ls), "
    NL, "      ^{"
    NL, "         block_fn(array, index + 1, ls, gs, res);"
    NL, "       });"
    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);"
    NL, ""
    NL, "    if(lid == 0)"
    NL, "    {"
    NL, "      size_t gid = get_group_id(0);"
    NL, "      res[gid] = 1;"
    NL, ""
    NL, "      for(int j = 0; j < BITS_DEPTH; j++)"
    NL, "      {"
    NL, "        for(int i = 0; i < ls; i++)"
    NL, "        {"
    NL, "          if(array[j * gs + ls * gid + i] != ((ls * gid + i) + j))"
    NL, "          {"
    NL, "            res[gid] = 2;"
    NL, "            break;"
    NL, "          }"
    NL, "        }"
    NL, "      }"
    NL, "    }"
    NL, "  }"
    NL, "}"
    NL, ""
    NL, "kernel void enqueue_flags_wait_kernel_simple(__global int* res, __global int* array)"
    NL, "{"
    NL, "  size_t ls  = get_local_size(0);"
    NL, "  size_t gs  = get_global_size(0);"
    NL, "  size_t tid  = get_global_id(0);"
    NL, ""
    NL, "  res[tid] = 0;"
    NL, "  array[tid] = tid;"
    NL, ""
    NL, "  if(tid == 0)"
    NL, "  {"
    NL, "    enqueue_kernel(get_default_queue(), CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange_1D(gs, ls), "
    NL, "    ^{"
    NL, "       block_fn(array, 1, ls, gs, res);"
    NL, "     });"
    NL, "  }"
    NL, "}"
    NL
};

static const char* enqueue_flags_wait_kernel_event[] =
{
    NL, "#define BITS_DEPTH " STRINGIFY_VALUE(BITS_DEPTH)
    NL, ""
    NL, "void block_fn(__global int* array, int index, size_t ls, size_t gs, __global int* res)"
    NL, "{"
    NL, "  int val = 0;"
    NL, "  size_t lid = get_local_id(0);"
    NL, "  size_t tid = get_global_id(0);"
    NL, ""
    NL, "  if(tid == 0)"
    NL, "  {"
    NL, "    if((index + 1) < BITS_DEPTH)"
    NL, "    {"
    NL, "      clk_event_t block_evt;"
    NL, "      clk_event_t user_evt = create_user_event();"
    NL, "      enqueue_kernel(get_default_queue(), CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange_1D(gs, ls), 1, &user_evt, &block_evt, "
    NL, "      ^{"
    NL, "         block_fn(array, index + 1, ls, gs, res);"
    NL, "       });"
    NL, "      set_user_event_status(user_evt, CL_COMPLETE);"
    NL, "      release_event(user_evt);"
    NL, "      release_event(block_evt);"
    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);"
    NL, ""
    NL, "    if(lid == 0)"
    NL, "    {"
    NL, "      size_t gid = get_group_id(0);"
    NL, "      res[gid] = 1;"
    NL, ""
    NL, "      for(int j = 0; j < BITS_DEPTH; j++)"
    NL, "      {"
    NL, "        for(int i = 0; i < ls; i++)"
    NL, "        {"
    NL, "          if(array[j * gs + ls * gid + i] != ((ls * gid + i) + j))"
    NL, "          {"
    NL, "            res[gid] = 2;"
    NL, "            break;"
    NL, "          }"
    NL, "        }"
    NL, "      }"
    NL, "    }"
    NL, "  }"
    NL, "}"
    NL, ""
    NL, "kernel void enqueue_flags_wait_kernel_event(__global int* res, __global int* array)"
    NL, "{"
    NL, "  size_t tid  = get_global_id(0);"
    NL, "  size_t gs = get_global_size(0);"
    NL, "  size_t ls = get_local_size(0);"
    NL, ""
    NL, "  res[tid] = 0;"
    NL, "  array[tid] = tid;"
    NL, ""
    NL, "  if(tid == 0)"
    NL, "  {"
    NL, "    clk_event_t block_evt;"
    NL, "    clk_event_t user_evt = create_user_event();"
    NL, "    enqueue_kernel(get_default_queue(), CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange_1D(gs, ls), 1, &user_evt, &block_evt, "
    NL, "    ^{"
    NL, "       block_fn(array, 1, ls, gs, res);"
    NL, "     });"
    NL, "    set_user_event_status(user_evt, CL_COMPLETE);"
    NL, "    release_event(user_evt);"
    NL, "    release_event(block_evt);"
    NL, "  }"
    NL, "}"
    NL
};

static const char* enqueue_flags_wait_kernel_local[] =
{
    NL, "#define BITS_DEPTH " STRINGIFY_VALUE(BITS_DEPTH)
    NL, ""
    NL, "void block_fn(__global int* array, int index, size_t ls, size_t gs, __global int* res, __local int* sub_array)"
    NL, "{"
    NL, "  int val = 0;"
    NL, "  size_t gid = get_group_id(0);"
    NL, "  size_t lid = get_local_id(0);"
    NL, "  size_t tid = get_global_id(0);"
    NL, ""
    NL, "  sub_array[lid] = array[(index - 1) * gs + tid];"
    NL, "  barrier(CLK_LOCAL_MEM_FENCE);"
    NL, ""
    NL, "  for(int i = 0; i < ls; i++)"
    NL, "  {"
    NL, "    int id = gid * ls + i;"
    NL, "    val += sub_array[i];"
    NL, "    val -= (tid == id)? 0: (id + index - 1);"
    NL, "  }"
    NL, ""
    NL, "  if(tid == 0)"
    NL, "  {"
    NL, "    if((index + 1) < BITS_DEPTH)"
    NL, "    {"
    NL, "      enqueue_kernel(get_default_queue(), CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange_1D(gs, ls), "
    NL, "      ^(__local void* sub_array){"
    NL, "        block_fn(array, index + 1, ls, gs, res, sub_array);"
    NL, "      }, ls * sizeof(int));"
    NL, "    }"
    NL, "  }"
    NL, ""
    NL, "  array[index * gs + tid] = val + 1;"
    NL, ""
    NL, "  if((index + 1) == BITS_DEPTH)"
    NL, "  {"
    NL, "    barrier(CLK_GLOBAL_MEM_FENCE);"
    NL, ""
    NL, "    if(lid == 0)"
    NL, "    {"
    NL, "      res[gid] = 1;"
    NL, ""
    NL, "      for(int j = 0; j < BITS_DEPTH; j++)"
    NL, "      {"
    NL, "        for(int i = 0; i < ls; i++)"
    NL, "        {"
    NL, "          if(array[j * gs + ls * gid + i] != ((ls * gid + i) + j))"
    NL, "          {"
    NL, "            res[gid] = 2;"
    NL, "            break;"
    NL, "          }"
    NL, "        }"
    NL, "      }"
    NL, "    }"
    NL, "  }"
    NL, "}"
    NL, ""
    NL, "kernel void enqueue_flags_wait_kernel_local(__global int* res, __global int* array)"
    NL, "{"
    NL, "  size_t ls  = get_local_size(0);"
    NL, "  size_t gs  = get_global_size(0);"
    NL, "  size_t tid  = get_global_id(0);"
    NL, ""
    NL, "  res[tid] = 0;"
    NL, "  array[tid] = tid;"
    NL, ""
    NL, "  if(tid == 0)"
    NL, "  {"
    NL, "    enqueue_kernel(get_default_queue(), CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange_1D(gs, ls), "
    NL, "    ^(__local void* sub_array){"
    NL, "      block_fn(array, 1, ls, gs, res, sub_array);"
    NL, "    }, ls * sizeof(int));"
    NL, "  }"
    NL, "}"
    NL
};

static const char* enqueue_flags_wait_kernel_event_local[] =
{
    NL, "#define BITS_DEPTH " STRINGIFY_VALUE(BITS_DEPTH)
    NL, ""
    NL, "void block_fn(__global int* array, int index, size_t ls, size_t gs, __global int* res, __local int* sub_array)"
    NL, "{"
    NL, "  int val = 0;"
    NL, "  size_t gid = get_group_id(0);"
    NL, "  size_t lid = get_local_id(0);"
    NL, "  size_t tid = get_global_id(0);"
    NL, ""
    NL, "  sub_array[lid] = array[(index - 1) * gs + tid];"
    NL, "  barrier(CLK_LOCAL_MEM_FENCE);"
    NL, ""
    NL, "  for(int i = 0; i < ls; i++)"
    NL, "  {"
    NL, "    int id = gid * ls + i;"
    NL, "    val += sub_array[i];"
    NL, "    val -= (tid == id)? 0: (id + index - 1);"
    NL, "  }"
    NL, ""
    NL, "  if(tid == 0)"
    NL, "  {"
    NL, "    if((index + 1) < BITS_DEPTH)"
    NL, "    {"
    NL, "      clk_event_t block_evt;"
    NL, "      clk_event_t user_evt = create_user_event();"
    NL, "      enqueue_kernel(get_default_queue(), CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange_1D(gs, ls), 1, &user_evt, &block_evt, "
    NL, "      ^(__local void* sub_array){"
    NL, "        block_fn(array, index + 1, ls, gs, res, sub_array);"
    NL, "      }, ls * sizeof(int));"
    NL, "      set_user_event_status(user_evt, CL_COMPLETE);"
    NL, "      release_event(user_evt);"
    NL, "      release_event(block_evt);"
    NL, "    }"
    NL, "  }"
    NL, ""
    NL, "  array[index * gs + tid] = val + 1;"
    NL, ""
    NL, "  if((index + 1) == BITS_DEPTH)"
    NL, "  {"
    NL, "    barrier(CLK_GLOBAL_MEM_FENCE);"
    NL, ""
    NL, "    if(lid == 0)"
    NL, "    {"
    NL, "      res[gid] = 1;"
    NL, ""
    NL, "      for(int j = 0; j < BITS_DEPTH; j++)"
    NL, "      {"
    NL, "        for(int i = 0; i < ls; i++)"
    NL, "        {"
    NL, "          if(array[j * gs + ls * gid + i] != ((ls * gid + i) + j))"
    NL, "          {"
    NL, "            res[gid] = 2;"
    NL, "            break;"
    NL, "          }"
    NL, "        }"
    NL, "      }"
    NL, "    }"
    NL, "  }"
    NL, "}"
    NL, ""
    NL, "kernel void enqueue_flags_wait_kernel_event_local(__global int* res, __global int* array)"
    NL, "{"
    NL, "  size_t ls  = get_local_size(0);"
    NL, "  size_t gs  = get_global_size(0);"
    NL, "  size_t tid  = get_global_id(0);"
    NL, ""
    NL, "  res[tid] = 0;"
    NL, "  array[tid] = tid;"
    NL, ""
    NL, "  if(tid == 0)"
    NL, "  {"
    NL, "    clk_event_t block_evt;"
    NL, "    clk_event_t user_evt = create_user_event();"
    NL, "    enqueue_kernel(get_default_queue(), CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange_1D(gs, ls), 1, &user_evt, &block_evt, "
    NL, "    ^(__local void* sub_array){"
    NL, "      block_fn(array, 1, ls, gs, res, sub_array);"
    NL, "    }, ls * sizeof(int));"
    NL, "    set_user_event_status(user_evt, CL_COMPLETE);"
    NL, "    release_event(user_evt);"
    NL, "    release_event(block_evt);"
    NL, "  }"
    NL, "}"
    NL
};

static const char* enqueue_flags_wait_work_group_simple[] =
{
    NL, "#define BITS_DEPTH " STRINGIFY_VALUE(BITS_DEPTH)
    NL, ""
    NL, "void block_fn(__global int* array, int index, size_t ls, __global int* res, int group_id)"
    NL, "{"
    NL, "  size_t tid = get_global_id(0);"
    NL, "  size_t lid = get_local_id(0);"
    NL, "  size_t gs = get_global_size(0);"
    NL, "  size_t gid = get_group_id(0);"
    NL, "  "
    NL, "  if(gid == group_id)"
    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), "
    NL, "      ^{"
    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 && gid == group_id)"
    NL, "    {"
    NL, "      res[gid] = 1;"
    NL, ""
    NL, "      for(int j = 0; j < BITS_DEPTH; j++)"
    NL, "      {"
    NL, "        for(int i = 0; i < ls; i++)"
    NL, "        {"
    NL, "          if(array[j * gs + ls * gid + i] != ((ls * gid + i) + j))"
    NL, "          {"
    NL, "            res[gid] = 2;"
    NL, "            break;"
    NL, "          }"
    NL, "        }"
    NL, "      }"
    NL, "    }"
    NL, "  }"
    NL, "}"
    NL, ""
    NL, "kernel void enqueue_flags_wait_work_group_simple(__global int* res, __global int* array)"
    NL, "{"
    NL, "  size_t ls = get_local_size(0);"
    NL, "  size_t gs = get_global_size(0);"
    NL, "  size_t tid = get_global_id(0);"
    NL, "  size_t gid = get_group_id(0);"
    NL, "  size_t lid = get_local_id(0);"
    NL, ""
    NL, "  res[tid] = 0;"
    NL, "  array[tid] = tid;"
    NL, ""
    NL, "  if(lid == 0)"
    NL, "  {"
    NL, "    enqueue_kernel(get_default_queue(), CLK_ENQUEUE_FLAGS_WAIT_WORK_GROUP, ndrange_1D(gs, ls), "
    NL, "    ^{"
    NL, "       block_fn(array, 1, ls, res, gid);"
    NL, "     });"
    NL, "  }"
    NL, "}"
    NL
};

static const char* enqueue_flags_wait_work_group_event[] =
{
    NL, "#define BITS_DEPTH " STRINGIFY_VALUE(BITS_DEPTH)
    NL, ""
    NL, "void block_fn(__global int* array, int index, size_t ls, __global int* res, int group_id)"
    NL, "{"
    NL, "  size_t tid = get_global_id(0);"
    NL, "  size_t lid = get_local_id(0);"
    NL, "  size_t gs = get_global_size(0);"
    NL, "  size_t gid = get_group_id(0);"
    NL, "  "
    NL, "  if(gid == group_id)"
    NL, "  {"
    NL, "    if((index + 1) < BITS_DEPTH && lid == 0)"
    NL, "    {"
    NL, "      clk_event_t block_evt;"
    NL, "      clk_event_t user_evt = create_user_event();"
    NL, "      enqueue_kernel(get_default_queue(), CLK_ENQUEUE_FLAGS_WAIT_WORK_GROUP, ndrange_1D(gs, ls), 1, &user_evt, &block_evt, "
    NL, "      ^{"
    NL, "         block_fn(array, index + 1, ls, res, gid);"
    NL, "       });"
    NL, "      set_user_event_status(user_evt, CL_COMPLETE);"
    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, ""
    NL, "  if((index + 1) == BITS_DEPTH)"
    NL, "  {"
    NL, "    barrier(CLK_GLOBAL_MEM_FENCE);"
    NL, ""
    NL, "    if(lid == 0 && gid == group_id)"
    NL, "    {"
    NL, "      res[gid] = 1;"
    NL, ""
    NL, "      for(int j = 0; j < BITS_DEPTH; j++)"
    NL, "      {"
    NL, "        for(int i = 0; i < ls; i++)"
    NL, "        {"
    NL, "          if(array[j * gs + ls * gid + i] != ((ls * gid + i) + j))"
    NL, "          {"
    NL, "            res[gid] = 2;"
    NL, "            break;"
    NL, "          }"
    NL, "        }"
    NL, "      }"
    NL, "    }"
    NL, "  }"
    NL, "}"
    NL, ""
    NL, "kernel void enqueue_flags_wait_work_group_event(__global int* res, __global int* array)"
    NL, "{"
    NL, "  size_t ls = get_local_size(0);"
    NL, "  size_t gs = get_global_size(0);"
    NL, "  size_t tid = get_global_id(0);"
    NL, "  size_t gid = get_group_id(0);"
    NL, "  size_t lid = get_local_id(0);"
    NL, ""
    NL, "  res[tid] = 0;"
    NL, "  array[tid] = tid;"
    NL, ""
    NL, "  if(lid == 0)"
    NL, "  {"
    NL, "    clk_event_t block_evt;"
    NL, "    clk_event_t user_evt = create_user_event();"
    NL, "    enqueue_kernel(get_default_queue(), CLK_ENQUEUE_FLAGS_WAIT_WORK_GROUP, ndrange_1D(gs, ls), 1, &user_evt, &block_evt, "
    NL, "    ^{"
    NL, "       block_fn(array, 1, ls, res, gid);"
    NL, "     });"
    NL, "    set_user_event_status(user_evt, CL_COMPLETE);"
    NL, "    release_event(user_evt);"
    NL, "    release_event(block_evt);"
    NL, "  }"
    NL, "}"
    NL
};

static const char* enqueue_flags_wait_work_group_local[] =
{
    NL, "#define BITS_DEPTH " STRINGIFY_VALUE(BITS_DEPTH)
    NL, ""
    NL, "void block_fn(__global int* array, int index, size_t ls, __global int* res, __local int* sub_array, int group_id)"
    NL, "{"
    NL, "  int val = 0;"
    NL, "  size_t gid = get_group_id(0);"
    NL, "  size_t lid = get_local_id(0);"
    NL, "  size_t tid = get_global_id(0);"
    NL, "  size_t gs = get_global_size(0);"
    NL, ""
    NL, "  sub_array[lid] = array[(index - 1) * gs + tid];"
    NL, "  barrier(CLK_LOCAL_MEM_FENCE);"
    NL, ""
    NL, "  for(int i = 0; i < ls; i++)"
    NL, "  {"
    NL, "    int id = gid * ls + i;"
    NL, "    val += sub_array[i];"
    NL, "    val -= (tid == id)? 0: (id + index - 1);"
    NL, "  }"
    NL, " "
    NL, "  if(gid == group_id)"
    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), "
    NL, "      ^(__local void* sub_array){"
    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, ""
    NL, "  if((index + 1) == BITS_DEPTH)"
    NL, "  {"
    NL, "    barrier(CLK_GLOBAL_MEM_FENCE);"
    NL, ""
    NL, "    if(lid == 0 && gid == group_id)"
    NL, "    {"
    NL, "      res[gid] = 1;"
    NL, ""
    NL, "      for(int j = 0; j < BITS_DEPTH; j++)"
    NL, "      {"
    NL, "        for(int i = 0; i < ls; i++)"
    NL, "        {"
    NL, "          if(array[j * gs + ls * gid + i] != ((ls * gid + i) + j))"
    NL, "          {"
    NL, "            res[gid] = 2;"
    NL, "            break;"
    NL, "          }"
    NL, "        }"
    NL, "      }"
    NL, "    }"
    NL, "  }"
    NL, "}"
    NL, ""
    NL, "kernel void enqueue_flags_wait_work_group_local(__global int* res, __global int* array)"
    NL, "{"
    NL, "  size_t ls  = get_local_size(0);"
    NL, "  size_t gs = get_global_size(0);"
    NL, "  size_t tid  = get_global_id(0);"
    NL, "  size_t gid = get_group_id(0);"
    NL, "  size_t lid = get_local_id(0);"
    NL, ""
    NL, "  res[tid] = 0;"
    NL, "  array[tid] = tid;"
    NL, ""
    NL, "  if(lid == 0)"
    NL, "  {"
    NL, "    enqueue_kernel(get_default_queue(), CLK_ENQUEUE_FLAGS_WAIT_WORK_GROUP, ndrange_1D(gs, ls), "
    NL, "    ^(__local void* sub_array){"
    NL, "      block_fn(array, 1, ls, res, sub_array, gid);"
    NL, "    }, ls * sizeof(int));"
    NL, "  }"
    NL, "}"
    NL
};

static const char* enqueue_flags_wait_work_group_event_local[] =
{
    NL, "#define BITS_DEPTH " STRINGIFY_VALUE(BITS_DEPTH)
    NL, ""
    NL, "void block_fn(__global int* array, int index, size_t ls, __global int* res, __local int* sub_array, int group_id)"
    NL, "{"
    NL, "  int val = 0;"
    NL, "  size_t gid = get_group_id(0);"
    NL, "  size_t lid = get_local_id(0);"
    NL, "  size_t tid = get_global_id(0);"
    NL, "  size_t gs = get_global_size(0);"
    NL, ""
    NL, "  sub_array[lid] = array[(index - 1) * gs + tid];"
    NL, "  barrier(CLK_LOCAL_MEM_FENCE);"
    NL, ""
    NL, "  for(int i = 0; i < ls; i++)"
    NL, "  {"
    NL, "    int id = gid * ls + i;"
    NL, "    val += sub_array[i];"
    NL, "    val -= (tid == id)? 0: (id + index - 1);"
    NL, "  }"
    NL, ""
    NL, "  if(gid == group_id)"
    NL, "  {"
    NL, "    if((index + 1) < BITS_DEPTH && lid == 0)"
    NL, "    {"
    NL, "      clk_event_t block_evt;"
    NL, "      clk_event_t user_evt = create_user_event();"
    NL, "      enqueue_kernel(get_default_queue(), CLK_ENQUEUE_FLAGS_WAIT_WORK_GROUP, ndrange_1D(gs, ls), 1, &user_evt, &block_evt, "
    NL, "      ^(__local void* sub_array){"
    NL, "        block_fn(array, index + 1, ls, res, sub_array, gid);"
    NL, "      }, ls * sizeof(int));"
    NL, "      set_user_event_status(user_evt, CL_COMPLETE);"
    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 && gid == group_id)"
    NL, "    {"
    NL, "      res[gid] = 1;"
    NL, ""
    NL, "      for(int j = 0; j < BITS_DEPTH; j++)"
    NL, "      {"
    NL, "        for(int i = 0; i < ls; i++)"
    NL, "        {"
    NL, "          if(array[j * gs + ls * gid + i] != ((ls * gid + i) + j))"
    NL, "          {"
    NL, "            res[gid] = 2;"
    NL, "            break;"
    NL, "          }"
    NL, "        }"
    NL, "      }"
    NL, "    }"
    NL, "  }"
    NL, "}"
    NL, ""
    NL, "kernel void enqueue_flags_wait_work_group_event_local(__global int* res, __global int* array)"
    NL, "{"
    NL, "  size_t ls  = get_local_size(0);"
    NL, "  size_t gs  = get_global_size(0);"
    NL, "  size_t tid  = get_global_id(0);"
    NL, "  size_t gid = get_group_id(0);"
    NL, "  size_t lid  = get_local_id(0);"
    NL, ""
    NL, "  res[tid] = 0;"
    NL, "  array[tid] = tid;"
    NL, ""
    NL, "  if(lid == 0)"
    NL, "  {"
    NL, "    clk_event_t block_evt;"
    NL, "    clk_event_t user_evt = create_user_event();"
    NL, "    enqueue_kernel(get_default_queue(), CLK_ENQUEUE_FLAGS_WAIT_WORK_GROUP, ndrange_1D(gs, ls), 1, &user_evt, &block_evt, "
    NL, "    ^(__local void* sub_array){"
    NL, "      block_fn(array, 1, ls, res, sub_array, gid);"
    NL, "    }, ls * sizeof(int));"
    NL, "    set_user_event_status(user_evt, CL_COMPLETE);"
    NL, "    release_event(user_evt);"
    NL, "    release_event(block_evt);"
    NL, "  }"
    NL, "}"
    NL
};

static const kernel_src sources_enqueue_block_flags[] =
{
    KERNEL(enqueue_flags_wait_kernel_simple),
    KERNEL(enqueue_flags_wait_kernel_event),
    KERNEL(enqueue_flags_wait_kernel_local),
    KERNEL(enqueue_flags_wait_kernel_event_local),
    KERNEL(enqueue_flags_wait_work_group_simple),
    KERNEL(enqueue_flags_wait_work_group_event),
    KERNEL(enqueue_flags_wait_work_group_local),
    KERNEL(enqueue_flags_wait_work_group_event_local)
};
static const size_t num_enqueue_block_flags = arr_size(sources_enqueue_block_flags);


int test_enqueue_flags(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
{
    cl_uint i;
    cl_int err_ret, res = 0;
    clCommandQueueWrapper dev_queue;
    cl_int kernel_results[MAX_GWS] = { -1 };
    int buff[MAX_GWS * BITS_DEPTH] = { 0 };

    size_t ret_len;
    size_t max_local_size = 1;
    cl_uint maxQueueSize = 0;
    err_ret = clGetDeviceInfo(device, CL_DEVICE_QUEUE_ON_DEVICE_MAX_SIZE, sizeof(maxQueueSize), &maxQueueSize, 0);
    test_error(err_ret, "clGetDeviceInfo(CL_DEVICE_QUEUE_ON_DEVICE_MAX_SIZE) failed");

    err_ret = clGetDeviceInfo(device, CL_DEVICE_MAX_WORK_GROUP_SIZE, sizeof(max_local_size), &max_local_size, &ret_len);
    test_error(err_ret, "clGetDeviceInfo(CL_DEVICE_MAX_WORK_GROUP_SIZE) failed");

    cl_queue_properties queue_prop_def[] =
    {
        CL_QUEUE_PROPERTIES, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE|CL_QUEUE_ON_DEVICE|CL_QUEUE_ON_DEVICE_DEFAULT,
        CL_QUEUE_SIZE, maxQueueSize,
        0
    };

    dev_queue = clCreateCommandQueueWithProperties(context, device, queue_prop_def, &err_ret);
    test_error(err_ret, "clCreateCommandQueueWithProperties(CL_QUEUE_DEVICE|CL_QUEUE_DEFAULT) failed");

    size_t global_size = MAX_GWS;
    size_t local_size = (max_local_size > global_size/16) ? global_size/16 : max_local_size;
    if(gWimpyMode)
    {
        global_size = 4;
        local_size = 2;
    }

    size_t failCnt = 0;
    for(i = 0; i < num_enqueue_block_flags; ++i)
    {
        if (!gKernelName.empty() && gKernelName != sources_enqueue_block_flags[i].kernel_name)
            continue;

        log_info("Running '%s' kernel (%d of %d) ...\n", sources_enqueue_block_flags[i].kernel_name, i + 1, num_enqueue_block_flags);

        clMemWrapper mem = clCreateBuffer(context, CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR, global_size * BITS_DEPTH * sizeof(cl_int), buff, &err_ret);
        test_error(err_ret, "clCreateBuffer() failed");

        kernel_arg args[] =
        {
            { sizeof(cl_mem),  &mem }
        };

        err_ret = run_n_kernel_args(context, queue, sources_enqueue_block_flags[i].lines, sources_enqueue_block_flags[i].num_lines, sources_enqueue_block_flags[i].kernel_name, local_size, global_size, kernel_results, sizeof(kernel_results), arr_size(args), args);
        if(check_error(err_ret, "'%s' kernel execution failed", sources_enqueue_block_flags[i].kernel_name)) { ++failCnt; res = -1; }
        else
        {
            int r = 0;
            for (int j=0; j<global_size; j++)
            {
                if (kernel_results[j] != 1 && j < (global_size / local_size) && check_error(-1, "'%s' kernel result[idx: %d] validation failed (test) %d != (expected) 1", sources_enqueue_block_flags[i].kernel_name, j, kernel_results[j]))
                {
                    r = -1;
                    break;
                }
                else if (kernel_results[j] != 0 && j >= (global_size / local_size) && check_error(-1, "'%s' kernel result[idx: %d] validation failed (test) %d != (expected) 0", sources_enqueue_block_flags[i].kernel_name, j, kernel_results[j]))
                {
                    r = -1;
                    break;
                }
            }
            if(r == 0) log_info("'%s' kernel is OK.\n", sources_enqueue_block_flags[i].kernel_name);
            else res = -1;
        }
    }

    if (failCnt > 0)
    {
        log_error("ERROR: %d of %d kernels failed.\n", failCnt, num_enqueue_block_flags);
    }

    return res;
}



#endif