//
// 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;
static const char* enqueue_simple_block[] =
{
    NL, "void block_fn(size_t tid, int mul, __global int* res)"
    NL, "{"
    NL, "  res[tid] = mul * 7 - 21;"
    NL, "}"
    NL, ""
    NL, "kernel void enqueue_simple_block(__global int* res)"
    NL, "{"
    NL, "  int multiplier = 3;"
    NL, "  size_t tid = get_global_id(0);"
    NL, ""
    NL, "  void (^kernelBlock)(void) = ^{ block_fn(tid, multiplier, res); };"
    NL, ""
    NL, "  res[tid] = -1;"
    NL, "  queue_t def_q = get_default_queue();"
    NL, "  ndrange_t ndrange = ndrange_1D(1);"
    NL, "  int enq_res = enqueue_kernel(def_q, CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange, kernelBlock);"
    NL, "  if(enq_res != CLK_SUCCESS) { res[tid] = -1; return; }"
    NL, "}"
    NL
};

static const char* enqueue_block_with_local_arg1[] =
{
    NL, "#define LOCAL_MEM_SIZE 10"
    NL, ""
    NL, "void block_fn_local_arg1(size_t tid, int mul, __global int* res, __local int* tmp)"
    NL, "{"
    NL, "  for(int i = 0; i < LOCAL_MEM_SIZE; i++)"
    NL, "  {"
    NL, "    tmp[i] = mul * 7 - 21;"
    NL, "    res[tid] += tmp[i];"
    NL, "  }"
    NL, "  res[tid] += 2;"
    NL, "}"
    NL, ""
    NL, "kernel void enqueue_block_with_local_arg1(__global int* res)"
    NL, "{"
    NL, "  int multiplier = 3;"
    NL, "  size_t tid = get_global_id(0);"
    NL, ""
    NL, "  void (^kernelBlock)(__local void*) = ^(__local void* buf){ block_fn_local_arg1(tid, multiplier, res, (local int*)buf); };"
    NL, ""
    NL, "  res[tid] = -2;"
    NL, "  queue_t def_q = get_default_queue();"
    NL, "  ndrange_t ndrange = ndrange_1D(1);"
    NL, "  int enq_res = enqueue_kernel(def_q, CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange, kernelBlock, (uint)(LOCAL_MEM_SIZE*sizeof(int)));"
    NL, "  if(enq_res != CLK_SUCCESS) { res[tid] = -1; return; }"
    NL, "}"
    NL
};

static const char* enqueue_block_with_local_arg2[] =
{
    NL, "#define LOCAL_MEM_SIZE 10"
    NL, ""
    NL, "void block_fn_local_arg1(size_t tid, int mul, __global int* res, __local int* tmp1, __local float4* tmp2)"
    NL, "{"
    NL, "  for(int i = 0; i < LOCAL_MEM_SIZE; i++)"
    NL, "  {"
    NL, "    tmp1[i]   = mul * 7 - 21;"
    NL, "    tmp2[i].x = (float)(mul * 7 - 21);"
    NL, "    tmp2[i].y = (float)(mul * 7 - 21);"
    NL, "    tmp2[i].z = (float)(mul * 7 - 21);"
    NL, "    tmp2[i].w = (float)(mul * 7 - 21);"
    NL, ""
    NL, "    res[tid] += tmp1[i];"
    NL, "    res[tid] += (int)(tmp2[i].x+tmp2[i].y+tmp2[i].z+tmp2[i].w);"
    NL, "  }"
    NL, "  res[tid] += 2;"
    NL, "}"
    NL, ""
    NL, "kernel void enqueue_block_with_local_arg2(__global int* res)"
    NL, "{"
    NL, "  int multiplier = 3;"
    NL, "  size_t tid = get_global_id(0);"
    NL, ""
    NL, "  void (^kernelBlock)(__local void*, __local void*) = ^(__local void* buf1, __local void* buf2)"
    NL, "    { block_fn_local_arg1(tid, multiplier, res, (local int*)buf1, (local float4*)buf2); };"
    NL, ""
    NL, "  res[tid] = -2;"
    NL, "  queue_t def_q = get_default_queue();"
    NL, "  ndrange_t ndrange = ndrange_1D(1);"
    NL, "  int enq_res = enqueue_kernel(def_q, CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange, kernelBlock, (uint)(LOCAL_MEM_SIZE*sizeof(int)), (uint)(LOCAL_MEM_SIZE*sizeof(float4)));"
    NL, "  if(enq_res != CLK_SUCCESS) { res[tid] = -1; return; }"
    NL, "}"
    NL
};

static const char* enqueue_block_with_wait_list[] =
{
    NL, "#define BLOCK_SUBMITTED 1"
    NL, "#define BLOCK_COMPLETED 2"
    NL, "#define CHECK_SUCCESS   0"
    NL, ""
    NL, "kernel void enqueue_block_with_wait_list(__global int* res)"
    NL, "{"
    NL, "  size_t tid = get_global_id(0);"
    NL, ""
    NL, "  clk_event_t user_evt = create_user_event();"
    NL, ""
    NL, "  res[tid] = BLOCK_SUBMITTED;"
    NL, "  queue_t def_q = get_default_queue();"
    NL, "  ndrange_t ndrange = ndrange_1D(1);"
    NL, "  clk_event_t block_evt;"
    NL, "  int enq_res = enqueue_kernel(def_q, CLK_ENQUEUE_FLAGS_NO_WAIT, ndrange, 1, &user_evt, &block_evt,"
    NL, "  ^{"
    NL, "      res[tid] = BLOCK_COMPLETED;"
    NL, "   });"
    NL, "  if(enq_res != CLK_SUCCESS) { res[tid] = -1; return; }"
    NL, ""
    NL, "  retain_event(block_evt);"
    NL, "  release_event(block_evt);"
    NL, ""
    NL, "  //check block is not started"
    NL, "  if(res[tid] == BLOCK_SUBMITTED)"
    NL, "  {"
    NL, "    clk_event_t my_evt;"
    NL, "    enqueue_kernel(def_q, CLK_ENQUEUE_FLAGS_NO_WAIT, ndrange, 1, &block_evt, &my_evt, "
    NL, "    ^{"
    NL, "       //check block is completed"
    NL, "       if(res[tid] == BLOCK_COMPLETED) res[tid] = CHECK_SUCCESS;"
    NL, "     });"
    NL, "    release_event(my_evt);"
    NL, "  }"
    NL, ""
    NL, "  set_user_event_status(user_evt, CL_COMPLETE);"
    NL, ""
    NL, "  release_event(user_evt);"
    NL, "  release_event(block_evt);"
    NL, "}"
    NL
};

static const char* enqueue_block_with_wait_list_and_local_arg[] =
{
    NL, "#define LOCAL_MEM_SIZE 10"
    NL, "#define BLOCK_COMPLETED 1"
    NL, "#define BLOCK_SUBMITTED 2"
    NL, "#define BLOCK_STARTED   3"
    NL, "#define CHECK_SUCCESS   0"
    NL, ""
    NL, "void block_fn_local_arg(size_t tid, int mul, __global int* res, __local int* tmp)"
    NL, "{"
    NL, "  res[tid] = BLOCK_STARTED;"
    NL, "  for(int i = 0; i < LOCAL_MEM_SIZE; i++)"
    NL, "  {"
    NL, "    tmp[i] = mul * 7 - 21;"
    NL, "    res[tid] += tmp[i];"
    NL, "  }"
    NL, "  if(res[tid] == BLOCK_STARTED) res[tid] = BLOCK_COMPLETED;"
    NL, "}"
    NL, ""
    NL, "kernel void enqueue_block_with_wait_list_and_local_arg(__global int* res)"
    NL, "{"
    NL, "  int multiplier = 3;"
    NL, "  size_t tid = get_global_id(0);"
    NL, "  clk_event_t user_evt = create_user_event();"
    NL, ""
    NL, "  res[tid] = BLOCK_SUBMITTED;"
    NL, "  queue_t def_q = get_default_queue();"
    NL, "  ndrange_t ndrange = ndrange_1D(1);"
    NL, "  clk_event_t block_evt;"
    NL, "  int enq_res = enqueue_kernel(def_q, CLK_ENQUEUE_FLAGS_NO_WAIT, ndrange, 1, &user_evt, &block_evt, "
    NL, "    ^(__local void* buf) {"
    NL, "       block_fn_local_arg(tid, multiplier, res, (__local int*)buf);"
    NL, "     }, LOCAL_MEM_SIZE*sizeof(int));"
    NL, "  if(enq_res != CLK_SUCCESS) { res[tid] = -1; return; }"
    NL, ""
    NL, "  retain_event(block_evt);"
    NL, "  release_event(block_evt);"
    NL, ""
    NL, "  //check block is not started"
    NL, "  if(res[tid] == BLOCK_SUBMITTED)"
    NL, "  {"
    NL, "    clk_event_t my_evt;"
    NL, "    enqueue_kernel(def_q, CLK_ENQUEUE_FLAGS_NO_WAIT, ndrange, 1, &block_evt, &my_evt, "
    NL, "    ^{"
    NL, "       //check block is completed"
    NL, "       if(res[tid] == BLOCK_COMPLETED) res[tid] = CHECK_SUCCESS;"
    NL, "     });"
    NL, "    release_event(my_evt);"
    NL, "  }"
    NL, ""
    NL, "  set_user_event_status(user_evt, CL_COMPLETE);"
    NL, ""
    NL, "  release_event(user_evt);"
    NL, "  release_event(block_evt);"
    NL, "}"
    NL
};

static const char* enqueue_block_get_kernel_work_group_size[] =
{
    NL, "void block_fn(size_t tid, int mul, __global int* res)"
    NL, "{"
    NL, "  res[tid] = mul * 7 - 21;"
    NL, "}"
    NL, ""
    NL, "kernel void enqueue_block_get_kernel_work_group_size(__global int* res)"
    NL, "{"
    NL, "    int multiplier = 3;"
    NL, "    size_t tid = get_global_id(0);"
    NL, ""
    NL, "    void (^kernelBlock)(void) = ^{ block_fn(tid, multiplier, res); };"
    NL, ""
    NL, "    size_t local_work_size = get_kernel_work_group_size(kernelBlock);"
    NL, "    if (local_work_size <= 0){ res[tid] = -1; return; }"
    NL, "    size_t global_work_size = local_work_size * 4;"
    NL, ""
    NL, "    res[tid] = -1;"
    NL, "    queue_t q1 = get_default_queue();"
    NL, "    ndrange_t ndrange = ndrange_1D(global_work_size, local_work_size);"
    NL, ""
    NL, "    int enq_res = enqueue_kernel(q1, CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange, kernelBlock);"
    NL, "    if(enq_res != CLK_SUCCESS) { res[tid] = -1; return; }"
    NL, "}"
};

static const char* enqueue_block_get_kernel_preferred_work_group_size_multiple[] =
{
    NL, "void block_fn(size_t tid, int mul, __global int* res)"
    NL, "{"
    NL, "  res[tid] = mul * 7 - 21;"
    NL, "}"
    NL, ""
    NL, "kernel void enqueue_block_get_kernel_preferred_work_group_size_multiple(__global int* res)"
    NL, "{"
    NL, "    int multiplier = 3;"
    NL, "    size_t tid = get_global_id(0);"
    NL, ""
    NL, "    void (^kernelBlock)(void) = ^{ block_fn(tid, multiplier, res); };"
    NL, ""
    NL, "    size_t local_work_size = get_kernel_preferred_work_group_size_multiple(kernelBlock);"
    NL, "    if (local_work_size <= 0){ res[tid] = -1; return; }"
    NL, "    size_t global_work_size = local_work_size * 4;"
    NL, ""
    NL, "    res[tid] = -1;"
    NL, "    queue_t q1 = get_default_queue();"
    NL, "    ndrange_t ndrange = ndrange_1D(global_work_size, local_work_size);"
    NL, ""
    NL, "    int enq_res = enqueue_kernel(q1, CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange, kernelBlock);"
    NL, "    if(enq_res != CLK_SUCCESS) { res[tid] = -1; return; }"
    NL, "}"
};

static const char* enqueue_block_capture_event_profiling_info_after_execution[] =
{
    NL, "#define MAX_GWS " STRINGIFY_VALUE(MAX_GWS)
    NL, ""
    NL, "__global ulong value[MAX_GWS*2] = {0};"
    NL, ""
    NL, "void block_fn(size_t tid, __global int* res)"
    NL, "{"
    NL, "    res[tid] = -2;"
    NL, "}"
    NL, ""
    NL, "void check_res(size_t tid, const clk_event_t evt, __global int* res)"
    NL, "{"
    NL, "    capture_event_profiling_info (evt, CLK_PROFILING_COMMAND_EXEC_TIME, &value[tid*2]);"
    NL, ""
    NL, "    if (value[tid*2] > 0 && value[tid*2+1] > 0) res[tid] =  0;"
    NL, "    else                                        res[tid] = -4;"
    NL, "    release_event(evt);"
    NL, "}"
    NL, ""
    NL, "kernel void enqueue_block_capture_event_profiling_info_after_execution(__global int* res)"
    NL, "{"
    NL, "    size_t tid = get_global_id(0);"
    NL, ""
    NL, "    res[tid] = -1;"
    NL, "    queue_t def_q = get_default_queue();"
    NL, "    ndrange_t ndrange = ndrange_1D(1);"
    NL, "    clk_event_t block_evt1;"
    NL, ""
    NL, "    void (^kernelBlock)(void)  = ^{ block_fn (tid, res);                   };"
    NL, ""
    NL, "    int enq_res = enqueue_kernel(def_q, CLK_ENQUEUE_FLAGS_NO_WAIT, ndrange, 0, NULL, &block_evt1, kernelBlock);"
    NL, "    if(enq_res != CLK_SUCCESS) { res[tid] = -1; return; }"
    NL, ""
    NL, "    void (^checkBlock) (void)  = ^{ check_res(tid, block_evt1, res);      };"
    NL, ""
    NL, "    enq_res = enqueue_kernel(def_q, CLK_ENQUEUE_FLAGS_NO_WAIT, ndrange, 1, &block_evt1, NULL, checkBlock);"
    NL, "    if(enq_res != CLK_SUCCESS) { res[tid] = -3; return; }"
    NL, "}"
    NL
};

static const char* enqueue_block_capture_event_profiling_info_before_execution[] =
{
    NL, "#define MAX_GWS " STRINGIFY_VALUE(MAX_GWS)
    NL, ""
    NL, "__global ulong value[MAX_GWS*2] = {0};"
    NL, ""
    NL, "void block_fn(size_t tid, __global int* res)"
    NL, "{"
    NL, "    res[tid] = -2;"
    NL, "}"
    NL, ""
    NL, "void check_res(size_t tid, const ulong *value, __global int* res)"
    NL, "{"
    NL, "    if (value[tid*2] > 0 && value[tid*2+1] > 0) res[tid] =  0;"
    NL, "    else                                        res[tid] = -4;"
    NL, "}"
    NL, ""
    NL, "kernel void enqueue_block_capture_event_profiling_info_before_execution(__global int* res)"
    NL, "{"
    NL, "    int multiplier = 3;"
    NL, "    size_t tid = get_global_id(0);"
    NL, "    clk_event_t user_evt = create_user_event();"
    NL, ""
    NL, "    res[tid] = -1;"
    NL, "    queue_t def_q = get_default_queue();"
    NL, "    ndrange_t ndrange = ndrange_1D(1);"
    NL, "    clk_event_t block_evt1;"
    NL, "    clk_event_t block_evt2;"
    NL, ""
    NL, "    void (^kernelBlock)(void)  = ^{ block_fn (tid, res);                   };"
    NL, ""
    NL, "    int enq_res = enqueue_kernel(def_q, CLK_ENQUEUE_FLAGS_NO_WAIT, ndrange, 1, &user_evt, &block_evt1, kernelBlock);"
    NL, "    if(enq_res != CLK_SUCCESS) { res[tid] = -1; return; }"
    NL, ""
    NL, "    capture_event_profiling_info (block_evt1, CLK_PROFILING_COMMAND_EXEC_TIME, &value[tid*2]);"
    NL, ""
    NL, "    set_user_event_status(user_evt, CL_COMPLETE);"
    NL, ""
    NL, "    void (^checkBlock) (void)  = ^{ check_res(tid, &value, res);      };"
    NL, ""
    NL, "    enq_res = enqueue_kernel(def_q, CLK_ENQUEUE_FLAGS_NO_WAIT, ndrange, 1, &block_evt1, &block_evt2, checkBlock);"
    NL, "    if(enq_res != CLK_SUCCESS) { res[tid] = -3; return; }"
    NL, ""
    NL, "    release_event(user_evt);"
    NL, "    release_event(block_evt1);"
    NL, "    release_event(block_evt2);"
    NL, "}"
    NL
};

static const char* enqueue_block_with_barrier[] =
{
    NL, "void block_fn(size_t tid, int mul, __global int* res)"
    NL, "{"
    NL, "  if(mul > 0) barrier(CLK_GLOBAL_MEM_FENCE);"
    NL, "  res[tid] = mul * 7 -21;"
    NL, "}"
    NL, ""
    NL, "void loop_fn(size_t tid, int n, __global int* res)"
    NL, "{"
    NL, "  while(n > 0)"
    NL, "  {"
    NL, "    barrier(CLK_GLOBAL_MEM_FENCE);"
    NL, "    res[tid] = 0;"
    NL, "    --n;"
    NL, "  }"
    NL, "}"
    NL, ""
    NL, "kernel void enqueue_block_with_barrier(__global int* res)"
    NL, "{"
    NL, "  int multiplier = 3;"
    NL, "  size_t tid = get_global_id(0);"
    NL, "  queue_t def_q = get_default_queue();"
    NL, "  res[tid] = -1;"
    NL, "  size_t n = 256;"
    NL, ""
    NL, "  void (^kernelBlock)(void) = ^{ block_fn(tid, multiplier, res); };"
    NL, ""
    NL, "  ndrange_t ndrange = ndrange_1D(n);"
    NL, "  int enq_res = enqueue_kernel(def_q, CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange, kernelBlock);"
    NL, "  if(enq_res != CLK_SUCCESS) { res[tid] = -1; return; }"
    NL, ""
    NL, "  void (^loopBlock)(void) = ^{ loop_fn(tid, n, res); };"
    NL, ""
    NL, "  enq_res = enqueue_kernel(def_q, CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange, loopBlock);"
    NL, "  if(enq_res != CLK_SUCCESS) { res[tid] = -1; return; }"
    NL, "}"
    NL
};

static const char* enqueue_marker_with_block_event[] =
{
    NL, "#define BLOCK_COMPLETED 1"
    NL, "#define BLOCK_SUBMITTED 2"
    NL, "#define CHECK_SUCCESS   0"
    NL, ""
    NL, "kernel void enqueue_marker_with_block_event(__global int* res)"
    NL, "{"
    NL, "  size_t tid = get_global_id(0);"
    NL, ""
    NL, "  clk_event_t user_evt = create_user_event();"
    NL, ""
    NL, "  res[tid] = BLOCK_SUBMITTED;"
    NL, "  queue_t def_q = get_default_queue();"
    NL, "  ndrange_t ndrange = ndrange_1D(1);"
    NL, ""
    NL, "  clk_event_t block_evt1;"
    NL, "  clk_event_t marker_evt;"
    NL, ""
    NL, "  int enq_res = enqueue_kernel(def_q, CLK_ENQUEUE_FLAGS_NO_WAIT, ndrange, 1, &user_evt, &block_evt1,"
    NL, "  ^{"
    NL, "     res[tid] = BLOCK_COMPLETED;"
    NL, "   });"
    NL, "  if(enq_res != CLK_SUCCESS) { res[tid] = -2; return; }"
    NL, ""
    NL, "  enq_res = enqueue_marker(def_q, 1, &block_evt1, &marker_evt);"
    NL, "  if(enq_res != CLK_SUCCESS) { res[tid] = -3; return; }"
    NL, ""
    NL, "  retain_event(marker_evt);"
    NL, "  release_event(marker_evt);"
    NL, ""
    NL, "  //check block is not started"
    NL, "  if(res[tid] == BLOCK_SUBMITTED)"
    NL, "  {"
    NL, "    clk_event_t my_evt;"
    NL, "    enqueue_kernel(def_q, CLK_ENQUEUE_FLAGS_NO_WAIT, ndrange, 1, &marker_evt, &my_evt, "
    NL, "    ^{"
    NL, "       //check block is completed"
    NL, "       if(res[tid] == BLOCK_COMPLETED) res[tid] = CHECK_SUCCESS;"
    NL, "     });"
    NL, "    release_event(my_evt);"
    NL, "  }"
    NL, ""
    NL, "  set_user_event_status(user_evt, CL_COMPLETE);"
    NL, ""
    NL, "  release_event(block_evt1);"
    NL, "  release_event(marker_evt);"
    NL, "  release_event(user_evt);"
    NL, "}"
    NL
};

static const char* enqueue_marker_with_user_event[] =
{
    NL, "#define BLOCK_COMPLETED 1"
    NL, "#define BLOCK_SUBMITTED 2"
    NL, "#define CHECK_SUCCESS   0"
    NL, ""
    NL, "kernel void enqueue_marker_with_user_event(__global int* res)"
    NL, "{"
    NL, "  size_t tid = get_global_id(0);"
    NL, "  uint multiplier = 7;"
    NL, ""
    NL, "  clk_event_t user_evt = create_user_event();"
    NL, ""
    NL, "  res[tid] = BLOCK_SUBMITTED;"
    NL, "  queue_t def_q = get_default_queue();"
    NL, "  ndrange_t ndrange = ndrange_1D(1);"
    NL, ""
    NL, "  clk_event_t marker_evt;"
    NL, "  clk_event_t block_evt;"
    NL, ""
    NL, "  int enq_res = enqueue_marker(def_q, 1, &user_evt, &marker_evt);"
    NL, "  if(enq_res != CLK_SUCCESS) { res[tid] = -1; return; }"
    NL, ""
    NL, "  retain_event(marker_evt);"
    NL, "  release_event(marker_evt);"
    NL, ""
    NL, "  enqueue_kernel(def_q, CLK_ENQUEUE_FLAGS_NO_WAIT, ndrange, 1, &marker_evt, &block_evt, "
    NL, "  ^{"
    NL, "     if(res[tid] == BLOCK_SUBMITTED) res[tid] = CHECK_SUCCESS;"
    NL, "   });"
    NL, ""
    NL, "  //check block is not started"
    NL, "  if(res[tid] != BLOCK_SUBMITTED)  { res[tid] = -2; return; }"
    NL, ""
    NL, "  set_user_event_status(user_evt, CL_COMPLETE);"
    NL, ""
    NL, "  release_event(block_evt);"
    NL, "  release_event(marker_evt);"
    NL, "  release_event(user_evt);"
    NL, "}"
    NL
};

static const char* enqueue_marker_with_mixed_events[] =
{
    NL, "#define BLOCK_COMPLETED 1"
    NL, "#define BLOCK_SUBMITTED 2"
    NL, "#define CHECK_SUCCESS   0"
    NL, ""
    NL, "kernel void enqueue_marker_with_mixed_events(__global int* res)"
    NL, "{"
    NL, "  size_t tid = get_global_id(0);"
    NL, ""
    NL, "  clk_event_t mix_ev[2];"
    NL, "  mix_ev[0] = create_user_event();"
    NL, ""
    NL, "  res[tid] = BLOCK_SUBMITTED;"
    NL, "  queue_t def_q = get_default_queue();"
    NL, "  ndrange_t ndrange = ndrange_1D(1);"
    NL, ""
    NL, "  int enq_res = enqueue_kernel(def_q, CLK_ENQUEUE_FLAGS_NO_WAIT, ndrange, 1, &mix_ev[0], &mix_ev[1],"
    NL, "  ^{"
    NL, "     res[tid] = BLOCK_COMPLETED;"
    NL, "   });"
    NL, "  if(enq_res != CLK_SUCCESS) { res[tid] = -2; return; }"
    NL, ""
    NL, "  clk_event_t marker_evt;"
    NL, ""
    NL, "  enq_res = enqueue_marker(def_q, 2, mix_ev, &marker_evt);"
    NL, "  if(enq_res != CLK_SUCCESS) { res[tid] = -3; return; }"
    NL, ""
    NL, "  retain_event(marker_evt);"
    NL, "  release_event(marker_evt);"
    NL, ""
    NL, "  //check block is not started"
    NL, "  if(res[tid] == BLOCK_SUBMITTED)"
    NL, "  {"
    NL, "    clk_event_t my_evt;"
    NL, "    enqueue_kernel(def_q, CLK_ENQUEUE_FLAGS_NO_WAIT, ndrange, 1, &marker_evt, &my_evt, "
    NL, "    ^{"
    NL, "       //check block is completed"
    NL, "       if(res[tid] == BLOCK_COMPLETED) res[tid] = CHECK_SUCCESS;"
    NL, "     });"
    NL, "    release_event(my_evt);"
    NL, "  }"
    NL, ""
    NL, "  set_user_event_status(mix_ev[0], CL_COMPLETE);"
    NL, ""
    NL, "  release_event(mix_ev[1]);"
    NL, "  release_event(marker_evt);"
    NL, "  release_event(mix_ev[0]);"
    NL, "}"
    NL
};

static const char* enqueue_block_with_mixed_events[] =
{
    NL, "kernel void enqueue_block_with_mixed_events(__global int* res)"
    NL, "{"
    NL, "  int enq_res;"
    NL, "  size_t tid = get_global_id(0);"
    NL, "  clk_event_t mix_ev[3];"
    NL, "  mix_ev[0] = create_user_event();"
    NL, "  queue_t def_q = get_default_queue();"
    NL, "  ndrange_t ndrange = ndrange_1D(1);"
    NL, "  res[tid] = -2;"
    NL, ""
    NL, "  enq_res = enqueue_kernel(def_q, CLK_ENQUEUE_FLAGS_NO_WAIT, ndrange, 1, &mix_ev[0], &mix_ev[1], ^{ res[tid]++; });"
    NL, "  if(enq_res != CLK_SUCCESS) { res[tid] = -1; return; }"
    NL, ""
    NL, "  enq_res = enqueue_marker(def_q, 1, &mix_ev[1], &mix_ev[2]);"
    NL, "  if(enq_res != CLK_SUCCESS) { res[tid] = -3; return; }"
    NL, ""
    NL, "  enq_res = enqueue_kernel(def_q, CLK_ENQUEUE_FLAGS_NO_WAIT, ndrange, sizeof(mix_ev)/sizeof(mix_ev[0]), mix_ev, NULL, ^{ res[tid]++; });"
    NL, "  if(enq_res != CLK_SUCCESS) { res[tid] = -4; return; }"
    NL, ""
    NL, "  set_user_event_status(mix_ev[0], CL_COMPLETE);"
    NL, ""
    NL, "  release_event(mix_ev[0]);"
    NL, "  release_event(mix_ev[1]);"
    NL, "  release_event(mix_ev[2]);"
    NL, "}"
    NL
};

static const kernel_src sources_enqueue_block[] =
{
    KERNEL(enqueue_simple_block),
    // Block with local mem
    KERNEL(enqueue_block_with_local_arg1),
    KERNEL(enqueue_block_with_local_arg2),
    KERNEL(enqueue_block_with_wait_list),
    KERNEL(enqueue_block_with_wait_list_and_local_arg),
    // WG size built-ins
    KERNEL(enqueue_block_get_kernel_work_group_size),
    KERNEL(enqueue_block_get_kernel_preferred_work_group_size_multiple),
    // Event profiling info
    KERNEL(enqueue_block_capture_event_profiling_info_after_execution),
    KERNEL(enqueue_block_capture_event_profiling_info_before_execution),
    // Marker
    KERNEL(enqueue_marker_with_block_event),
    KERNEL(enqueue_marker_with_user_event),
    // Mixed events
    KERNEL(enqueue_marker_with_mixed_events),
    KERNEL(enqueue_block_with_mixed_events),
    // Barrier
    KERNEL(enqueue_block_with_barrier),

};
static const size_t num_kernels_enqueue_block = arr_size(sources_enqueue_block);

static int check_kernel_results(cl_int* results, cl_int len)
{
    for(cl_int i = 0; i < len; ++i)
    {
        if(results[i] != 0) return i;
    }
    return -1;
}

int test_enqueue_block(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
{
    cl_uint i;
    cl_int n, err_ret, res = 0;
    clCommandQueueWrapper dev_queue;
    cl_int kernel_results[MAX_GWS] = {0};

    size_t ret_len;
    cl_uint max_queues = 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_ON_DEVICE_QUEUES, sizeof(max_queues), &max_queues, &ret_len);
    test_error(err_ret, "clGetDeviceInfo(CL_DEVICE_MAX_ON_DEVICE_QUEUES) failed");

    size_t max_local_size = 1;
    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_PROFILING_ENABLE,
        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_kernels_enqueue_block; ++i)
    {
        if (!gKernelName.empty() && gKernelName != sources_enqueue_block[i].kernel_name)
            continue;

        log_info("Running '%s' kernel (%d of %d) ...\n", sources_enqueue_block[i].kernel_name, i + 1, num_kernels_enqueue_block);
        err_ret = run_n_kernel_args(context, queue, sources_enqueue_block[i].lines, sources_enqueue_block[i].num_lines, sources_enqueue_block[i].kernel_name, local_size, global_size, kernel_results, sizeof(kernel_results), 0, NULL);
        if(check_error(err_ret, "'%s' kernel execution failed", sources_enqueue_block[i].kernel_name)) { ++failCnt; res = -1; }
        else if((n = check_kernel_results(kernel_results, arr_size(kernel_results))) >= 0 && check_error(-1, "'%s' kernel results validation failed: [%d] returned %d expected 0", sources_enqueue_block[i].kernel_name, n, kernel_results[n])) res = -1;
        else log_info("'%s' kernel is OK.\n", sources_enqueue_block[i].kernel_name);
    }

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

    return res;
}


#endif