| // |
| // 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 int nestingLevel = 3; |
| |
| static const char* enqueue_1D_wg_size_single[] = |
| { |
| NL, "void block_fn(int level, int maxGlobalWorkSize, __global int* rnd, __global int* res)" |
| NL, "{" |
| NL, " size_t tidX = get_global_id(0);" |
| NL, " queue_t def_q = get_default_queue();" |
| NL, " if(--level < 0) return;" |
| NL, "" |
| NL, " void (^kernelBlock)(void) = ^{ block_fn(level, maxGlobalWorkSize, rnd, res); };" |
| NL, " uint wg = get_kernel_work_group_size(kernelBlock);" |
| NL, "" |
| NL, " const size_t gs = 64 * 64 * 64;" |
| NL, " size_t ls = rnd[tidX % maxGlobalWorkSize] % wg % gs;" |
| NL, " ls = ls? ls: 1;" |
| NL, "" |
| NL, " ndrange_t ndrange = ndrange_1D(gs, ls);" |
| NL, "" |
| NL, " // Only 1 work-item enqueues block" |
| NL, " if(tidX == 0)" |
| NL, " {" |
| NL, " atomic_inc(&res[tidX % maxGlobalWorkSize]);" |
| NL, " int enq_res = enqueue_kernel(def_q, CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange, kernelBlock);" |
| NL, " if(enq_res != CLK_SUCCESS) { res[tidX % maxGlobalWorkSize] = -1; return; }" |
| NL, " }" |
| NL, "}" |
| NL, "" |
| NL, "kernel void enqueue_1D_wg_size_single(__global int* res, int level, int maxGlobalWorkSize, __global int* rnd)" |
| NL, "{" |
| NL, " block_fn(level, maxGlobalWorkSize, rnd, res);" |
| NL, "}" |
| NL |
| }; |
| |
| static int check_single(cl_int* results, cl_int len, cl_int nesting_level) |
| { |
| for(size_t i = 0; i < len; ++i) |
| { |
| if(i == 0 && results[i] != nestingLevel) |
| { |
| log_error("ERROR: Kernel returned %d vs. expected %d, index: %d\n", results[i], nestingLevel, i); |
| return (int)i; |
| } |
| |
| if(i > 0 && results[i] != 0) |
| { |
| log_error("ERROR: Kernel returned %d vs. expected 0, index: %d\n", results[i], i); |
| return (int)i; |
| } |
| } |
| |
| return -1; |
| } |
| |
| static const char* enqueue_1D_wg_size_some_eq[] = |
| { |
| NL, "void block_fn(__global int* res, int level, int maxGlobalWorkSize, __global int* rnd)" |
| NL, "{" |
| NL, " size_t tidX = get_global_id(0);" |
| NL, " queue_t def_q = get_default_queue();" |
| NL, " if(--level < 0) return;" |
| NL, "" |
| NL, " void (^kernelBlock)(void) = ^{ block_fn(res, level, maxGlobalWorkSize, rnd); };" |
| NL, " uint wg = get_kernel_work_group_size(kernelBlock);" |
| NL, "" |
| NL, " const size_t gs = 8 * 8 * 2;" |
| NL, " size_t ls = rnd[tidX % maxGlobalWorkSize] % wg % gs;" |
| NL, " ls = ls? ls: 1;" |
| NL, "" |
| NL, " ndrange_t ndrange = ndrange_1D(gs, ls);" |
| NL, "" |
| NL, " // Some work-items enqueues nested blocks with the same level" |
| NL, " if((tidX % (maxGlobalWorkSize / 8)) == 0)" |
| NL, " {" |
| NL, " atomic_inc(&res[tidX % maxGlobalWorkSize]);" |
| NL, " int enq_res = enqueue_kernel(def_q, CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange, kernelBlock);" |
| NL, " if(enq_res != CLK_SUCCESS) { res[tidX % maxGlobalWorkSize] = -1; return; }" |
| NL, " }" |
| NL, "}" |
| NL, "" |
| NL, "kernel void enqueue_1D_wg_size_some_eq(__global int* res, int level, int maxGlobalWorkSize, __global int* rnd)" |
| NL, "{" |
| NL, " block_fn(res, level, maxGlobalWorkSize, rnd);" |
| NL, "}" |
| NL |
| }; |
| |
| void generate_reference_results_some_eq_1D(std::vector<cl_int> &referenceResults, cl_int maxGlobalWorkSize, cl_int level) |
| { |
| size_t globalSize = (level == nestingLevel) ? maxGlobalWorkSize: (8 * 8 * 2); |
| if(--level < 0) |
| { |
| return; |
| } |
| |
| for (size_t tidX = 0; tidX < globalSize; ++tidX) |
| { |
| if ((tidX % (maxGlobalWorkSize / 8)) == 0) |
| { |
| ++referenceResults[tidX % maxGlobalWorkSize]; |
| generate_reference_results_some_eq_1D(referenceResults, maxGlobalWorkSize, level); |
| } |
| } |
| } |
| |
| static int check_some_eq_1D(cl_int* results, cl_int len, cl_int nesting_level) |
| { |
| std::vector<cl_int> referenceResults(len, 0); |
| generate_reference_results_some_eq_1D(referenceResults, len, nesting_level); |
| |
| for(size_t i = 0; i < len; ++i) |
| { |
| if (results[i] != referenceResults[i]) |
| { |
| log_error("ERROR: Kernel returned %d vs. expected %d, index: %d\n", results[i], referenceResults[i], i); |
| return (int)i; |
| } |
| } |
| |
| return -1; |
| } |
| |
| static const char* enqueue_1D_wg_size_some_diff[] = |
| { |
| NL, "void block_fn(__global int* res, int level, int maxGlobalWorkSize, __global int* rnd)" |
| NL, "{" |
| NL, " size_t tidX = get_global_id(0);" |
| NL, " queue_t def_q = get_default_queue();" |
| NL, " if(--level < 0) return;" |
| NL, "" |
| NL, " void (^kernelBlock)(void) = ^{ block_fn(res, level, maxGlobalWorkSize, rnd); };" |
| NL, " uint wg = get_kernel_work_group_size(kernelBlock);" |
| NL, "" |
| NL, " const size_t gs = 8 * 8 * 8;" |
| NL, " size_t ls = rnd[tidX % maxGlobalWorkSize] % wg % gs;" |
| NL, " ls = ls? ls: 1;" |
| NL, "" |
| NL, " ndrange_t ndrange = ndrange_1D(gs, ls);" |
| NL, "" |
| NL, " // Some work-items enqueues nested blocks with different levels" |
| NL, " if((tidX % 2) == 0)" |
| NL, " {" |
| NL, " atomic_inc(&res[tidX % maxGlobalWorkSize]);" |
| NL, " if(level >= tidX)" |
| NL, " {" |
| NL, " int enq_res = enqueue_kernel(def_q, CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange, kernelBlock);" |
| NL, " if(enq_res != CLK_SUCCESS) { res[tidX % maxGlobalWorkSize] = -1; return; }" |
| NL, " }" |
| NL, " }" |
| NL, "}" |
| NL, "" |
| NL, "kernel void enqueue_1D_wg_size_some_diff(__global int* res, int level, int maxGlobalWorkSize, __global int* rnd)" |
| NL, "{" |
| NL, " block_fn(res, level, maxGlobalWorkSize, rnd);" |
| NL, "}" |
| NL |
| }; |
| |
| void generate_reference_results_some_diff_1D(std::vector<cl_int> &referenceResults, cl_int maxGlobalWorkSize, cl_int level) |
| { |
| size_t globalSize = (level == nestingLevel) ? maxGlobalWorkSize: (8 * 8 * 8); |
| if(--level < 0) |
| { |
| return; |
| } |
| |
| for (size_t tidX = 0; tidX < globalSize; ++tidX) |
| { |
| if ((tidX % 2) == 0) |
| { |
| ++referenceResults[tidX % maxGlobalWorkSize]; |
| if (level >= tidX) |
| { |
| generate_reference_results_some_diff_1D(referenceResults, maxGlobalWorkSize, level); |
| } |
| } |
| } |
| } |
| |
| static int check_some_diff_1D(cl_int* results, cl_int maxGlobalWorkSize, cl_int nesting_level) |
| { |
| std::vector<cl_int> referenceResults(maxGlobalWorkSize, 0); |
| generate_reference_results_some_diff_1D(referenceResults, maxGlobalWorkSize, nesting_level); |
| |
| for(size_t i = 0; i < maxGlobalWorkSize; ++i) |
| { |
| if (results[i] != referenceResults[i]) |
| { |
| log_error("ERROR: Kernel returned %d vs. expected %d, index: %d\n", results[i], referenceResults[i], i); |
| return (int)i; |
| } |
| } |
| |
| return -1; |
| } |
| |
| static const char* enqueue_1D_wg_size_all_eq[] = |
| { |
| NL, "void block_fn(int level, int maxGlobalWorkSize, __global int* rnd, __global int* res)" |
| NL, "{" |
| NL, " size_t tidX = get_global_id(0);" |
| NL, " queue_t def_q = get_default_queue();" |
| NL, " if(--level < 0) return;" |
| NL, "" |
| NL, " void (^kernelBlock)(void) = ^{ block_fn(level, maxGlobalWorkSize, rnd, res); };" |
| NL, " uint wg = get_kernel_work_group_size(kernelBlock);" |
| NL, "" |
| NL, " const size_t gs = 8;" |
| NL, " size_t ls = rnd[tidX % maxGlobalWorkSize] % wg % gs;" |
| NL, " ls = ls? ls: 1;" |
| NL, "" |
| NL, " ndrange_t ndrange = ndrange_1D(gs, ls);" |
| NL, "" |
| NL, " // All work-items enqueues nested blocks with the same level" |
| NL, " atomic_inc(&res[tidX % maxGlobalWorkSize]);" |
| NL, " int enq_res = enqueue_kernel(def_q, CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange, kernelBlock);" |
| NL, " if(enq_res != CLK_SUCCESS) { res[tidX % maxGlobalWorkSize] = -1; return; }" |
| NL, "}" |
| NL, "" |
| NL, "kernel void enqueue_1D_wg_size_all_eq(__global int* res, int level, int maxGlobalWorkSize, __global int* rnd)" |
| NL, "{" |
| NL, " block_fn(level, maxGlobalWorkSize, rnd, res);" |
| NL, "}" |
| NL |
| }; |
| |
| void generate_reference_results_all_eq_1D(std::vector<cl_int> &referenceResults, cl_int len, cl_int level) |
| { |
| size_t globalSize = (level == nestingLevel) ? len: 8; |
| if(--level < 0) |
| { |
| return; |
| } |
| |
| for (size_t tidX = 0; tidX < globalSize; ++tidX) |
| { |
| ++referenceResults[tidX % len]; |
| generate_reference_results_all_eq_1D(referenceResults, len, level); |
| } |
| } |
| |
| static int check_all_eq_1D(cl_int* results, cl_int len, cl_int nesting_level) |
| { |
| std::vector<cl_int> referenceResults(len, 0); |
| generate_reference_results_all_eq_1D(referenceResults, len, nesting_level); |
| |
| for(size_t i = 0; i < len; ++i) |
| { |
| if (results[i] != referenceResults[i]) |
| { |
| log_error("ERROR: Kernel returned %d vs. expected %d, index: %d\n", results[i], referenceResults[i], i); |
| return (int)i; |
| } |
| } |
| |
| return -1; |
| } |
| |
| static const char* enqueue_1D_wg_size_all_diff[] = |
| { |
| NL, "void block_fn(__global int* res, int level, int maxGlobalWorkSize, __global int* rnd)" |
| NL, "{" |
| NL, " size_t tidX = get_global_id(0);" |
| NL, " queue_t def_q = get_default_queue();" |
| NL, " if((--level) < 0) return;" |
| NL, "" |
| NL, " void (^kernelBlock)(void) = ^{ block_fn(res, level, maxGlobalWorkSize, rnd); };" |
| NL, " uint wg = get_kernel_work_group_size(kernelBlock);" |
| NL, "" |
| NL, " const size_t gs = 8 * 8 * 8;" |
| NL, " size_t ls = rnd[tidX % maxGlobalWorkSize] % wg % gs;" |
| NL, " ls = ls? ls: 1;" |
| NL, "" |
| NL, " ndrange_t ndrange = ndrange_1D(gs, ls);" |
| NL, "" |
| NL, " // All work-items enqueues nested blocks with different levels" |
| NL, " atomic_inc(&res[tidX % maxGlobalWorkSize]);" |
| NL, " if(level >= tidX)" |
| NL, " {" |
| NL, " int enq_res = enqueue_kernel(def_q, CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange, kernelBlock);" |
| NL, " if(enq_res != CLK_SUCCESS) { res[tidX % maxGlobalWorkSize] = -1; return; }" |
| NL, " }" |
| NL, "}" |
| NL, "" |
| NL, "kernel void enqueue_1D_wg_size_all_diff(__global int* res, int level, int maxGlobalWorkSize, __global int* rnd)" |
| NL, "{" |
| NL, " block_fn(res, level, maxGlobalWorkSize, rnd);" |
| NL, "}" |
| NL |
| }; |
| |
| void generate_reference_results_all_diff_1D(std::vector<cl_int> &referenceResults, cl_int len, cl_int level) |
| { |
| size_t globalSize = (level == nestingLevel) ? len: (8 * 8 * 8); |
| if((--level) < 0) |
| { |
| return; |
| } |
| |
| for (size_t threadIdx = 0; threadIdx < globalSize; ++threadIdx) |
| { |
| ++referenceResults[threadIdx % len]; |
| if (level >= threadIdx) |
| { |
| generate_reference_results_all_diff_1D(referenceResults, len, level); |
| } |
| } |
| } |
| |
| static int check_all_diff_1D(cl_int* results, cl_int len, cl_int nesting_level) |
| { |
| std::vector<cl_int> referenceResults(len, 0); |
| generate_reference_results_all_diff_1D(referenceResults, len, nesting_level); |
| |
| for(size_t i = 0; i < len; ++i) |
| { |
| if (results[i] != referenceResults[i]) |
| { |
| log_error("ERROR: Kernel returned %d vs. expected %d, index: %d\n", results[i], referenceResults[i], i); |
| return (int)i; |
| } |
| } |
| |
| return -1; |
| } |
| |
| static const char* enqueue_2D_wg_size_single[] = |
| { |
| NL, "void block_fn(int level, int maxGlobalWorkSize, __global int* rnd, __global int* res)" |
| NL, "{" |
| NL, " size_t tidX = get_global_id(0);" |
| NL, " size_t tidY = get_global_id(1);" |
| NL, " size_t linearId = get_global_linear_id();" |
| NL, " queue_t def_q = get_default_queue();" |
| NL, " if(--level < 0) return;" |
| NL, "" |
| NL, " void (^kernelBlock)(void) = ^{ block_fn(level, maxGlobalWorkSize, rnd, res); };" |
| NL, " uint wg = get_kernel_work_group_size(kernelBlock);" |
| NL, "" |
| NL, " const size_t gs[] = { 64, 64 * 64 };" |
| NL, " size_t ls[] = { 1, rnd[tidY % maxGlobalWorkSize] % wg % gs[1] };" |
| NL, " ls[1] = ls[1]? ls[1]: 1;" |
| NL, " " |
| NL, " ndrange_t ndrange = ndrange_2D(gs, ls);" |
| NL, "" |
| NL, " // Only 1 work-item enqueues block" |
| NL, " if(tidX == 0 && tidY == 0)" |
| NL, " {" |
| NL, " atomic_inc(&res[linearId % maxGlobalWorkSize]);" |
| NL, " int enq_res = enqueue_kernel(def_q, CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange, kernelBlock);" |
| NL, " if(enq_res != CLK_SUCCESS) { res[linearId % maxGlobalWorkSize] = -1; return; }" |
| NL, " }" |
| NL, "}" |
| NL, "" |
| NL, "kernel void enqueue_2D_wg_size_single(__global int* res, int level, int maxGlobalWorkSize, __global int* rnd)" |
| NL, "{" |
| NL, " block_fn(level, maxGlobalWorkSize, rnd, res);" |
| NL, "}" |
| NL |
| }; |
| |
| static const char* enqueue_2D_wg_size_some_eq[] = |
| { |
| NL, "void block_fn(int level, int maxGlobalWorkSize, __global int* rnd, __global int* res)" |
| NL, "{" |
| NL, " size_t tidX = get_global_id(0);" |
| NL, " size_t tidY = get_global_id(1);" |
| NL, " size_t linearId = get_global_linear_id();" |
| NL, " queue_t def_q = get_default_queue();" |
| NL, " if(--level < 0) return;" |
| NL, "" |
| NL, " void (^kernelBlock)(void) = ^{ block_fn(level, maxGlobalWorkSize, rnd, res); };" |
| NL, " uint wg = get_kernel_work_group_size(kernelBlock);" |
| NL, "" |
| NL, " const size_t gs[] = { 4, 4 };" |
| NL, " size_t ls[] = { 1, rnd[tidY % maxGlobalWorkSize] % wg % gs[1] };" |
| NL, " ls[1] = ls[1]? ls[1]: 1;" |
| NL, " " |
| NL, " ndrange_t ndrange = ndrange_2D(gs, ls);" |
| NL, "" |
| NL, " // Some work-items enqueues nested blocks with the same level" |
| NL, " if((tidX < (get_global_size(0) >> 1)) && ((tidY < (get_global_size(1) >> 1)) || get_global_size(1) == 1))" |
| NL, " {" |
| NL, " atomic_inc(&res[linearId % maxGlobalWorkSize]);" |
| NL, " int enq_res = enqueue_kernel(def_q, CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange, kernelBlock);" |
| NL, " if(enq_res != CLK_SUCCESS) { res[linearId % maxGlobalWorkSize] = -1; return; }" |
| NL, " }" |
| NL, "}" |
| NL, "" |
| NL, "kernel void enqueue_2D_wg_size_some_eq(__global int* res, int level, int maxGlobalWorkSize, __global int* rnd)" |
| NL, "{" |
| NL, " block_fn(level, maxGlobalWorkSize, rnd, res);" |
| NL, "}" |
| NL |
| }; |
| |
| void generate_reference_results_some_eq_2D(std::vector<cl_int> &referenceResults, cl_int len, cl_int level) |
| { |
| size_t globalSizeX = (level == nestingLevel) ? len: 4; |
| size_t globalSizeY = (level == nestingLevel) ? 1: 4; |
| if(--level < 0) |
| { |
| return; |
| } |
| |
| for (size_t tidY = 0; tidY < globalSizeY; ++tidY) |
| { |
| for (size_t tidX = 0; tidX < globalSizeX; ++tidX) |
| { |
| if ((tidX < (globalSizeX >> 1)) && ((tidY < (globalSizeY >> 1)) || globalSizeY == 1)) |
| { |
| ++referenceResults[(globalSizeX * tidY + tidX) % len]; |
| generate_reference_results_some_eq_2D(referenceResults, len, level); |
| } |
| } |
| } |
| } |
| |
| static int check_some_eq_2D(cl_int* results, cl_int len, cl_int nesting_level) |
| { |
| std::vector<cl_int> referenceResults(len, 0); |
| generate_reference_results_some_eq_2D(referenceResults, len, nesting_level); |
| |
| for(size_t i = 0; i < len; ++i) |
| { |
| if (results[i] != referenceResults[i]) |
| { |
| log_error("ERROR: Kernel returned %d vs. expected %d, index: %d\n", results[i], referenceResults[i], i); |
| return (int)i; |
| } |
| } |
| |
| return -1; |
| } |
| |
| static const char* enqueue_2D_wg_size_some_diff[] = |
| { |
| NL, "void block_fn(int level, int maxGlobalWorkSize, __global int* rnd, __global int* res)" |
| NL, "{" |
| NL, " size_t tidX = get_global_id(0);" |
| NL, " size_t tidY = get_global_id(1);" |
| NL, " size_t linearId = get_global_linear_id();" |
| NL, " queue_t def_q = get_default_queue();" |
| NL, " if(--level < 0) return;" |
| NL, "" |
| NL, " void (^kernelBlock)(void) = ^{ block_fn(level, maxGlobalWorkSize, rnd, res); };" |
| NL, " uint wg = get_kernel_work_group_size(kernelBlock);" |
| NL, "" |
| NL, " const size_t gs[] = { 8, 8 };" |
| NL, " size_t ls[] = { 1, rnd[tidY % maxGlobalWorkSize] % wg % gs[1] };" |
| NL, " ls[1] = ls[1]? ls[1]: 1;" |
| NL, " " |
| NL, " ndrange_t ndrange = ndrange_2D(gs, ls);" |
| NL, "" |
| NL, " // Some work-items enqueues nested blocks with different levels" |
| NL, " if((tidX % 2) == 0 && (tidY % 2) == 0)" |
| NL, " {" |
| NL, " atomic_inc(&res[linearId % maxGlobalWorkSize]);" |
| NL, " if(level >= tidX && level >= tidY)" |
| NL, " {" |
| NL, " int enq_res = enqueue_kernel(def_q, CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange, kernelBlock);" |
| NL, " if(enq_res != CLK_SUCCESS) { res[linearId % maxGlobalWorkSize] = -1; return; }" |
| NL, " }" |
| NL, " }" |
| NL, "}" |
| NL, "" |
| NL, "kernel void enqueue_2D_wg_size_some_diff(__global int* res, int level, int maxGlobalWorkSize, __global int* rnd)" |
| NL, "{" |
| NL, " block_fn(level, maxGlobalWorkSize, rnd, res);" |
| NL, "}" |
| NL |
| }; |
| |
| void generate_reference_results_some_diff_2D(std::vector<cl_int> &referenceResults, cl_int len, cl_int level) |
| { |
| size_t globalSizeX = (level == nestingLevel) ? len: 8; |
| size_t globalSizeY = (level == nestingLevel) ? 1: 8; |
| if(--level < 0) |
| { |
| return; |
| } |
| |
| for (size_t tidY = 0; tidY < globalSizeY; ++tidY) |
| { |
| for (size_t tidX = 0; tidX < globalSizeX; ++tidX) |
| { |
| if ((tidX % 2) == 0 && (tidY % 2) == 0) |
| { |
| ++referenceResults[(globalSizeX * tidY + tidX) % len]; |
| if (level >= tidX && level >= tidY) |
| { |
| generate_reference_results_some_diff_2D(referenceResults, len, level); |
| } |
| } |
| } |
| } |
| } |
| |
| static int check_some_diff_2D(cl_int* results, cl_int len, cl_int nesting_level) |
| { |
| std::vector<cl_int> referenceResults(len, 0); |
| generate_reference_results_some_diff_2D(referenceResults, len, nesting_level); |
| |
| for(size_t i = 0; i < len; ++i) |
| { |
| if (results[i] != referenceResults[i]) |
| { |
| log_error("ERROR: Kernel returned %d vs. expected %d, index: %d\n", results[i], referenceResults[i], i); |
| return (int)i; |
| } |
| } |
| |
| return -1; |
| } |
| |
| static const char* enqueue_2D_wg_size_all_eq[] = |
| { |
| NL, "void block_fn(int level, int maxGlobalWorkSize, __global int* rnd, __global int* res)" |
| NL, "{" |
| NL, " size_t tidX = get_global_id(0);" |
| NL, " size_t tidY = get_global_id(1);" |
| NL, " size_t linearId = get_global_linear_id();" |
| NL, " queue_t def_q = get_default_queue();" |
| NL, " if(--level < 0) return;" |
| NL, "" |
| NL, " void (^kernelBlock)(void) = ^{ block_fn(level, maxGlobalWorkSize, rnd, res); };" |
| NL, " uint wg = get_kernel_work_group_size(kernelBlock);" |
| NL, "" |
| NL, " const size_t gs[] = { 2, 2 };" |
| NL, " size_t ls[] = { 1, rnd[tidY % maxGlobalWorkSize] % wg % gs[1] };" |
| NL, " ls[1] = ls[1]? ls[1]: 1;" |
| NL, " " |
| NL, " ndrange_t ndrange = ndrange_2D(gs, ls);" |
| NL, "" |
| NL, " // All work-items enqueues nested blocks with the same level" |
| NL, " atomic_inc(&res[linearId % maxGlobalWorkSize]);" |
| NL, " int enq_res = enqueue_kernel(def_q, CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange, kernelBlock);" |
| NL, " if(enq_res != CLK_SUCCESS) { res[linearId % maxGlobalWorkSize] = -1; return; }" |
| NL, "}" |
| NL, "" |
| NL, "kernel void enqueue_2D_wg_size_all_eq(__global int* res, int level, int maxGlobalWorkSize, __global int* rnd)" |
| NL, "{" |
| NL, " block_fn(level, maxGlobalWorkSize, rnd, res);" |
| NL, "}" |
| NL |
| }; |
| |
| void generate_reference_results_all_eq_2D(std::vector<cl_int> &referenceResults, cl_int len, cl_int level) |
| { |
| size_t globalSizeX = (level == nestingLevel) ? len: 2; |
| size_t globalSizeY = (level == nestingLevel) ? 1: 2; |
| if(--level < 0) |
| { |
| return; |
| } |
| |
| for (size_t tidY = 0; tidY < globalSizeY; ++tidY) |
| { |
| for (size_t tidX = 0; tidX < globalSizeX; ++tidX) |
| { |
| ++referenceResults[(globalSizeX * tidY + tidX) % len]; |
| generate_reference_results_all_eq_2D(referenceResults, len, level); |
| } |
| } |
| } |
| |
| static int check_all_eq_2D(cl_int* results, cl_int len, cl_int nesting_level) |
| { |
| std::vector<cl_int> referenceResults(len, 0); |
| generate_reference_results_all_eq_2D(referenceResults, len, nesting_level); |
| |
| for(size_t i = 0; i < len; ++i) |
| { |
| if (results[i] != referenceResults[i]) |
| { |
| log_error("ERROR: Kernel returned %d vs. expected %d, index: %d\n", results[i], referenceResults[i], i); |
| return (int)i; |
| } |
| } |
| |
| return -1; |
| } |
| |
| static const char* enqueue_2D_wg_size_all_diff[] = |
| { |
| NL, "void block_fn(int level, int maxGlobalWorkSize, __global int* rnd, __global int* res)" |
| NL, "{" |
| NL, " size_t tidX = get_global_id(0);" |
| NL, " size_t tidY = get_global_id(1);" |
| NL, " size_t linearId = get_global_linear_id();" |
| NL, " queue_t def_q = get_default_queue();" |
| NL, " if(--level < 0) return;" |
| NL, "" |
| NL, " void (^kernelBlock)(void) = ^{ block_fn(level, maxGlobalWorkSize, rnd, res); };" |
| NL, " uint wg = get_kernel_work_group_size(kernelBlock);" |
| NL, "" |
| NL, " size_t gs[] = { 8, 8 * 8 };" |
| NL, " size_t ls[] = { 1, rnd[tidY % maxGlobalWorkSize] % wg % gs[1] };" |
| NL, " ls[1] = ls[1]? ls[1]: 1;" |
| NL, " " |
| NL, " ndrange_t ndrange = ndrange_2D(gs, ls);" |
| NL, "" |
| NL, " // All work-items enqueues nested blocks with different levels" |
| NL, " atomic_inc(&res[linearId % maxGlobalWorkSize]);" |
| NL, " if(level >= tidX && level >= tidY)" |
| NL, " {" |
| NL, " int enq_res = enqueue_kernel(def_q, CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange, kernelBlock);" |
| NL, " if(enq_res != CLK_SUCCESS) { res[linearId % maxGlobalWorkSize] = -1; return; }" |
| NL, " }" |
| NL, "}" |
| NL, "" |
| NL, "kernel void enqueue_2D_wg_size_all_diff(__global int* res, int level, int maxGlobalWorkSize, __global int* rnd)" |
| NL, "{" |
| NL, " block_fn(level, maxGlobalWorkSize, rnd, res);" |
| NL, "}" |
| NL |
| }; |
| |
| void generate_reference_results_all_diff_2D(std::vector<cl_int> &referenceResults, cl_int len, cl_int level) |
| { |
| size_t globalSizeX = (level == nestingLevel) ? len: 8; |
| size_t globalSizeY = (level == nestingLevel) ? 1: (8 * 8); |
| if(--level < 0) |
| { |
| return; |
| } |
| |
| for (size_t tidY = 0; tidY < globalSizeY; ++tidY) |
| { |
| for (size_t tidX = 0; tidX < globalSizeX; ++tidX) |
| { |
| ++referenceResults[(globalSizeX * tidY + tidX) % len]; |
| if (level >= tidX && level >= tidY) |
| { |
| generate_reference_results_all_diff_2D(referenceResults, len, level); |
| } |
| } |
| } |
| } |
| |
| static int check_all_diff_2D(cl_int* results, cl_int len, cl_int nesting_level) |
| { |
| std::vector<cl_int> referenceResults(len, 0); |
| generate_reference_results_all_diff_2D(referenceResults, len, nesting_level); |
| |
| for(size_t i = 0; i < len; ++i) |
| { |
| if (results[i] != referenceResults[i]) |
| { |
| log_error("ERROR: Kernel returned %d vs. expected %d, index: %d\n", results[i], referenceResults[i], i); |
| return (int)i; |
| } |
| } |
| |
| return -1; |
| } |
| |
| static const char* enqueue_3D_wg_size_single[] = |
| { |
| NL, "void block_fn(int level, int maxGlobalWorkSize, __global int* rnd, __global int* res)" |
| NL, "{" |
| NL, " size_t tidX = get_global_id(0);" |
| NL, " size_t tidY = get_global_id(1);" |
| NL, " size_t tidZ = get_global_id(2);" |
| NL, " size_t linearId = get_global_linear_id();" |
| NL, " queue_t def_q = get_default_queue();" |
| NL, " if(--level < 0) return;" |
| NL, "" |
| NL, " void (^kernelBlock)(void) = ^{ block_fn(level, maxGlobalWorkSize, rnd, res); };" |
| NL, " uint wg = get_kernel_work_group_size(kernelBlock);" |
| NL, "" |
| NL, " const size_t gs[] = { 64, 64, 64 };" |
| NL, " size_t ls[] = { 1, 1, rnd[tidZ % maxGlobalWorkSize] % wg % gs[2] };" |
| NL, " ls[2] = ls[2]? ls[2]: 1;" |
| NL, " " |
| NL, " ndrange_t ndrange = ndrange_3D(gs, ls);" |
| NL, "" |
| NL, " // Only 1 work-item enqueues block" |
| NL, " if(tidX == 0 && tidY == 0 && tidZ == 0)" |
| NL, " {" |
| NL, " atomic_inc(&res[linearId % maxGlobalWorkSize]);" |
| NL, " int enq_res = enqueue_kernel(def_q, CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange, kernelBlock);" |
| NL, " if(enq_res != CLK_SUCCESS) { res[linearId % maxGlobalWorkSize] = -1; return; }" |
| NL, " }" |
| NL, "}" |
| NL, "" |
| NL, "kernel void enqueue_3D_wg_size_single(__global int* res, int level, int maxGlobalWorkSize, __global int* rnd)" |
| NL, "{" |
| NL, " block_fn(level, maxGlobalWorkSize, rnd, res);" |
| NL, "}" |
| NL |
| }; |
| |
| static const char* enqueue_3D_wg_size_some_eq[] = |
| { |
| NL, "void block_fn(int level, int maxGlobalWorkSize, __global int* rnd, __global int* res)" |
| NL, "{" |
| NL, " size_t tidX = get_global_id(0);" |
| NL, " size_t tidY = get_global_id(1);" |
| NL, " size_t tidZ = get_global_id(2);" |
| NL, " size_t linearId = get_global_linear_id();" |
| NL, " queue_t def_q = get_default_queue();" |
| NL, " if(--level < 0) return;" |
| NL, "" |
| NL, " void (^kernelBlock)(void) = ^{ block_fn(level, maxGlobalWorkSize, rnd, res); };" |
| NL, " uint wg = get_kernel_work_group_size(kernelBlock);" |
| NL, "" |
| NL, " const size_t gs[] = { 4, 4, 4 };" |
| NL, " size_t ls[] = { 1, 1, rnd[tidZ % maxGlobalWorkSize] % wg % gs[2] };" |
| NL, " ls[2] = ls[2]? ls[2]: 1;" |
| NL, " " |
| NL, " ndrange_t ndrange = ndrange_3D(gs, ls);" |
| NL, "" |
| NL, " // Some work-items enqueues nested blocks with the same level" |
| NL, " if((tidX < (get_global_size(0) >> 1)) && " |
| NL, " ((tidY < (get_global_size(1) >> 1)) || get_global_size(1) == 1) &&" |
| NL, " ((tidZ < (get_global_size(2) >> 1)) || get_global_size(2) == 1))" |
| NL, " {" |
| NL, " atomic_inc(&res[linearId % maxGlobalWorkSize]);" |
| NL, " int enq_res = enqueue_kernel(def_q, CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange, kernelBlock);" |
| NL, " if(enq_res != CLK_SUCCESS) { res[linearId % maxGlobalWorkSize] = -1; return; }" |
| NL, " }" |
| NL, "}" |
| NL, "" |
| NL, "kernel void enqueue_3D_wg_size_some_eq(__global int* res, int level, int maxGlobalWorkSize, __global int* rnd)" |
| NL, "{" |
| NL, " block_fn(level, maxGlobalWorkSize, rnd, res);" |
| NL, "}" |
| NL |
| }; |
| |
| void generate_reference_results_some_eq_3D(std::vector<cl_int> &referenceResults, cl_int len, cl_int level) |
| { |
| size_t globalSizeX = (level == nestingLevel) ? len: 4; |
| size_t globalSizeY = (level == nestingLevel) ? 1: 4; |
| size_t globalSizeZ = (level == nestingLevel) ? 1: 4; |
| if(--level < 0) |
| { |
| return; |
| } |
| |
| for (size_t tidZ = 0; tidZ < globalSizeZ; ++tidZ) |
| { |
| for (size_t tidY = 0; tidY < globalSizeY; ++tidY) |
| { |
| for (size_t tidX = 0; tidX < globalSizeX; ++tidX) |
| { |
| if ((tidX < (globalSizeX >> 1)) && ((tidY < (globalSizeY >> 1)) || globalSizeY == 1) && ((tidZ < (globalSizeZ >> 1)) || globalSizeZ == 1)) |
| { |
| ++referenceResults[(globalSizeX * globalSizeY * tidZ + globalSizeX * tidY + tidX) % len]; |
| generate_reference_results_some_eq_3D(referenceResults, len, level); |
| } |
| } |
| } |
| } |
| } |
| |
| static int check_some_eq_3D(cl_int* results, cl_int len, cl_int nesting_level) |
| { |
| std::vector<cl_int> referenceResults(len, 0); |
| generate_reference_results_some_eq_3D(referenceResults, len, nesting_level); |
| |
| for(size_t i = 0; i < len; ++i) |
| { |
| if (results[i] != referenceResults[i]) |
| { |
| log_error("ERROR: Kernel returned %d vs. expected %d, index: %d\n", results[i], referenceResults[i], i); |
| return (int)i; |
| } |
| } |
| |
| return -1; |
| } |
| |
| static const char* enqueue_3D_wg_size_some_diff[] = |
| { |
| NL, "void block_fn(int level, int maxGlobalWorkSize, __global int* rnd, __global int* res)" |
| NL, "{" |
| NL, " size_t tidX = get_global_id(0);" |
| NL, " size_t tidY = get_global_id(1);" |
| NL, " size_t tidZ = get_global_id(2);" |
| NL, " size_t linearId = get_global_linear_id();" |
| NL, " queue_t def_q = get_default_queue();" |
| NL, " if(--level < 0) return;" |
| NL, "" |
| NL, " void (^kernelBlock)(void) = ^{ block_fn(level, maxGlobalWorkSize, rnd, res); };" |
| NL, " uint wg = get_kernel_work_group_size(kernelBlock);" |
| NL, "" |
| NL, " const size_t gs[] = { 8, 8, 8 };" |
| NL, " size_t ls[] = { 1, 1, rnd[tidZ % maxGlobalWorkSize] % wg % gs[2] };" |
| NL, " ls[2] = ls[2]? ls[2]: 1;" |
| NL, " " |
| NL, " ndrange_t ndrange = ndrange_3D(gs, ls);" |
| NL, "" |
| NL, " // Some work-items enqueues nested blocks with different levels" |
| NL, " if((tidX % 2) == 0 && (tidY % 2) == 0 && (tidZ % 2) == 0)" |
| NL, " {" |
| NL, " atomic_inc(&res[linearId % maxGlobalWorkSize]);" |
| NL, " if(level >= tidX && level >= tidY && level >= tidZ)" |
| NL, " {" |
| NL, " int enq_res = enqueue_kernel(def_q, CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange, kernelBlock);" |
| NL, " if(enq_res != CLK_SUCCESS) { res[linearId % maxGlobalWorkSize] = -1; return; }" |
| NL, " }" |
| NL, " }" |
| NL, "}" |
| NL, "" |
| NL, "kernel void enqueue_3D_wg_size_some_diff(__global int* res, int level, int maxGlobalWorkSize, __global int* rnd)" |
| NL, "{" |
| NL, " block_fn(level, maxGlobalWorkSize, rnd, res);" |
| NL, "}" |
| NL |
| }; |
| |
| void generate_reference_results_some_diff_3D(std::vector<cl_int> &referenceResults, cl_int len, cl_int level) |
| { |
| size_t globalSizeX = (level == nestingLevel) ? len: 8; |
| size_t globalSizeY = (level == nestingLevel) ? 1: 8; |
| size_t globalSizeZ = (level == nestingLevel) ? 1: 8; |
| if(--level < 0) |
| { |
| return; |
| } |
| |
| for (size_t tidZ = 0; tidZ < globalSizeZ; ++tidZ) |
| { |
| for (size_t tidY = 0; tidY < globalSizeY; ++tidY) |
| { |
| for (size_t tidX = 0; tidX < globalSizeX; ++tidX) |
| { |
| if ((tidX % 2) == 0 && (tidY % 2) == 0 && (tidZ % 2) == 0) |
| { |
| ++referenceResults[(globalSizeX * globalSizeY * tidZ + globalSizeX * tidY + tidX) % len]; |
| if (level >= tidX && level >= tidY && level >= tidZ) |
| { |
| generate_reference_results_some_diff_3D(referenceResults, len, level); |
| } |
| } |
| } |
| } |
| } |
| } |
| |
| static int check_some_diff_3D(cl_int* results, cl_int len, cl_int nesting_level) |
| { |
| std::vector<cl_int> referenceResults(len, 0); |
| generate_reference_results_some_diff_3D(referenceResults, len, nesting_level); |
| |
| for(size_t i = 0; i < len; ++i) |
| { |
| if (results[i] != referenceResults[i]) |
| { |
| log_error("ERROR: Kernel returned %d vs. expected %d, index: %d\n", results[i], referenceResults[i], i); |
| return (int)i; |
| } |
| } |
| |
| return -1; |
| } |
| |
| static const char* enqueue_3D_wg_size_all_eq[] = |
| { |
| NL, "void block_fn(int level, int maxGlobalWorkSize, __global int* rnd, __global int* res)" |
| NL, "{" |
| NL, " size_t tidX = get_global_id(0);" |
| NL, " size_t tidY = get_global_id(1);" |
| NL, " size_t tidZ = get_global_id(2);" |
| NL, " size_t linearId = get_global_linear_id();" |
| NL, " queue_t def_q = get_default_queue();" |
| NL, " if(--level < 0) return;" |
| NL, "" |
| NL, " void (^kernelBlock)(void) = ^{ block_fn(level, maxGlobalWorkSize, rnd, res); };" |
| NL, " uint wg = get_kernel_work_group_size(kernelBlock);" |
| NL, "" |
| NL, " const size_t gs[] = { 2, 2, 2 };" |
| NL, " size_t ls[] = { 1, 1, rnd[tidZ % maxGlobalWorkSize] % wg % gs[2] };" |
| NL, " ls[2] = ls[2]? ls[2]: 1;" |
| NL, " " |
| NL, " ndrange_t ndrange = ndrange_3D(gs, ls);" |
| NL, "" |
| NL, " // All work-items enqueues nested blocks with the same level" |
| NL, " atomic_inc(&res[linearId % maxGlobalWorkSize]);" |
| NL, " int enq_res = enqueue_kernel(def_q, CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange, kernelBlock);" |
| NL, " if(enq_res != CLK_SUCCESS) { res[linearId % maxGlobalWorkSize] = -1; return; }" |
| NL, "}" |
| NL, "" |
| NL, "kernel void enqueue_3D_wg_size_all_eq(__global int* res, int level, int maxGlobalWorkSize, __global int* rnd)" |
| NL, "{" |
| NL, " block_fn(level, maxGlobalWorkSize, rnd, res);" |
| NL, "}" |
| NL |
| }; |
| |
| void generate_reference_results_all_eq_3D(std::vector<cl_int> &referenceResults, cl_int len, cl_int level) |
| { |
| size_t globalSizeX = (level == nestingLevel) ? len: 2; |
| size_t globalSizeY = (level == nestingLevel) ? 1: 2; |
| size_t globalSizeZ = (level == nestingLevel) ? 1: 2; |
| if(--level < 0) |
| { |
| return; |
| } |
| |
| for (size_t tidZ = 0; tidZ < globalSizeZ; ++tidZ) |
| { |
| for (size_t tidY = 0; tidY < globalSizeY; ++tidY) |
| { |
| for (size_t tidX = 0; tidX < globalSizeX; ++tidX) |
| { |
| ++referenceResults[(globalSizeX * globalSizeY * tidZ + globalSizeX * tidY + tidX) % len]; |
| generate_reference_results_all_eq_3D(referenceResults, len, level); |
| } |
| } |
| } |
| } |
| |
| static int check_all_eq_3D(cl_int* results, cl_int len, cl_int nesting_level) |
| { |
| std::vector<cl_int> referenceResults(len, 0); |
| generate_reference_results_all_eq_3D(referenceResults, len, nesting_level); |
| |
| for(size_t i = 0; i < len; ++i) |
| { |
| if (results[i] != referenceResults[i]) |
| { |
| log_error("ERROR: Kernel returned %d vs. expected %d, index: %d\n", results[i], referenceResults[i], i); |
| return (int)i; |
| } |
| } |
| |
| return -1; |
| } |
| |
| static const char* enqueue_3D_wg_size_all_diff[] = |
| { |
| NL, "void block_fn(int level, int maxGlobalWorkSize, __global int* rnd, __global int* res)" |
| NL, "{" |
| NL, " size_t tidX = get_global_id(0);" |
| NL, " size_t tidY = get_global_id(1);" |
| NL, " size_t tidZ = get_global_id(2);" |
| NL, " size_t linearId = get_global_linear_id();" |
| NL, " queue_t def_q = get_default_queue();" |
| NL, " if(--level < 0) return;" |
| NL, "" |
| NL, " void (^kernelBlock)(void) = ^{ block_fn(level, maxGlobalWorkSize, rnd, res); };" |
| NL, " uint wg = get_kernel_work_group_size(kernelBlock);" |
| NL, "" |
| NL, " const size_t gs[] = { 8, 8, 8 };" |
| NL, " size_t ls[] = { 1, 1, rnd[tidZ % maxGlobalWorkSize] % wg % gs[2] };" |
| NL, " ls[2] = ls[2]? ls[2]: 1;" |
| NL, " " |
| NL, " ndrange_t ndrange = ndrange_3D(gs, ls);" |
| NL, "" |
| NL, " // All work-items enqueues nested blocks with different levels" |
| NL, " atomic_inc(&res[linearId % maxGlobalWorkSize]);" |
| NL, " if(level >= tidX && level >= tidY && level >= tidZ)" |
| NL, " {" |
| NL, " int enq_res = enqueue_kernel(def_q, CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange, kernelBlock);" |
| NL, " if(enq_res != CLK_SUCCESS) { res[linearId % maxGlobalWorkSize] = -1; return; }" |
| NL, " }" |
| NL, "}" |
| NL, "" |
| NL, "kernel void enqueue_3D_wg_size_all_diff(__global int* res, int level, int maxGlobalWorkSize, __global int* rnd)" |
| NL, "{" |
| NL, " block_fn(level, maxGlobalWorkSize, rnd, res);" |
| NL, "}" |
| NL |
| }; |
| |
| void generate_reference_results_all_diff_3D(std::vector<cl_int> &referenceResults, cl_int len, cl_int level) |
| { |
| size_t globalSizeX = (level == nestingLevel) ? len: 8; |
| size_t globalSizeY = (level == nestingLevel) ? 1: 8; |
| size_t globalSizeZ = (level == nestingLevel) ? 1: 8; |
| if(--level < 0) |
| { |
| return; |
| } |
| |
| for (size_t tidZ = 0; tidZ < globalSizeZ; ++tidZ) |
| { |
| for (size_t tidY = 0; tidY < globalSizeY; ++tidY) |
| { |
| for (size_t tidX = 0; tidX < globalSizeX; ++tidX) |
| { |
| ++referenceResults[(globalSizeX * globalSizeY * tidZ + globalSizeX * tidY + tidX) % len]; |
| if (level >= tidX && level >= tidY && level >= tidZ) |
| { |
| generate_reference_results_all_diff_3D(referenceResults, len, level); |
| } |
| } |
| } |
| } |
| } |
| |
| static int check_all_diff_3D(cl_int* results, cl_int len, cl_int nesting_level) |
| { |
| std::vector<cl_int> referenceResults(len, 0); |
| generate_reference_results_all_diff_3D(referenceResults, len, nesting_level); |
| |
| for(size_t i = 0; i < len; ++i) |
| { |
| if (results[i] != referenceResults[i]) |
| { |
| log_error("ERROR: Kernel returned %d vs. expected %d, index: %d\n", results[i], referenceResults[i], i); |
| return (int)i; |
| } |
| } |
| |
| return -1; |
| } |
| |
| static const char* enqueue_mix_wg_size_single[] = |
| { |
| NL, "void block_fn(int level, int maxGlobalWorkSize, __global int* rnd, __global int* res)" |
| NL, "{" |
| NL, " size_t tidX = get_global_id(0);" |
| NL, " size_t tidY = get_global_id(1);" |
| NL, " size_t tidZ = get_global_id(2);" |
| NL, " size_t linearId = get_global_linear_id();" |
| NL, " queue_t def_q = get_default_queue();" |
| NL, " if(--level < 0) return;" |
| NL, "" |
| NL, " void (^kernelBlock)(void) = ^{ block_fn(level, maxGlobalWorkSize, rnd, res); };" |
| NL, " uint wg = get_kernel_work_group_size(kernelBlock);" |
| NL, "" |
| NL, " ndrange_t ndrange;" |
| NL, " switch((linearId + level) % 3)" |
| NL, " {" |
| NL, " case 0:" |
| NL, " {" |
| NL, " const size_t gs = 64 * 64 * 64;" |
| NL, " size_t ls = rnd[tidX % maxGlobalWorkSize] % wg % gs;" |
| NL, " ls = ls? ls: 1;" |
| NL, " ndrange = ndrange_1D(gs, ls);" |
| NL, " }" |
| NL, " break;" |
| NL, " case 1:" |
| NL, " {" |
| NL, " const size_t gs[] = { 64, 64 * 64 };" |
| NL, " size_t ls[] = { 1, rnd[tidY % maxGlobalWorkSize] % wg % gs[1] };" |
| NL, " ls[1] = ls[1]? ls[1]: 1;" |
| NL, " ndrange = ndrange_2D(gs, ls);" |
| NL, " }" |
| NL, " break;" |
| NL, " case 2:" |
| NL, " {" |
| NL, " const size_t gs[] = { 64, 64, 64 };" |
| NL, " size_t ls[] = { 1, 1, rnd[tidZ % maxGlobalWorkSize] % wg % gs[2] };" |
| NL, " ls[2] = ls[2]? ls[2]: 1;" |
| NL, " ndrange = ndrange_3D(gs, ls);" |
| NL, " }" |
| NL, " break;" |
| NL, " default:" |
| NL, " break;" |
| NL, " }" |
| NL, "" |
| NL, " // Only 1 work-item enqueues block" |
| NL, " if(tidX == 0 && (tidY == 0 || get_global_size(1) == 1) && (tidZ == 0 || get_global_size(2) == 1))" |
| NL, " {" |
| NL, " atomic_inc(&res[linearId % maxGlobalWorkSize]);" |
| NL, " int enq_res = enqueue_kernel(def_q, CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange, kernelBlock);" |
| NL, " if(enq_res != CLK_SUCCESS) { res[linearId % maxGlobalWorkSize] = -1; return; }" |
| NL, " }" |
| NL, "}" |
| NL, "" |
| NL, "kernel void enqueue_mix_wg_size_single(__global int* res, int level, int maxGlobalWorkSize, __global int* rnd)" |
| NL, "{" |
| NL, " block_fn(level, maxGlobalWorkSize, rnd, res);" |
| NL, "}" |
| NL |
| }; |
| |
| static const char* enqueue_mix_wg_size_some_eq[] = |
| { |
| NL, "void block_fn(int level, int maxGlobalWorkSize, __global int* rnd, __global int* res)" |
| NL, "{" |
| NL, " queue_t def_q = get_default_queue();" |
| NL, " size_t tidX = get_global_id(0);" |
| NL, " size_t tidY = get_global_id(1);" |
| NL, " size_t tidZ = get_global_id(2);" |
| NL, " size_t linearId = get_global_linear_id();" |
| NL, " if(--level < 0) return;" |
| NL, "" |
| NL, " void (^kernelBlock)(void) = ^{ block_fn(level, maxGlobalWorkSize, rnd, res); };" |
| NL, " uint wg = get_kernel_work_group_size(kernelBlock);" |
| NL, "" |
| NL, " ndrange_t ndrange;" |
| NL, " switch((linearId + level) % 3)" |
| NL, " {" |
| NL, " case 0:" |
| NL, " {" |
| NL, " const size_t gs = 2 * 4 * 4;" |
| NL, " size_t ls = rnd[tidX % maxGlobalWorkSize] % wg % gs;" |
| NL, " ls = ls? ls: 1;" |
| NL, " ndrange = ndrange_1D(gs, ls);" |
| NL, " }" |
| NL, " break;" |
| NL, " case 1:" |
| NL, " {" |
| NL, " const size_t gs[] = { 2, 4 * 4 };" |
| NL, " size_t ls[] = { 1, rnd[tidY % maxGlobalWorkSize] % wg % gs[1] };" |
| NL, " ls[1] = ls[1]? ls[1]: 1;" |
| NL, " ndrange = ndrange_2D(gs, ls);" |
| NL, " }" |
| NL, " break;" |
| NL, " case 2:" |
| NL, " {" |
| NL, " const size_t gs[] = { 2, 4, 4 };" |
| NL, " size_t ls[] = { 1, 1, rnd[tidZ % maxGlobalWorkSize] % wg % gs[2] };" |
| NL, " ls[2] = ls[2]? ls[2]: 1;" |
| NL, " ndrange = ndrange_3D(gs, ls);" |
| NL, " }" |
| NL, " break;" |
| NL, " default:" |
| NL, " break;" |
| NL, " }" |
| NL, "" |
| NL, " // Some work-items enqueues nested blocks with the same level" |
| NL, " size_t globalSizeX = get_global_size(0);" |
| NL, " size_t globalSizeY = get_global_size(1);" |
| NL, " size_t globalSizeZ = get_global_size(2);" |
| NL, " if((tidX < (globalSizeX >> 1)) && ((tidY < (globalSizeY >> 1)) || globalSizeY == 1) && ((tidZ < (globalSizeZ >> 1)) || globalSizeZ == 1))" |
| NL, " {" |
| NL, " atomic_inc(&res[linearId % maxGlobalWorkSize]);" |
| NL, " int enq_res = enqueue_kernel(def_q, CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange, kernelBlock);" |
| NL, " if(enq_res != CLK_SUCCESS) { res[linearId % maxGlobalWorkSize] = -1; return; }" |
| NL, " }" |
| NL, "}" |
| NL, "" |
| NL, "kernel void enqueue_mix_wg_size_some_eq(__global int* res, int level, int maxGlobalWorkSize, __global int* rnd)" |
| NL, "{" |
| NL, " block_fn(level, maxGlobalWorkSize, rnd, res);" |
| NL, "}" |
| NL |
| }; |
| |
| void generate_reference_results_some_eq_mix(std::vector<cl_int> &referenceResults, cl_int len, cl_int level, cl_int dim) |
| { |
| size_t globalSizeX = 1, globalSizeY = 1, globalSizeZ = 1; |
| switch (dim) |
| { |
| case 0: |
| globalSizeX = (level == nestingLevel) ? len: (2 * 4 * 4); |
| break; |
| case 1: |
| globalSizeX = 2; |
| globalSizeY = 4 * 4; |
| break; |
| case 2: |
| globalSizeX = 2; |
| globalSizeY = 4; |
| globalSizeZ = 4; |
| break; |
| default: |
| break; |
| } |
| |
| if(--level < 0) |
| { |
| return; |
| } |
| |
| for (size_t tidZ = 0; tidZ < globalSizeZ; ++tidZ) |
| { |
| for (size_t tidY = 0; tidY < globalSizeY; ++tidY) |
| { |
| for (size_t tidX = 0; tidX < globalSizeX; ++tidX) |
| { |
| size_t linearID = globalSizeX * globalSizeY * tidZ + globalSizeX * tidY + tidX; |
| cl_int nextDim = (linearID + level) % 3; |
| if ((tidX < (globalSizeX >> 1)) && ((tidY < (globalSizeY >> 1)) || globalSizeY == 1) && ((tidZ < (globalSizeZ >> 1)) || globalSizeZ == 1)) |
| { |
| ++referenceResults[linearID % len]; |
| generate_reference_results_some_eq_mix(referenceResults, len, level, nextDim); |
| } |
| } |
| } |
| } |
| } |
| |
| static int check_some_eq_mix(cl_int* results, cl_int len, cl_int nesting_level) |
| { |
| std::vector<cl_int> referenceResults(len, 0); |
| generate_reference_results_some_eq_mix(referenceResults, len, nesting_level, 0); |
| |
| for(size_t i = 0; i < len; ++i) |
| { |
| if (results[i] != referenceResults[i]) |
| { |
| log_error("ERROR: Kernel returned %d vs. expected %d, index: %d\n", results[i], referenceResults[i], i); |
| return (int)i; |
| } |
| } |
| |
| return -1; |
| } |
| |
| static const char* enqueue_mix_wg_size_some_diff[] = |
| { |
| NL, "void block_fn(int level, int maxGlobalWorkSize, __global int* rnd, __global int* res)" |
| NL, "{" |
| NL, " queue_t def_q = get_default_queue();" |
| NL, " size_t tidX = get_global_id(0);" |
| NL, " size_t tidY = get_global_id(1);" |
| NL, " size_t tidZ = get_global_id(2);" |
| NL, " size_t linearId = get_global_linear_id();" |
| NL, " if(--level < 0) return;" |
| NL, "" |
| NL, " void (^kernelBlock)(void) = ^{ block_fn(level, maxGlobalWorkSize, rnd, res); };" |
| NL, " uint wg = get_kernel_work_group_size(kernelBlock);" |
| NL, "" |
| NL, " ndrange_t ndrange;" |
| NL, " switch((linearId + level) % 3)" |
| NL, " {" |
| NL, " case 0:" |
| NL, " {" |
| NL, " const size_t gs = 8 * 8 * 8;" |
| NL, " size_t ls = rnd[tidX % maxGlobalWorkSize] % wg % gs;" |
| NL, " ls = ls? ls: 1;" |
| NL, " ndrange = ndrange_1D(gs, ls);" |
| NL, " }" |
| NL, " break;" |
| NL, " case 1:" |
| NL, " {" |
| NL, " const size_t gs[] = { 8, 8 * 8 };" |
| NL, " size_t ls[] = { 1, rnd[tidY % maxGlobalWorkSize] % wg % gs[1] };" |
| NL, " ls[1] = ls[1]? ls[1]: 1;" |
| NL, " ndrange = ndrange_2D(gs, ls);" |
| NL, " }" |
| NL, " break;" |
| NL, " case 2:" |
| NL, " {" |
| NL, " const size_t gs[] = { 8, 8, 8 };" |
| NL, " size_t ls[] = { 1, 1, rnd[tidZ % maxGlobalWorkSize] % wg % gs[2] };" |
| NL, " ls[2] = ls[2]? ls[2]: 1;" |
| NL, " ndrange = ndrange_3D(gs, ls);" |
| NL, " }" |
| NL, " break;" |
| NL, " default:" |
| NL, " break;" |
| NL, " }" |
| NL, "" |
| NL, " // Some work-items enqueues nested blocks with different levels" |
| NL, " if((tidX % 2) == 0 && (tidY % 2) == 0 && (tidZ % 2) == 0)" |
| NL, " {" |
| NL, " atomic_inc(&res[linearId % maxGlobalWorkSize]);" |
| NL, " if(level >= tidX && level >= tidY && level >= tidZ)" |
| NL, " {" |
| NL, " int enq_res = enqueue_kernel(def_q, CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange, kernelBlock);" |
| NL, " if(enq_res != CLK_SUCCESS) { res[linearId % maxGlobalWorkSize] = -1; return; }" |
| NL, " }" |
| NL, " }" |
| NL, "}" |
| NL, "" |
| NL, "kernel void enqueue_mix_wg_size_some_diff(__global int* res, int level, int maxGlobalWorkSize, __global int* rnd)" |
| NL, "{" |
| NL, " block_fn(level, maxGlobalWorkSize, rnd, res);" |
| NL, "}" |
| NL |
| }; |
| |
| void generate_reference_results_some_diff_mix(std::vector<cl_int> &referenceResults, cl_int len, cl_int level, cl_int dim) |
| { |
| size_t globalSizeX = 1, globalSizeY = 1, globalSizeZ = 1; |
| switch (dim) |
| { |
| case 0: |
| globalSizeX = (level == nestingLevel) ? len: (8 * 8 * 8); |
| break; |
| case 1: |
| globalSizeX = 8; |
| globalSizeY = 8 * 8; |
| break; |
| case 2: |
| globalSizeX = 8; |
| globalSizeY = 8; |
| globalSizeZ = 8; |
| break; |
| default: |
| return; |
| break; |
| } |
| |
| if(--level < 0) |
| { |
| return; |
| } |
| |
| for (size_t tidZ = 0; tidZ < globalSizeZ; ++tidZ) |
| { |
| for (size_t tidY = 0; tidY < globalSizeY; ++tidY) |
| { |
| for (size_t tidX = 0; tidX < globalSizeX; ++tidX) |
| { |
| size_t linearID = globalSizeX * globalSizeY * tidZ + globalSizeX * tidY + tidX; |
| cl_int nextDim = (linearID + level) % 3; |
| if ((tidX % 2) == 0 && (tidY % 2) == 0 && (tidZ % 2) == 0) |
| { |
| ++referenceResults[linearID % len]; |
| if (level >= tidX && level >= tidY && level >= tidZ) |
| { |
| generate_reference_results_some_diff_mix(referenceResults, len, level, nextDim); |
| } |
| } |
| } |
| } |
| } |
| } |
| |
| static int check_some_diff_mix(cl_int* results, cl_int len, cl_int nesting_level) |
| { |
| std::vector<cl_int> referenceResults(len, 0); |
| generate_reference_results_some_diff_mix(referenceResults, len, nesting_level, 0); |
| |
| for(size_t i = 0; i < len; ++i) |
| { |
| if (results[i] != referenceResults[i]) |
| { |
| log_error("ERROR: Kernel returned %d vs. expected %d, index: %d\n", results[i], referenceResults[i], i); |
| return (int)i; |
| } |
| } |
| |
| return -1; |
| } |
| |
| static const char* enqueue_mix_wg_size_all_eq[] = |
| { |
| NL, "void block_fn(int level, int maxGlobalWorkSize, __global int* rnd, __global int* res)" |
| NL, "{" |
| NL, " queue_t def_q = get_default_queue();" |
| NL, " size_t tidX = get_global_id(0);" |
| NL, " size_t tidY = get_global_id(1);" |
| NL, " size_t tidZ = get_global_id(2);" |
| NL, " size_t linearId = get_global_linear_id();" |
| NL, " if(--level < 0) return;" |
| NL, "" |
| NL, " void (^kernelBlock)(void) = ^{ block_fn(level, maxGlobalWorkSize, rnd, res); };" |
| NL, " uint wg = get_kernel_work_group_size(kernelBlock);" |
| NL, "" |
| NL, " ndrange_t ndrange;" |
| NL, " switch((linearId + level) % 3)" |
| NL, " {" |
| NL, " case 0:" |
| NL, " {" |
| NL, " const size_t gs = 2 * 2 * 2;" |
| NL, " size_t ls = rnd[tidX % maxGlobalWorkSize] % wg % gs;" |
| NL, " ls = ls? ls: 1;" |
| NL, " ndrange = ndrange_1D(gs, ls);" |
| NL, " }" |
| NL, " break;" |
| NL, " case 1:" |
| NL, " {" |
| NL, " const size_t gs[] = { 2, 2 * 2 };" |
| NL, " size_t ls[] = { 1, rnd[tidY % maxGlobalWorkSize] % wg % gs[1] };" |
| NL, " ls[1] = ls[1]? ls[1]: 1;" |
| NL, " ndrange = ndrange_2D(gs, ls);" |
| NL, " }" |
| NL, " break;" |
| NL, " case 2:" |
| NL, " {" |
| NL, " const size_t gs[] = { 2, 2, 2 };" |
| NL, " size_t ls[] = { 1, 1, rnd[tidZ % maxGlobalWorkSize] % wg % gs[2] };" |
| NL, " ls[2] = ls[2]? ls[2]: 1;" |
| NL, " ndrange = ndrange_3D(gs, ls);" |
| NL, " }" |
| NL, " break;" |
| NL, " default:" |
| NL, " break;" |
| NL, " }" |
| NL, "" |
| NL, " // All work-items enqueues nested blocks with the same level" |
| NL, " atomic_inc(&res[linearId % maxGlobalWorkSize]);" |
| NL, " int enq_res = enqueue_kernel(def_q, CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange, kernelBlock);" |
| NL, " if(enq_res != CLK_SUCCESS) { res[linearId % maxGlobalWorkSize] = -1; return; }" |
| NL, "}" |
| NL, "" |
| NL, "kernel void enqueue_mix_wg_size_all_eq(__global int* res, int level, int maxGlobalWorkSize, __global int* rnd)" |
| NL, "{" |
| NL, " block_fn(level, maxGlobalWorkSize, rnd, res);" |
| NL, "}" |
| NL |
| }; |
| |
| void generate_reference_results_all_eq_mix(std::vector<cl_int> &referenceResults, cl_int len, cl_int level, cl_int dim) |
| { |
| size_t globalSizeX = 1, globalSizeY = 1, globalSizeZ = 1; |
| switch (dim) |
| { |
| case 0: |
| globalSizeX = (level == nestingLevel) ? len: (2 * 2 * 2); |
| break; |
| case 1: |
| globalSizeX = 2; |
| globalSizeY = 2 * 2; |
| break; |
| case 2: |
| globalSizeX = 2; |
| globalSizeY = 2; |
| globalSizeZ = 2; |
| break; |
| default: |
| break; |
| } |
| |
| if(--level < 0) |
| { |
| return; |
| } |
| |
| for (size_t tidZ = 0; tidZ < globalSizeZ; ++tidZ) |
| { |
| for (size_t tidY = 0; tidY < globalSizeY; ++tidY) |
| { |
| for (size_t tidX = 0; tidX < globalSizeX; ++tidX) |
| { |
| size_t linearID = globalSizeX * globalSizeY * tidZ + globalSizeX * tidY + tidX; |
| cl_int nextDim = (linearID + level) % 3; |
| ++referenceResults[linearID % len]; |
| generate_reference_results_all_eq_mix(referenceResults, len, level, nextDim); |
| } |
| } |
| } |
| } |
| |
| static int check_all_eq_mix(cl_int* results, cl_int len, cl_int nesting_level) |
| { |
| std::vector<cl_int> referenceResults(len, 0); |
| generate_reference_results_all_eq_mix(referenceResults, len, nesting_level, 0); |
| |
| for(size_t i = 0; i < len; ++i) |
| { |
| if (results[i] != referenceResults[i]) |
| { |
| log_error("ERROR: Kernel returned %d vs. expected %d, index: %d\n", results[i], referenceResults[i], i); |
| return (int)i; |
| } |
| } |
| |
| return -1; |
| } |
| |
| static const char* enqueue_mix_wg_size_all_diff[] = |
| { |
| NL, "void block_fn(int level, int maxGlobalWorkSize, __global int* rnd, __global int* res)" |
| NL, "{" |
| NL, " queue_t def_q = get_default_queue();" |
| NL, " size_t tidX = get_global_id(0);" |
| NL, " size_t tidY = get_global_id(1);" |
| NL, " size_t tidZ = get_global_id(2);" |
| NL, " size_t linearId = get_global_linear_id();" |
| NL, " if(--level < 0) return;" |
| NL, "" |
| NL, " void (^kernelBlock)(void) = ^{ block_fn(level, maxGlobalWorkSize, rnd, res); };" |
| NL, " uint wg = get_kernel_work_group_size(kernelBlock);" |
| NL, "" |
| NL, " ndrange_t ndrange;" |
| NL, " switch((linearId + level) % 3)" |
| NL, " {" |
| NL, " case 0:" |
| NL, " {" |
| NL, " const size_t gs = 8 * 8 * 8;" |
| NL, " size_t ls = rnd[tidX % maxGlobalWorkSize] % wg % gs;" |
| NL, " ls = ls? ls: 1;" |
| NL, " ndrange = ndrange_1D(gs, ls);" |
| NL, " }" |
| NL, " break;" |
| NL, " case 1:" |
| NL, " {" |
| NL, " const size_t gs[] = { 8, 8 * 8 };" |
| NL, " size_t ls[] = { 1, rnd[tidY % maxGlobalWorkSize] % wg % gs[1] };" |
| NL, " ls[1] = ls[1]? ls[1]: 1;" |
| NL, " ndrange = ndrange_2D(gs, ls);" |
| NL, " }" |
| NL, " break;" |
| NL, " case 2:" |
| NL, " {" |
| NL, " const size_t gs[] = { 8, 8, 8 };" |
| NL, " size_t ls[] = { 1, 1, rnd[tidZ % maxGlobalWorkSize] % wg % gs[2] };" |
| NL, " ls[2] = ls[2]? ls[2]: 1;" |
| NL, " ndrange = ndrange_3D(gs, ls);" |
| NL, " }" |
| NL, " break;" |
| NL, " default:" |
| NL, " break;" |
| NL, " }" |
| NL, "" |
| NL, " // All work-items enqueues nested blocks with different levels" |
| NL, " atomic_inc(&res[linearId % maxGlobalWorkSize]);" |
| NL, " if(level >= tidX && level >= tidY && level >= tidZ)" |
| NL, " {" |
| NL, " int enq_res = enqueue_kernel(def_q, CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange, kernelBlock);" |
| NL, " if(enq_res != CLK_SUCCESS) { res[linearId % maxGlobalWorkSize] = -1; return; }" |
| NL, " }" |
| NL, "}" |
| NL, "" |
| NL, "kernel void enqueue_mix_wg_size_all_diff(__global int* res, int level, int maxGlobalWorkSize, __global int* rnd)" |
| NL, "{" |
| NL, " block_fn(level, maxGlobalWorkSize, rnd, res);" |
| NL, "}" |
| NL |
| }; |
| |
| void generate_reference_results_all_diff_mix(std::vector<cl_int> &referenceResults, cl_int len, cl_int level, cl_int dim) |
| { |
| size_t globalSizeX = 1, globalSizeY = 1, globalSizeZ = 1; |
| switch (dim) |
| { |
| case 0: |
| globalSizeX = (level == nestingLevel) ? len: (8 * 8 * 8); |
| break; |
| case 1: |
| globalSizeX = 8; |
| globalSizeY = 8 * 8; |
| break; |
| case 2: |
| globalSizeX = 8; |
| globalSizeY = 8; |
| globalSizeZ = 8; |
| break; |
| default: |
| break; |
| } |
| |
| if(--level < 0) |
| { |
| return; |
| } |
| |
| for (size_t tidZ = 0; tidZ < globalSizeZ; ++tidZ) |
| { |
| for (size_t tidY = 0; tidY < globalSizeY; ++tidY) |
| { |
| for (size_t tidX = 0; tidX < globalSizeX; ++tidX) |
| { |
| size_t linearID = globalSizeX * globalSizeY * tidZ + globalSizeX * tidY + tidX; |
| cl_int nextDim = (linearID + level) % 3; |
| ++referenceResults[linearID % len]; |
| if (level >= tidX && level >= tidY && level >= tidZ) |
| { |
| generate_reference_results_all_diff_mix(referenceResults, len, level, nextDim); |
| } |
| } |
| } |
| } |
| } |
| |
| static int check_all_diff_mix(cl_int* results, cl_int len, cl_int nesting_level) |
| { |
| std::vector<cl_int> referenceResults(len, 0); |
| generate_reference_results_all_diff_mix(referenceResults, len, nesting_level, 0); |
| |
| for(size_t i = 0; i < len; ++i) |
| { |
| if (results[i] != referenceResults[i]) |
| { |
| log_error("ERROR: Kernel returned %d vs. expected %d, index: %d\n", results[i], referenceResults[i], i); |
| return (int)i; |
| } |
| } |
| |
| return -1; |
| } |
| |
| static const kernel_src_check sources_enqueue_wg_size[] = |
| { |
| { KERNEL(enqueue_1D_wg_size_single), check_single }, |
| { KERNEL(enqueue_1D_wg_size_some_eq), check_some_eq_1D }, |
| { KERNEL(enqueue_1D_wg_size_some_diff), check_some_diff_1D }, |
| { KERNEL(enqueue_1D_wg_size_all_eq), check_all_eq_1D }, |
| { KERNEL(enqueue_1D_wg_size_all_diff), check_all_diff_1D }, |
| |
| { KERNEL(enqueue_2D_wg_size_single), check_single }, |
| { KERNEL(enqueue_2D_wg_size_some_eq), check_some_eq_2D }, |
| { KERNEL(enqueue_2D_wg_size_some_diff), check_some_diff_2D }, |
| { KERNEL(enqueue_2D_wg_size_all_eq), check_all_eq_2D }, |
| { KERNEL(enqueue_2D_wg_size_all_diff), check_all_diff_2D }, |
| |
| { KERNEL(enqueue_3D_wg_size_single), check_single }, |
| { KERNEL(enqueue_3D_wg_size_some_eq), check_some_eq_3D }, |
| { KERNEL(enqueue_3D_wg_size_some_diff), check_some_diff_3D }, |
| { KERNEL(enqueue_3D_wg_size_all_eq), check_all_eq_3D }, |
| { KERNEL(enqueue_3D_wg_size_all_diff), check_all_diff_3D }, |
| |
| { KERNEL(enqueue_mix_wg_size_single), check_single }, |
| { KERNEL(enqueue_mix_wg_size_some_eq), check_some_eq_mix }, |
| { KERNEL(enqueue_mix_wg_size_some_diff), check_some_diff_mix }, |
| { KERNEL(enqueue_mix_wg_size_all_eq), check_all_eq_mix }, |
| { KERNEL(enqueue_mix_wg_size_all_diff), check_all_diff_mix } |
| }; |
| |
| int test_enqueue_wg_size(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements) |
| { |
| MTdata d; |
| cl_uint i, k; |
| cl_int err_ret, res = 0; |
| clCommandQueueWrapper dev_queue; |
| const cl_int MAX_GLOBAL_WORK_SIZE = MAX_GWS / 4; |
| cl_int kernel_results[MAX_GLOBAL_WORK_SIZE] = { 0 }; |
| cl_uint vrnd[MAX_GLOBAL_WORK_SIZE] = { 0 }; |
| |
| size_t ret_len; |
| cl_uint max_queues = 1; |
| cl_uint maxQueueSize = 0; |
| d = init_genrand(gRandomSeed); |
| |
| if(gWimpyMode) |
| { |
| nestingLevel = 2; |
| vlog( "*** WARNING: Testing in Wimpy mode! ***\n" ); |
| vlog( "*** Wimpy mode is not sufficient to verify correctness. ***\n" ); |
| } |
| |
| 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_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 failCnt = 0; |
| for(k = 0; k < arr_size(sources_enqueue_wg_size); ++k) |
| { |
| if (!gKernelName.empty() && gKernelName != sources_enqueue_wg_size[k].src.kernel_name) |
| continue; |
| |
| log_info("Running '%s' kernel (%d of %d) ...\n", sources_enqueue_wg_size[k].src.kernel_name, k + 1, arr_size(sources_enqueue_wg_size)); |
| for(i = 0; i < MAX_GLOBAL_WORK_SIZE; ++i) |
| { |
| kernel_results[i] = 0; |
| vrnd[i] = genrand_int32(d); |
| } |
| |
| // Fill some elements with prime numbers |
| cl_uint prime[] = { 3, 5, 7, 11, 13, 17, 19, 23, |
| 29, 31, 37, 41, 43, 47, 53, 59, |
| 61, 67, 71, 73, 79, 83, 89, 97, |
| 101, 103, 107, 109, 113, 127 }; |
| |
| for(i = 0; i < arr_size(prime); ++i) |
| { |
| vrnd[genrand_int32(d) % MAX_GLOBAL_WORK_SIZE] = prime[i]; |
| } |
| |
| clMemWrapper mem; |
| mem = clCreateBuffer(context, CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR, sizeof(vrnd), vrnd, &err_ret); |
| test_error(err_ret, "clCreateBuffer() failed"); |
| |
| kernel_arg args[] = |
| { |
| { sizeof(cl_uint), &nestingLevel }, |
| { sizeof(cl_uint), &MAX_GLOBAL_WORK_SIZE }, |
| { sizeof(cl_mem), &mem } |
| }; |
| |
| size_t global_size = MAX_GLOBAL_WORK_SIZE; |
| size_t local_size = (max_local_size > global_size) ? global_size : max_local_size; |
| |
| err_ret = run_n_kernel_args(context, queue, sources_enqueue_wg_size[k].src.lines, sources_enqueue_wg_size[k].src.num_lines, sources_enqueue_wg_size[k].src.kernel_name, local_size, global_size, kernel_results, sizeof(kernel_results), arr_size(args), args); |
| |
| //check results |
| int fail = sources_enqueue_wg_size[k].check(kernel_results, global_size, nestingLevel); |
| |
| if(check_error(err_ret, "'%s' kernel execution failed", sources_enqueue_wg_size[k].src.kernel_name)) { ++failCnt; res = -1; continue; } |
| else if(fail >= 0 && check_error(-1, "'%s' kernel results validation failed: [%d]", sources_enqueue_wg_size[k].src.kernel_name, fail)) { ++failCnt; res = -1; continue; } |
| else log_info("'%s' kernel is OK.\n", sources_enqueue_wg_size[k].src.kernel_name); |
| } |
| |
| if (failCnt > 0) |
| { |
| log_error("ERROR: %d of %d kernels failed.\n", failCnt, arr_size(sources_enqueue_wg_size)); |
| } |
| |
| free_mtdata(d); |
| |
| return res; |
| } |
| |
| #endif |
| |
