From 869d5a1032aab856e2351231c65f4943111423d1 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?K=C3=A9vin=20Petit?= Date: Tue, 8 Sep 2020 18:16:08 +0100 Subject: [PATCH] Merge vec_step and vec_align binaries into a new vectors binary (#926) * Merge vec_step and vec_align binaries into a new vectors binary This change just merges the two suites to reuse common functions and definitions. There is probably room for further code consolidation but probably best done as a separate change. Contributes to #683 Signed-off-by: Kevin Petit * format fixes --- test_conformance/CMakeLists.txt | 3 +- .../opencl_conformance_tests_12_full.csv | 3 +- ...e_tests_12_full_no_math_or_conversions.csv | 3 +- .../opencl_conformance_tests_20_full.csv | 3 +- ...e_tests_20_full_no_math_or_conversions.csv | 3 +- ...opencl_conformance_tests_21_full_spirv.csv | 3 +- ...encl_conformance_tests_21_legacy_wimpy.csv | 3 +- .../opencl_conformance_tests_full.csv | 3 +- ...ance_tests_full_no_math_or_conversions.csv | 3 +- test_conformance/vec_align/globals.cpp | 59 --- test_conformance/vec_align/main.cpp | 41 -- test_conformance/vec_align/structs.h | 73 --- test_conformance/vec_align/type_replacer.h | 23 - test_conformance/vec_step/CMakeLists.txt | 11 - test_conformance/vec_step/defines.h | 41 -- test_conformance/vec_step/globals.cpp | 52 --- test_conformance/vec_step/procs.h | 43 -- test_conformance/vec_step/structs.cpp | 285 ------------ test_conformance/vec_step/structs.h | 67 --- test_conformance/vec_step/testBase.h | 28 -- test_conformance/vec_step/type_replacer.cpp | 115 ----- .../{vec_align => vectors}/CMakeLists.txt | 3 +- .../{vec_align => vectors}/defines.h | 6 +- test_conformance/vectors/globals.cpp | 46 ++ .../{vec_step => vectors}/main.cpp | 20 +- .../{vec_align => vectors}/procs.h | 38 +- .../{vec_align => vectors}/structs.cpp | 208 +++++---- test_conformance/vectors/structs.h | 75 ++++ .../{vec_align => vectors}/testBase.h | 2 +- .../{vec_step => vectors}/test_step.cpp | 130 +++--- .../{vec_align => vectors}/test_vec_align.cpp | 417 +++++++++--------- .../{vec_align => vectors}/type_replacer.cpp | 69 +-- .../{vec_step => vectors}/type_replacer.h | 12 +- 33 files changed, 618 insertions(+), 1273 deletions(-) delete mode 100644 test_conformance/vec_align/globals.cpp delete mode 100644 test_conformance/vec_align/main.cpp delete mode 100644 test_conformance/vec_align/structs.h delete mode 100644 test_conformance/vec_align/type_replacer.h delete mode 100644 test_conformance/vec_step/CMakeLists.txt delete mode 100644 test_conformance/vec_step/defines.h delete mode 100644 test_conformance/vec_step/globals.cpp delete mode 100644 test_conformance/vec_step/procs.h delete mode 100644 test_conformance/vec_step/structs.cpp delete mode 100644 test_conformance/vec_step/structs.h delete mode 100644 test_conformance/vec_step/testBase.h delete mode 100644 test_conformance/vec_step/type_replacer.cpp rename test_conformance/{vec_align => vectors}/CMakeLists.txt (77%) rename test_conformance/{vec_align => vectors}/defines.h (92%) create mode 100644 test_conformance/vectors/globals.cpp rename test_conformance/{vec_step => vectors}/main.cpp (67%) rename test_conformance/{vec_align => vectors}/procs.h (50%) rename test_conformance/{vec_align => vectors}/structs.cpp (57%) create mode 100644 test_conformance/vectors/structs.h rename test_conformance/{vec_align => vectors}/testBase.h (99%) rename test_conformance/{vec_step => vectors}/test_step.cpp (65%) rename test_conformance/{vec_align => vectors}/test_vec_align.cpp (51%) rename test_conformance/{vec_align => vectors}/type_replacer.cpp (62%) rename test_conformance/{vec_step => vectors}/type_replacer.h (62%) diff --git a/test_conformance/CMakeLists.txt b/test_conformance/CMakeLists.txt index b5125bea..83d18ea1 100644 --- a/test_conformance/CMakeLists.txt +++ b/test_conformance/CMakeLists.txt @@ -39,8 +39,7 @@ add_subdirectory( profiling ) add_subdirectory( relationals ) add_subdirectory( select ) add_subdirectory( thread_dimensions ) -add_subdirectory( vec_align ) -add_subdirectory( vec_step ) +add_subdirectory( vectors ) add_subdirectory( c11_atomics ) add_subdirectory( device_execution ) add_subdirectory( non_uniform_work_group ) diff --git a/test_conformance/opencl_conformance_tests_12_full.csv b/test_conformance/opencl_conformance_tests_12_full.csv index 083fcfaa..bb732c40 100644 --- a/test_conformance/opencl_conformance_tests_12_full.csv +++ b/test_conformance/opencl_conformance_tests_12_full.csv @@ -31,8 +31,7 @@ Profiling,profiling/test_profiling Events,events/test_events Allocations (single maximum),allocations/test_allocations single 5 all Allocations (total maximum),allocations/test_allocations multiple 5 all -VecAlign, vec_align/test_vecalign -VecStep, vec_step/test_vecstep +Vectors, vectors/test_vectors Printf,printf/test_printf Device Partitioning,device_partition/test_device_partition diff --git a/test_conformance/opencl_conformance_tests_12_full_no_math_or_conversions.csv b/test_conformance/opencl_conformance_tests_12_full_no_math_or_conversions.csv index 2f8d6531..fca9af4b 100644 --- a/test_conformance/opencl_conformance_tests_12_full_no_math_or_conversions.csv +++ b/test_conformance/opencl_conformance_tests_12_full_no_math_or_conversions.csv @@ -31,8 +31,7 @@ Profiling,profiling/test_profiling Events,events/test_events Allocations (single maximum),allocations/test_allocations single 5 all Allocations (total maximum),allocations/test_allocations multiple 5 all -VecAlign, vec_align/test_vecalign -VecStep, vec_step/test_vecstep +Vectors, vectors/test_vectors # ######################################### # Buffers and images diff --git a/test_conformance/opencl_conformance_tests_20_full.csv b/test_conformance/opencl_conformance_tests_20_full.csv index e1c0ecdd..e5320bb4 100644 --- a/test_conformance/opencl_conformance_tests_20_full.csv +++ b/test_conformance/opencl_conformance_tests_20_full.csv @@ -31,8 +31,7 @@ Profiling,profiling/test_profiling Events,events/test_events Allocations (single maximum),allocations/test_allocations single 5 all Allocations (total maximum),allocations/test_allocations multiple 5 all -VecAlign, vec_align/test_vecalign -VecStep, vec_step/test_vecstep +Vectors, vectors/test_vectors Printf,printf/test_printf Device Partitioning,device_partition/test_device_partition diff --git a/test_conformance/opencl_conformance_tests_20_full_no_math_or_conversions.csv b/test_conformance/opencl_conformance_tests_20_full_no_math_or_conversions.csv index 2f8d6531..fca9af4b 100644 --- a/test_conformance/opencl_conformance_tests_20_full_no_math_or_conversions.csv +++ b/test_conformance/opencl_conformance_tests_20_full_no_math_or_conversions.csv @@ -31,8 +31,7 @@ Profiling,profiling/test_profiling Events,events/test_events Allocations (single maximum),allocations/test_allocations single 5 all Allocations (total maximum),allocations/test_allocations multiple 5 all -VecAlign, vec_align/test_vecalign -VecStep, vec_step/test_vecstep +Vectors, vectors/test_vectors # ######################################### # Buffers and images diff --git a/test_conformance/opencl_conformance_tests_21_full_spirv.csv b/test_conformance/opencl_conformance_tests_21_full_spirv.csv index 4b84596c..1c2b7499 100644 --- a/test_conformance/opencl_conformance_tests_21_full_spirv.csv +++ b/test_conformance/opencl_conformance_tests_21_full_spirv.csv @@ -31,8 +31,7 @@ Profiling,profiling/test_profiling --compilation-mode spir-v --compilation-cache Events,events/test_events --compilation-mode spir-v --compilation-cache-path . Allocations (single maximum),allocations/test_allocations single 5 all --compilation-mode spir-v --compilation-cache-path . Allocations (total maximum),allocations/test_allocations multiple 5 all --compilation-mode spir-v --compilation-cache-path . -VecAlign, vec_align/test_vecalign --compilation-mode spir-v --compilation-cache-path . -VecStep, vec_step/test_vecstep --compilation-mode spir-v --compilation-cache-path . +Vectors, vectors/test_vectors --compilation-mode spir-v --compilation-cache-path . Printf,printf/test_printf --compilation-mode spir-v --compilation-cache-path . Device Partitioning,device_partition/test_device_partition --compilation-mode spir-v --compilation-cache-path . diff --git a/test_conformance/opencl_conformance_tests_21_legacy_wimpy.csv b/test_conformance/opencl_conformance_tests_21_legacy_wimpy.csv index bbde1834..610a2945 100644 --- a/test_conformance/opencl_conformance_tests_21_legacy_wimpy.csv +++ b/test_conformance/opencl_conformance_tests_21_legacy_wimpy.csv @@ -33,8 +33,7 @@ Profiling,profiling/test_profiling Events,events/test_events Allocations (single maximum),allocations/test_allocations single 5 all Allocations (total maximum),allocations/test_allocations multiple 5 all -VecAlign, vec_align/test_vecalign -VecStep, vec_step/test_vecstep +Vectors, vectors/test_vectors Printf,printf/test_printf Device Partitioning,device_partition/test_device_partition diff --git a/test_conformance/opencl_conformance_tests_full.csv b/test_conformance/opencl_conformance_tests_full.csv index 05555279..efd004e6 100644 --- a/test_conformance/opencl_conformance_tests_full.csv +++ b/test_conformance/opencl_conformance_tests_full.csv @@ -31,8 +31,7 @@ Profiling,profiling/test_profiling Events,events/test_events Allocations (single maximum),allocations/test_allocations single 5 all Allocations (total maximum),allocations/test_allocations multiple 5 all -VecAlign, vec_align/test_vecalign -VecStep, vec_step/test_vecstep +Vectors, vectors/test_vectors Printf,printf/test_printf Device Partitioning,device_partition/test_device_partition diff --git a/test_conformance/opencl_conformance_tests_full_no_math_or_conversions.csv b/test_conformance/opencl_conformance_tests_full_no_math_or_conversions.csv index ecc1314c..f14e0991 100644 --- a/test_conformance/opencl_conformance_tests_full_no_math_or_conversions.csv +++ b/test_conformance/opencl_conformance_tests_full_no_math_or_conversions.csv @@ -31,8 +31,7 @@ Profiling,profiling/test_profiling Events,events/test_events Allocations (single maximum),allocations/test_allocations single 5 all Allocations (total maximum),allocations/test_allocations multiple 5 all -VecAlign, vec_align/test_vecalign -VecStep, vec_step/test_vecstep +Vectors, vectors/test_vectors Printf,printf/test_printf Device Partitioning,device_partition/test_device_partition diff --git a/test_conformance/vec_align/globals.cpp b/test_conformance/vec_align/globals.cpp deleted file mode 100644 index 3deecadf..00000000 --- a/test_conformance/vec_align/globals.cpp +++ /dev/null @@ -1,59 +0,0 @@ -// -// 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 "defines.h" - - -// 1,2,3,4,8,16 or -// 1,2,4,8,16,3 -int g_arrVecSizes[NUM_VECTOR_SIZES] = {1,2,3,4,8,16}; -int g_arrVecSteps[NUM_VECTOR_SIZES] = {1,2,4,4,8,16}; -const char * g_arrVecSizeNames[NUM_VECTOR_SIZES] = {"", "2","3","4","8","16"}; -size_t g_arrVecAlignMasks[NUM_VECTOR_SIZES] = {(size_t)0, - (size_t)0x1, // 2 - (size_t)0x3, // 3 - (size_t)0x3, // 4 - (size_t)0x7, // 8 - (size_t)0xf // 16 -}; - -bool g_wimpyMode = false; - -ExplicitType types[] = { kChar, kUChar, - kShort, kUShort, - kInt, kUInt, - kLong, kULong, - kFloat, kDouble, - kNumExplicitTypes }; - - -const char *g_arrTypeNames[] = - { - "char", "uchar", - "short", "ushort", - "int", "uint", - "long", "ulong", - "float", "double" - }; - -extern const size_t g_arrTypeSizes[] = - { - 1, 1, - 2, 2, - 4, 4, - 8, 8, - 4, 8 - }; - diff --git a/test_conformance/vec_align/main.cpp b/test_conformance/vec_align/main.cpp deleted file mode 100644 index 6894895a..00000000 --- a/test_conformance/vec_align/main.cpp +++ /dev/null @@ -1,41 +0,0 @@ -// -// 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 "harness/compat.h" - -#include -#include -#include "procs.h" -#include "harness/testHarness.h" - -#if !defined(_WIN32) -#include -#endif - -test_definition test_list[] = { - ADD_TEST( vec_align_array ), - ADD_TEST( vec_align_struct ), - ADD_TEST( vec_align_packed_struct ), - ADD_TEST( vec_align_struct_arr ), - ADD_TEST( vec_align_packed_struct_arr ), -}; - -const int test_num = ARRAY_SIZE( test_list ); - -int main(int argc, const char *argv[]) -{ - return runTestHarness( argc, argv, test_num, test_list, false, false, 0 ); -} - diff --git a/test_conformance/vec_align/structs.h b/test_conformance/vec_align/structs.h deleted file mode 100644 index e26b8105..00000000 --- a/test_conformance/vec_align/structs.h +++ /dev/null @@ -1,73 +0,0 @@ -// -// 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 "testBase.h" - - -#include "harness/conversions.h" -#include "harness/typeWrappers.h" - -typedef struct _clState -{ - cl_device_id m_device; - cl_context m_context; - cl_command_queue m_queue; - - cl_program m_program; - cl_kernel m_kernel; - size_t m_numThreads; -} clState; - -clState * newClState(cl_device_id device, cl_context context, cl_command_queue queue); -clState * destroyClState(clState * pState); - -int clStateMakeProgram(clState * pState, const char * prog, - const char * kernelName); -void clStateDestroyProgramAndKernel(clState * pState); - -int runKernel(clState * pState, size_t numThreads); - -typedef struct _bufferStruct -{ - void * m_pIn; - void * m_pOut; - - cl_mem m_outBuffer; - cl_mem m_inBuffer; - - size_t m_bufSizeIn, m_bufSizeOut; - - int m_bufferUploaded; -} bufferStruct; - - -bufferStruct * newBufferStruct(size_t inSize, size_t outSize, clState * pClState); - -bufferStruct * destroyBufferStruct(bufferStruct * destroyMe, clState * pClState); - -void initContents(bufferStruct * pBufferStruct, clState * pClState, - size_t typeSize, - size_t vecWidth); - -int pushArgs(bufferStruct * pBufferStruct, clState * pClState); -int retrieveResults(bufferStruct * pBufferStruct, clState * pClState); - -// vecSizeIdx indexes into g_arrVecAlignMasks, g_arrVecSizeNames -// and g_arrVecSizes -int checkCorrectness(bufferStruct * pBufferStruct, clState * pClState, - size_t minAlign); - -int checkPackedCorrectness(bufferStruct * pBufferStruct, clState * pClState, - size_t totSize, size_t beforeSize); diff --git a/test_conformance/vec_align/type_replacer.h b/test_conformance/vec_align/type_replacer.h deleted file mode 100644 index f50b08d7..00000000 --- a/test_conformance/vec_align/type_replacer.h +++ /dev/null @@ -1,23 +0,0 @@ -// -// 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 - -size_t doReplace(char * dest, size_t destLength, const char * source, - const char * stringToReplace1, const char * replaceWith1, - const char * stringToReplace2, const char * replaceWith2); - -size_t doSingleReplace(char * dest, size_t destLength, const char * source, - const char * stringToReplace, const char * replaceWith); diff --git a/test_conformance/vec_step/CMakeLists.txt b/test_conformance/vec_step/CMakeLists.txt deleted file mode 100644 index c18fc1bc..00000000 --- a/test_conformance/vec_step/CMakeLists.txt +++ /dev/null @@ -1,11 +0,0 @@ -set(MODULE_NAME VECSTEP) - -set(${MODULE_NAME}_SOURCES - globals.cpp - test_step.cpp - main.cpp - structs.cpp - type_replacer.cpp -) - -include(../CMakeCommon.txt) diff --git a/test_conformance/vec_step/defines.h b/test_conformance/vec_step/defines.h deleted file mode 100644 index 5f364e41..00000000 --- a/test_conformance/vec_step/defines.h +++ /dev/null @@ -1,41 +0,0 @@ -// -// 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 "harness/errorHelpers.h" -#include "harness/kernelHelpers.h" -#include "harness/threadTesting.h" -#include "harness/typeWrappers.h" -#include "harness/conversions.h" -#include "harness/mt19937.h" - - -// 1,2,3,4,8,16 or -// 1,2,4,8,16,3 -#define NUM_VECTOR_SIZES 6 - -extern int g_arrVecSizes[NUM_VECTOR_SIZES]; -extern int g_arrVecSteps[NUM_VECTOR_SIZES]; -extern bool g_wimpyMode; - -extern const char * g_arrVecSizeNames[NUM_VECTOR_SIZES]; - -// Define the buffer size that we want to block our test with -#define BUFFER_SIZE (1024*1024) -#define KPAGESIZE 4096 - -extern ExplicitType types[]; - -extern const char *g_arrTypeNames[]; -extern const size_t g_arrTypeSizes[]; diff --git a/test_conformance/vec_step/globals.cpp b/test_conformance/vec_step/globals.cpp deleted file mode 100644 index 029ae1c6..00000000 --- a/test_conformance/vec_step/globals.cpp +++ /dev/null @@ -1,52 +0,0 @@ -// -// 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 "defines.h" - - -// 1,2,3,4,8,16 or -// 1,2,4,8,16,3 -int g_arrVecSizes[NUM_VECTOR_SIZES] = {1,2,3,4,8,16}; -int g_arrVecSteps[NUM_VECTOR_SIZES] = {1,2,4,4,8,16}; -const char * g_arrVecSizeNames[NUM_VECTOR_SIZES] = {"", "2","3","4","8","16"}; - -bool g_wimpyMode = false; - -ExplicitType types[] = { kChar, kUChar, - kShort, kUShort, - kInt, kUInt, - kLong, kULong, - kFloat, kDouble, - kNumExplicitTypes }; - - -const char *g_arrTypeNames[] = - { - "char", "uchar", - "short", "ushort", - "int", "uint", - "long", "ulong", - "float", "double" - }; - -extern const size_t g_arrTypeSizes[] = - { - 1, 1, - 2, 2, - 4, 4, - 8, 8, - 4, 8 - }; - diff --git a/test_conformance/vec_step/procs.h b/test_conformance/vec_step/procs.h deleted file mode 100644 index 382a36b7..00000000 --- a/test_conformance/vec_step/procs.h +++ /dev/null @@ -1,43 +0,0 @@ -// -// 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 "harness/errorHelpers.h" -#include "harness/kernelHelpers.h" -#include "harness/threadTesting.h" -#include "harness/typeWrappers.h" -#include "harness/conversions.h" -#include "harness/mt19937.h" - -// The number of errors to print out for each test in the shuffle tests -#define MAX_ERRORS_TO_PRINT 1 - - -extern int create_program_and_kernel(const char *source, const char *kernel_name, cl_program *program_ret, cl_kernel *kernel_ret); - - -/* - test_step_type, - test_step_var, - test_step_typedef_type, - test_step_typedef_var, -*/ - -extern int test_step_type(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); - -extern int test_step_var(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); - -extern int test_step_typedef_type(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); - -extern int test_step_typedef_var(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); diff --git a/test_conformance/vec_step/structs.cpp b/test_conformance/vec_step/structs.cpp deleted file mode 100644 index b36e892f..00000000 --- a/test_conformance/vec_step/structs.cpp +++ /dev/null @@ -1,285 +0,0 @@ -// -// 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 "structs.h" - - -#include "defines.h" - -/** typedef struct _bufferStruct - { - void * m_pIn; - void * m_pOut; - - cl_mem m_outBuffer; - cl_mem m_inBuffer; - - size_t m_bufSize; - } bufferStruct; - */ - - -clState * newClState(cl_device_id device, cl_context context, cl_command_queue queue) -{ - clState * pResult = (clState *)malloc(sizeof(clState)); - - pResult->m_device = device; - pResult->m_context = context; - pResult->m_queue = queue; - - pResult->m_kernel = NULL; pResult->m_program = NULL; - return pResult; -} - -clState * destroyClState(clState * pState) -{ - clStateDestroyProgramAndKernel(pState); - free(pState); - return NULL; -} - - -int clStateMakeProgram(clState * pState, const char * prog, - const char * kernelName) -{ - const char * srcArr[1] = {NULL}; - srcArr[0] = prog; - int err = create_single_kernel_helper(pState->m_context, - &(pState->m_program), - &(pState->m_kernel), - 1, srcArr, kernelName ); - return err; -} - -int runKernel(clState * pState, size_t numThreads) { - int err; - pState->m_numThreads = numThreads; - err = clEnqueueNDRangeKernel(pState->m_queue, pState->m_kernel, - 1, NULL, &(pState->m_numThreads), - NULL, 0, NULL, NULL); - if(err != CL_SUCCESS) - { - log_error("clEnqueueNDRangeKernel returned %d (%x)\n", - err, err); - return -1; - } - return 0; -} - - -void clStateDestroyProgramAndKernel(clState * pState) -{ - if(pState->m_kernel != NULL) { - clReleaseKernel( pState->m_kernel ); - pState->m_kernel = NULL; - } - if(pState->m_program != NULL) { - clReleaseProgram( pState->m_program ); - pState->m_program = NULL; - } -} - -bufferStruct * newBufferStruct(size_t inSize, size_t outSize, clState * pClState) { - int error; - bufferStruct * pResult = (bufferStruct *)malloc(sizeof(bufferStruct)); - - pResult->m_bufSizeIn = inSize; - pResult->m_bufSizeOut = outSize; - - pResult->m_pIn = malloc(inSize); - pResult->m_pOut = malloc(outSize); - - pResult->m_inBuffer = clCreateBuffer(pClState->m_context, CL_MEM_READ_ONLY, - inSize, NULL, &error); - if( pResult->m_inBuffer == NULL ) - { - vlog_error( "clCreateArray failed for input (%d)\n", error ); - return destroyBufferStruct(pResult, pClState); - } - - pResult->m_outBuffer = clCreateBuffer( pClState->m_context, - CL_MEM_WRITE_ONLY, - outSize, - NULL, - &error ); - if( pResult->m_outBuffer == NULL ) - { - vlog_error( "clCreateArray failed for output (%d)\n", error ); - return destroyBufferStruct(pResult, pClState); - } - - return pResult; -} - -bufferStruct * destroyBufferStruct(bufferStruct * destroyMe, clState * pClState) { - if(destroyMe) - { - if(destroyMe->m_outBuffer != NULL) { - clReleaseMemObject(destroyMe->m_outBuffer); - destroyMe->m_outBuffer = NULL; - } - if(destroyMe->m_inBuffer != NULL) { - clReleaseMemObject(destroyMe->m_inBuffer); - destroyMe->m_inBuffer = NULL; - } - if(destroyMe->m_pIn != NULL) { - free(destroyMe->m_pIn); - destroyMe->m_pIn = NULL; - } - if(destroyMe->m_pOut != NULL) { - free(destroyMe->m_pOut); - destroyMe->m_pOut = NULL; - } - - free((void *)destroyMe); - destroyMe = NULL; - } - return destroyMe; -} - -void initContents(bufferStruct * pBufferStruct, clState * pClState, - size_t typeSize, - size_t countIn, size_t countOut ) -{ - size_t i; - - uint64_t start = 0; - - switch(typeSize) - { - case 1: { - uint8_t* ub = (uint8_t *)(pBufferStruct->m_pIn); - for (i=0; i < countIn; ++i) - { - ub[i] = (uint8_t)start++; - } - break; - } - case 2: { - uint16_t* us = (uint16_t *)(pBufferStruct->m_pIn); - for (i=0; i < countIn; ++i) - { - us[i] = (uint16_t)start++; - } - break; - } - case 4: { - if (!g_wimpyMode) { - uint32_t* ui = (uint32_t *)(pBufferStruct->m_pIn); - for (i=0; i < countIn; ++i) { - ui[i] = (uint32_t)start++; - } - } - else { - // The short test doesn't iterate over the entire 32 bit space so - // we alternate between positive and negative values - int32_t* ui = (int32_t *)(pBufferStruct->m_pIn); - int32_t sign = 1; - for (i=0; i < countIn; ++i, ++start) { - ui[i] = (int32_t)start*sign; - sign = sign * -1; - } - } - break; - } - case 8: { - // We don't iterate over the entire space of 64 bit so for the - // selects, we want to test positive and negative values - int64_t* ll = (int64_t *)(pBufferStruct->m_pIn); - int64_t sign = 1; - for (i=0; i < countIn; ++i, ++start) { - ll[i] = start*sign; - sign = sign * -1; - } - break; - } - default: { - log_error("invalid type size %x\n", (int)typeSize); - } - } - // pBufferStruct->m_bufSizeIn - // pBufferStruct->m_bufSizeOut -} - -int pushArgs(bufferStruct * pBufferStruct, clState * pClState) -{ - int err; - err = clEnqueueWriteBuffer(pClState->m_queue, pBufferStruct->m_inBuffer, - CL_TRUE, 0, pBufferStruct->m_bufSizeIn, - pBufferStruct->m_pIn, 0, NULL, NULL); - if(err != CL_SUCCESS) - { - log_error("clEnqueueWriteBuffer failed\n"); - return -1; - } - - err = clSetKernelArg(pClState->m_kernel, 0, - sizeof(pBufferStruct->m_inBuffer), // pBufferStruct->m_bufSizeIn, - &(pBufferStruct->m_inBuffer)); - if(err != CL_SUCCESS) - { - log_error("clSetKernelArgs failed, first arg (0)\n"); - return -1; - } - - err = clSetKernelArg(pClState->m_kernel, 1, - sizeof(pBufferStruct->m_outBuffer), // pBufferStruct->m_bufSizeOut, - &(pBufferStruct->m_outBuffer)); - if(err != CL_SUCCESS) - { - log_error("clSetKernelArgs failed, second arg (1)\n"); - return -1; - } - - return 0; -} - -int retrieveResults(bufferStruct * pBufferStruct, clState * pClState) -{ - int err; - err = clEnqueueReadBuffer(pClState->m_queue, pBufferStruct->m_outBuffer, - CL_TRUE, 0, pBufferStruct->m_bufSizeOut, - pBufferStruct->m_pOut, 0, NULL, NULL); - if(err != CL_SUCCESS) - { - log_error("clEnqueueReadBuffer failed\n"); - return -1; - } - return 0; -} - -int checkCorrectness(bufferStruct * pBufferStruct, clState * pClState, - size_t typeSize, - size_t vecWidth) -{ - size_t i; - cl_int targetSize = (cl_int) vecWidth; - cl_int * targetArr = (cl_int *)(pBufferStruct->m_pOut); - if(targetSize == 3) - { - targetSize = 4; // hack for 4-aligned vec3 types - } - for(i = 0; i < pClState->m_numThreads; ++i) - { - if(targetArr[i] != targetSize) - { - vlog_error("Error %ld (of %ld). Expected %d, got %d\n", - i, pClState->m_numThreads, - targetSize, targetArr[i]); - return -1; - } - } - return 0; -} diff --git a/test_conformance/vec_step/structs.h b/test_conformance/vec_step/structs.h deleted file mode 100644 index 37e5524d..00000000 --- a/test_conformance/vec_step/structs.h +++ /dev/null @@ -1,67 +0,0 @@ -// -// 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 "testBase.h" - - -#include "harness/conversions.h" -#include "harness/typeWrappers.h" - -typedef struct _clState -{ - cl_device_id m_device; - cl_context m_context; - cl_command_queue m_queue; - - cl_program m_program; - cl_kernel m_kernel; - size_t m_numThreads; -} clState; - -clState * newClState(cl_device_id device, cl_context context, cl_command_queue queue); -clState * destroyClState(clState * pState); - -int clStateMakeProgram(clState * pState, const char * prog, - const char * kernelName); -void clStateDestroyProgramAndKernel(clState * pState); - -int runKernel(clState * pState, size_t numThreads); - -typedef struct _bufferStruct -{ - void * m_pIn; - void * m_pOut; - - cl_mem m_outBuffer; - cl_mem m_inBuffer; - - size_t m_bufSizeIn, m_bufSizeOut; -} bufferStruct; - - -bufferStruct * newBufferStruct(size_t inSize, size_t outSize, clState * pClState); - -bufferStruct * destroyBufferStruct(bufferStruct * destroyMe, clState * pClState); - -void initContents(bufferStruct * pBufferStruct, clState * pClState, - size_t typeSize, - size_t vecWidth); - -int pushArgs(bufferStruct * pBufferStruct, clState * pClState); -int retrieveResults(bufferStruct * pBufferStruct, clState * pClState); - -int checkCorrectness(bufferStruct * pBufferStruct, clState * pClState, - size_t typeSize, - size_t vecWidth); diff --git a/test_conformance/vec_step/testBase.h b/test_conformance/vec_step/testBase.h deleted file mode 100644 index bd72e841..00000000 --- a/test_conformance/vec_step/testBase.h +++ /dev/null @@ -1,28 +0,0 @@ -// -// 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. -// -#ifndef _testBase_h -#define _testBase_h - -#include "harness/compat.h" - -#include -#include -#include -#include - -#include "procs.h" - -#endif // _testBase_h diff --git a/test_conformance/vec_step/type_replacer.cpp b/test_conformance/vec_step/type_replacer.cpp deleted file mode 100644 index 74967b2c..00000000 --- a/test_conformance/vec_step/type_replacer.cpp +++ /dev/null @@ -1,115 +0,0 @@ -// -// 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 -#if !defined(_MSC_VER) -#include -#endif // !_MSC_VER - -size_t doReplace(char * dest, size_t destLength, const char * source, - const char * stringToReplace1, const char * replaceWith1, - const char * stringToReplace2, const char * replaceWith2) -{ - size_t copyCount = 0; - const char * sourcePtr = source; - char * destPtr = dest; - const char * ptr1; - const char * ptr2; - size_t nJump; - size_t len1, len2; - size_t lenReplace1, lenReplace2; - len1 = strlen(stringToReplace1); - len2 = strlen(stringToReplace2); - lenReplace1 = strlen(replaceWith1); - lenReplace2 = strlen(replaceWith2); - for(;copyCount < destLength && *sourcePtr; ) - { - ptr1 = strstr(sourcePtr, stringToReplace1); - ptr2 = strstr(sourcePtr, stringToReplace2); - if(ptr1 != NULL && (ptr2 == NULL || ptr2 > ptr1)) - { - nJump = ptr1-sourcePtr; - if(((uintptr_t)ptr1-(uintptr_t)sourcePtr) > destLength-copyCount) { return -1; } - copyCount += nJump; - strncpy(destPtr, sourcePtr, nJump); - destPtr += nJump; - sourcePtr += nJump + len1; - strcpy(destPtr, replaceWith1); - destPtr += lenReplace1; - } - else if(ptr2 != NULL && (ptr1 == NULL || ptr1 >= ptr2)) - { - nJump = ptr2-sourcePtr; - if(nJump > destLength-copyCount) { return -2; } - copyCount += nJump; - strncpy(destPtr, sourcePtr, nJump); - destPtr += nJump; - sourcePtr += nJump + len2; - strcpy(destPtr, replaceWith2); - destPtr += lenReplace2; - } - else - { - nJump = strlen(sourcePtr); - if(nJump > destLength-copyCount) { return -3; } - copyCount += nJump; - strcpy(destPtr, sourcePtr); - destPtr += nJump; - sourcePtr += nJump; - } - } - *destPtr = '\0'; - return copyCount; -} - -size_t doSingleReplace(char * dest, size_t destLength, const char * source, - const char * stringToReplace, const char * replaceWith) -{ - size_t copyCount = 0; - const char * sourcePtr = source; - char * destPtr = dest; - const char * ptr; - size_t nJump; - size_t len; - size_t lenReplace; - len = strlen(stringToReplace); - lenReplace = strlen(replaceWith); - for(;copyCount < destLength && *sourcePtr; ) - { - ptr = strstr(sourcePtr, stringToReplace); - if(ptr != NULL) - { - nJump = ptr-sourcePtr; - if(((uintptr_t)ptr-(uintptr_t)sourcePtr) > destLength-copyCount) { return -1; } - copyCount += nJump; - strncpy(destPtr, sourcePtr, nJump); - destPtr += nJump; - sourcePtr += nJump + len; - strcpy(destPtr, replaceWith); - destPtr += lenReplace; - } - else - { - nJump = strlen(sourcePtr); - if(nJump > destLength-copyCount) { return -3; } - copyCount += nJump; - strcpy(destPtr, sourcePtr); - destPtr += nJump; - sourcePtr += nJump; - } - } - *destPtr = '\0'; - return copyCount; -} diff --git a/test_conformance/vec_align/CMakeLists.txt b/test_conformance/vectors/CMakeLists.txt similarity index 77% rename from test_conformance/vec_align/CMakeLists.txt rename to test_conformance/vectors/CMakeLists.txt index f1a42e66..278de203 100644 --- a/test_conformance/vec_align/CMakeLists.txt +++ b/test_conformance/vectors/CMakeLists.txt @@ -1,9 +1,10 @@ -set(MODULE_NAME VECALIGN) +set(MODULE_NAME VECTORS) set(${MODULE_NAME}_SOURCES globals.cpp main.cpp structs.cpp + test_step.cpp test_vec_align.cpp type_replacer.cpp ) diff --git a/test_conformance/vec_align/defines.h b/test_conformance/vectors/defines.h similarity index 92% rename from test_conformance/vec_align/defines.h rename to test_conformance/vectors/defines.h index f2bf9e78..c96c3dad 100644 --- a/test_conformance/vec_align/defines.h +++ b/test_conformance/vectors/defines.h @@ -1,6 +1,6 @@ // // 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 @@ -29,11 +29,11 @@ extern int g_arrVecSizes[NUM_VECTOR_SIZES]; extern int g_arrVecSteps[NUM_VECTOR_SIZES]; extern bool g_wimpyMode; -extern const char * g_arrVecSizeNames[NUM_VECTOR_SIZES]; +extern const char *g_arrVecSizeNames[NUM_VECTOR_SIZES]; extern size_t g_arrVecAlignMasks[NUM_VECTOR_SIZES]; // Define the buffer size that we want to block our test with -#define BUFFER_SIZE (1024*1024) +#define BUFFER_SIZE (1024 * 1024) #define KPAGESIZE 4096 extern ExplicitType types[]; diff --git a/test_conformance/vectors/globals.cpp b/test_conformance/vectors/globals.cpp new file mode 100644 index 00000000..6dee6d96 --- /dev/null +++ b/test_conformance/vectors/globals.cpp @@ -0,0 +1,46 @@ +// +// 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 "defines.h" + + +// 1,2,3,4,8,16 or +// 1,2,4,8,16,3 +int g_arrVecSizes[NUM_VECTOR_SIZES] = { 1, 2, 3, 4, 8, 16 }; +int g_arrVecSteps[NUM_VECTOR_SIZES] = { 1, 2, 4, 4, 8, 16 }; +const char *g_arrVecSizeNames[NUM_VECTOR_SIZES] = { + "", "2", "3", "4", "8", "16" +}; +size_t g_arrVecAlignMasks[NUM_VECTOR_SIZES] = { + (size_t)0, + (size_t)0x1, // 2 + (size_t)0x3, // 3 + (size_t)0x3, // 4 + (size_t)0x7, // 8 + (size_t)0xf // 16 +}; + +bool g_wimpyMode = false; + +ExplicitType types[] = { + kChar, kUChar, kShort, kUShort, kInt, kUInt, kLong, + kULong, kFloat, kDouble, kNumExplicitTypes +}; + + +const char *g_arrTypeNames[] = { "char", "uchar", "short", "ushort", "int", + "uint", "long", "ulong", "float", "double" }; + +extern const size_t g_arrTypeSizes[] = { 1, 1, 2, 2, 4, 4, 8, 8, 4, 8 }; diff --git a/test_conformance/vec_step/main.cpp b/test_conformance/vectors/main.cpp similarity index 67% rename from test_conformance/vec_step/main.cpp rename to test_conformance/vectors/main.cpp index 9742a28f..98f278c3 100644 --- a/test_conformance/vec_step/main.cpp +++ b/test_conformance/vectors/main.cpp @@ -1,6 +1,6 @@ // // 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 @@ -25,16 +25,20 @@ #endif test_definition test_list[] = { - ADD_TEST( step_type ), - ADD_TEST( step_var ), - ADD_TEST( step_typedef_type ), - ADD_TEST( step_typedef_var ), + ADD_TEST(step_type), + ADD_TEST(step_var), + ADD_TEST(step_typedef_type), + ADD_TEST(step_typedef_var), + ADD_TEST(vec_align_array), + ADD_TEST(vec_align_struct), + ADD_TEST(vec_align_packed_struct), + ADD_TEST(vec_align_struct_arr), + ADD_TEST(vec_align_packed_struct_arr), }; -const int test_num = ARRAY_SIZE( test_list ); +const int test_num = ARRAY_SIZE(test_list); int main(int argc, const char *argv[]) { - return runTestHarness( argc, argv, test_num, test_list, false, false, 0 ); + return runTestHarness(argc, argv, test_num, test_list, false, false, 0); } - diff --git a/test_conformance/vec_align/procs.h b/test_conformance/vectors/procs.h similarity index 50% rename from test_conformance/vec_align/procs.h rename to test_conformance/vectors/procs.h index 7b60e1eb..db423a6a 100644 --- a/test_conformance/vec_align/procs.h +++ b/test_conformance/vectors/procs.h @@ -1,6 +1,6 @@ // // 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 @@ -24,16 +24,32 @@ #define MAX_ERRORS_TO_PRINT 1 -extern int create_program_and_kernel(const char *source, const char *kernel_name, cl_program *program_ret, cl_kernel *kernel_ret); +extern int create_program_and_kernel(const char *source, + const char *kernel_name, + cl_program *program_ret, + cl_kernel *kernel_ret); + +extern int test_step_type(cl_device_id deviceID, cl_context context, + cl_command_queue queue, int num_elements); +extern int test_step_var(cl_device_id deviceID, cl_context context, + cl_command_queue queue, int num_elements); +extern int test_step_typedef_type(cl_device_id deviceID, cl_context context, + cl_command_queue queue, int num_elements); +extern int test_step_typedef_var(cl_device_id deviceID, cl_context context, + cl_command_queue queue, int num_elements); + +int test_vec_align_array(cl_device_id deviceID, cl_context context, + cl_command_queue queue, int num_elements); + +int test_vec_align_struct(cl_device_id deviceID, cl_context context, + cl_command_queue queue, int num_elements); + +int test_vec_align_packed_struct(cl_device_id deviceID, cl_context context, + cl_command_queue queue, int num_elements); -int test_vec_align_array(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); +int test_vec_align_struct_arr(cl_device_id deviceID, cl_context context, + cl_command_queue queue, int num_elements); -int test_vec_align_struct(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); - -int test_vec_align_packed_struct(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); - - -int test_vec_align_struct_arr(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); - -int test_vec_align_packed_struct_arr(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); +int test_vec_align_packed_struct_arr(cl_device_id deviceID, cl_context context, + cl_command_queue queue, int num_elements); diff --git a/test_conformance/vec_align/structs.cpp b/test_conformance/vectors/structs.cpp similarity index 57% rename from test_conformance/vec_align/structs.cpp rename to test_conformance/vectors/structs.cpp index 2e15e36f..9bfa389b 100644 --- a/test_conformance/vec_align/structs.cpp +++ b/test_conformance/vectors/structs.cpp @@ -1,6 +1,6 @@ // // 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 @@ -33,9 +33,10 @@ */ -clState * newClState(cl_device_id device, cl_context context, cl_command_queue queue) +clState *newClState(cl_device_id device, cl_context context, + cl_command_queue queue) { - clState * pResult = (clState *)malloc(sizeof(clState)); + clState *pResult = (clState *)malloc(sizeof(clState)); #if DEBUG_MEM_ALLOC log_info("malloc clState * %x\n", pResult); #endif @@ -44,11 +45,12 @@ clState * newClState(cl_device_id device, cl_context context, cl_command_queue q pResult->m_context = context; pResult->m_queue = queue; - pResult->m_kernel = NULL; pResult->m_program = NULL; + pResult->m_kernel = NULL; + pResult->m_program = NULL; return pResult; } -clState * destroyClState(clState * pState) +clState *destroyClState(clState *pState) { clStateDestroyProgramAndKernel(pState); #if DEBUG_MEM_ALLOC @@ -59,55 +61,56 @@ clState * destroyClState(clState * pState) } -int clStateMakeProgram(clState * pState, const char * prog, - const char * kernelName) +int clStateMakeProgram(clState *pState, const char *prog, + const char *kernelName) { - const char * srcArr[1] = {NULL}; + const char *srcArr[1] = { NULL }; srcArr[0] = prog; - int err = create_single_kernel_helper(pState->m_context, - &(pState->m_program), - &(pState->m_kernel), - 1, srcArr, kernelName ); + int err = + create_single_kernel_helper(pState->m_context, &(pState->m_program), + &(pState->m_kernel), 1, srcArr, kernelName); #if DEBUG_MEM_ALLOC log_info("create program and kernel\n"); #endif return err; } -int runKernel(clState * pState, size_t numThreads) { +int runKernel(clState *pState, size_t numThreads) +{ int err; pState->m_numThreads = numThreads; - err = clEnqueueNDRangeKernel(pState->m_queue, pState->m_kernel, - 1, NULL, &(pState->m_numThreads), - NULL, 0, NULL, NULL); - if(err != CL_SUCCESS) + err = clEnqueueNDRangeKernel(pState->m_queue, pState->m_kernel, 1, NULL, + &(pState->m_numThreads), NULL, 0, NULL, NULL); + if (err != CL_SUCCESS) { - log_error("clEnqueueNDRangeKernel returned %d (%x)\n", - err, err); + log_error("clEnqueueNDRangeKernel returned %d (%x)\n", err, err); return -1; } return 0; } -void clStateDestroyProgramAndKernel(clState * pState) +void clStateDestroyProgramAndKernel(clState *pState) { #if DEBUG_MEM_ALLOC log_info("destroy program and kernel\n"); #endif - if(pState->m_kernel != NULL) { - clReleaseKernel( pState->m_kernel ); + if (pState->m_kernel != NULL) + { + clReleaseKernel(pState->m_kernel); pState->m_kernel = NULL; } - if(pState->m_program != NULL) { - clReleaseProgram( pState->m_program ); + if (pState->m_program != NULL) + { + clReleaseProgram(pState->m_program); pState->m_program = NULL; } } -bufferStruct * newBufferStruct(size_t inSize, size_t outSize, clState * pClState) { +bufferStruct *newBufferStruct(size_t inSize, size_t outSize, clState *pClState) +{ int error; - bufferStruct * pResult = (bufferStruct *)malloc(sizeof(bufferStruct)); + bufferStruct *pResult = (bufferStruct *)malloc(sizeof(bufferStruct)); #if DEBUG_MEM_ALLOC log_info("malloc bufferStruct * %x\n", pResult); #endif @@ -124,23 +127,20 @@ bufferStruct * newBufferStruct(size_t inSize, size_t outSize, clState * pClState pResult->m_inBuffer = clCreateBuffer(pClState->m_context, CL_MEM_READ_ONLY, inSize, NULL, &error); - if( pResult->m_inBuffer == NULL ) + if (pResult->m_inBuffer == NULL) { - vlog_error( "clCreateArray failed for input (%d)\n", error ); + vlog_error("clCreateArray failed for input (%d)\n", error); return destroyBufferStruct(pResult, pClState); } #if DEBUG_MEM_ALLOC log_info("clCreateBuffer %x\n", pResult->m_inBuffer); #endif - pResult->m_outBuffer = clCreateBuffer( pClState->m_context, - CL_MEM_WRITE_ONLY, - outSize, - NULL, - &error ); - if( pResult->m_outBuffer == NULL ) + pResult->m_outBuffer = clCreateBuffer( + pClState->m_context, CL_MEM_WRITE_ONLY, outSize, NULL, &error); + if (pResult->m_outBuffer == NULL) { - vlog_error( "clCreateArray failed for output (%d)\n", error ); + vlog_error("clCreateArray failed for output (%d)\n", error); return destroyBufferStruct(pResult, pClState); } #if DEBUG_MEM_ALLOC @@ -152,31 +152,36 @@ bufferStruct * newBufferStruct(size_t inSize, size_t outSize, clState * pClState return pResult; } -bufferStruct * destroyBufferStruct(bufferStruct * destroyMe, clState * pClState) { - if(destroyMe) +bufferStruct *destroyBufferStruct(bufferStruct *destroyMe, clState *pClState) +{ + if (destroyMe) { - if(destroyMe->m_outBuffer != NULL) { + if (destroyMe->m_outBuffer != NULL) + { #if DEBUG_MEM_ALLOC log_info("clReleaseMemObject %x\n", destroyMe->m_outBuffer); #endif clReleaseMemObject(destroyMe->m_outBuffer); destroyMe->m_outBuffer = NULL; } - if(destroyMe->m_inBuffer != NULL) { + if (destroyMe->m_inBuffer != NULL) + { #if DEBUG_MEM_ALLOC log_info("clReleaseMemObject %x\n", destroyMe->m_outBuffer); #endif clReleaseMemObject(destroyMe->m_inBuffer); destroyMe->m_inBuffer = NULL; } - if(destroyMe->m_pIn != NULL) { + if (destroyMe->m_pIn != NULL) + { #if DEBUG_MEM_ALLOC log_info("delete (free) m_pIn %x\n", destroyMe->m_pIn); #endif free(destroyMe->m_pIn); destroyMe->m_pIn = NULL; } - if(destroyMe->m_pOut != NULL) { + if (destroyMe->m_pOut != NULL) + { #if DEBUG_MEM_ALLOC log_info("delete (free) m_pOut %x\n", destroyMe->m_pOut); #endif @@ -192,46 +197,49 @@ bufferStruct * destroyBufferStruct(bufferStruct * destroyMe, clState * pClState) return destroyMe; } -void initContents(bufferStruct * pBufferStruct, clState * pClState, - size_t typeSize, - size_t countIn, size_t countOut ) +void initContents(bufferStruct *pBufferStruct, clState *pClState, + size_t typeSize, size_t countIn, size_t countOut) { size_t i; uint64_t start = 0; - switch(typeSize) + switch (typeSize) { case 1: { - uint8_t* ub = (uint8_t *)(pBufferStruct->m_pIn); - for (i=0; i < countIn; ++i) + uint8_t *ub = (uint8_t *)(pBufferStruct->m_pIn); + for (i = 0; i < countIn; ++i) { ub[i] = (uint8_t)start++; } break; } case 2: { - uint16_t* us = (uint16_t *)(pBufferStruct->m_pIn); - for (i=0; i < countIn; ++i) + uint16_t *us = (uint16_t *)(pBufferStruct->m_pIn); + for (i = 0; i < countIn; ++i) { us[i] = (uint16_t)start++; } break; } case 4: { - if (!g_wimpyMode) { - uint32_t* ui = (uint32_t *)(pBufferStruct->m_pIn); - for (i=0; i < countIn; ++i) { + if (!g_wimpyMode) + { + uint32_t *ui = (uint32_t *)(pBufferStruct->m_pIn); + for (i = 0; i < countIn; ++i) + { ui[i] = (uint32_t)start++; } } - else { - // The short test doesn't iterate over the entire 32 bit space so - // we alternate between positive and negative values - int32_t* ui = (int32_t *)(pBufferStruct->m_pIn); + else + { + // The short test doesn't iterate over the entire 32 bit space + // so we alternate between positive and negative values + int32_t *ui = (int32_t *)(pBufferStruct->m_pIn); int32_t sign = 1; - for (i=0; i < countIn; ++i, ++start) { - ui[i] = (int32_t)start*sign; + for (i = 0; i < countIn; ++i, ++start) + { + ui[i] = (int32_t)start * sign; sign = sign * -1; } } @@ -240,10 +248,11 @@ void initContents(bufferStruct * pBufferStruct, clState * pClState, case 8: { // We don't iterate over the entire space of 64 bit so for the // selects, we want to test positive and negative values - int64_t* ll = (int64_t *)(pBufferStruct->m_pIn); + int64_t *ll = (int64_t *)(pBufferStruct->m_pIn); int64_t sign = 1; - for (i=0; i < countIn; ++i, ++start) { - ll[i] = start*sign; + for (i = 0; i < countIn; ++i, ++start) + { + ll[i] = start * sign; sign = sign * -1; } break; @@ -256,10 +265,10 @@ void initContents(bufferStruct * pBufferStruct, clState * pClState, // pBufferStruct->m_bufSizeOut } -int pushArgs(bufferStruct * pBufferStruct, clState * pClState) +int pushArgs(bufferStruct *pBufferStruct, clState *pClState) { int err; - if( !pBufferStruct->m_bufferUploaded ) + if (!pBufferStruct->m_bufferUploaded) { err = clEnqueueWriteBuffer(pClState->m_queue, pBufferStruct->m_inBuffer, CL_TRUE, 0, pBufferStruct->m_bufSizeIn, @@ -267,7 +276,7 @@ int pushArgs(bufferStruct * pBufferStruct, clState * pClState) #if DEBUG_MEM_ALLOC log_info("clEnqueueWriteBuffer %x\n", pBufferStruct->m_inBuffer); #endif - if(err != CL_SUCCESS) + if (err != CL_SUCCESS) { log_error("clEnqueueWriteBuffer failed\n"); return -1; @@ -275,22 +284,24 @@ int pushArgs(bufferStruct * pBufferStruct, clState * pClState) pBufferStruct->m_bufferUploaded = true; } - err = clSetKernelArg(pClState->m_kernel, 0, - sizeof(pBufferStruct->m_inBuffer), // pBufferStruct->m_bufSizeIn, - &(pBufferStruct->m_inBuffer)); + err = clSetKernelArg( + pClState->m_kernel, 0, + sizeof(pBufferStruct->m_inBuffer), // pBufferStruct->m_bufSizeIn, + &(pBufferStruct->m_inBuffer)); #if DEBUG_MEM_ALLOC // log_info("clSetKernelArg 0, %x\n", pBufferStruct->m_inBuffer); #endif - if(err != CL_SUCCESS) + if (err != CL_SUCCESS) { log_error("clSetKernelArgs failed, first arg (0)\n"); return -1; } - err = clSetKernelArg(pClState->m_kernel, 1, - sizeof(pBufferStruct->m_outBuffer), // pBufferStruct->m_bufSizeOut, - &(pBufferStruct->m_outBuffer)); - if(err != CL_SUCCESS) + err = clSetKernelArg( + pClState->m_kernel, 1, + sizeof(pBufferStruct->m_outBuffer), // pBufferStruct->m_bufSizeOut, + &(pBufferStruct->m_outBuffer)); + if (err != CL_SUCCESS) { log_error("clSetKernelArgs failed, second arg (1)\n"); return -1; @@ -303,13 +314,13 @@ int pushArgs(bufferStruct * pBufferStruct, clState * pClState) return 0; } -int retrieveResults(bufferStruct * pBufferStruct, clState * pClState) +int retrieveResults(bufferStruct *pBufferStruct, clState *pClState) { int err; err = clEnqueueReadBuffer(pClState->m_queue, pBufferStruct->m_outBuffer, CL_TRUE, 0, pBufferStruct->m_bufSizeOut, pBufferStruct->m_pOut, 0, NULL, NULL); - if(err != CL_SUCCESS) + if (err != CL_SUCCESS) { log_error("clEnqueueReadBuffer failed\n"); return -1; @@ -319,19 +330,17 @@ int retrieveResults(bufferStruct * pBufferStruct, clState * pClState) // vecSizeIdx indexes into g_arrVecAlignMasks, g_arrVecSizeNames // and g_arrVecSizes -int checkCorrectness(bufferStruct * pBufferStruct, clState * pClState, - size_t minAlign) +int checkCorrectnessAlign(bufferStruct *pBufferStruct, clState *pClState, + size_t minAlign) { size_t i; - cl_uint * targetArr = (cl_uint *)(pBufferStruct->m_pOut); - for(i = 0; i < pClState->m_numThreads; ++i) + cl_uint *targetArr = (cl_uint *)(pBufferStruct->m_pOut); + for (i = 0; i < pClState->m_numThreads; ++i) { - if((targetArr[i])%minAlign != (cl_uint)0) + if ((targetArr[i]) % minAlign != (cl_uint)0) { vlog_error("Error %d (of %d). Expected a multple of %x, got %x\n", - i, pClState->m_numThreads, - minAlign, - targetArr[i]); + i, pClState->m_numThreads, minAlign, targetArr[i]); return -1; } } @@ -345,21 +354,42 @@ int checkCorrectness(bufferStruct * pBufferStruct, clState * pClState, return 0; } +int checkCorrectnessStep(bufferStruct *pBufferStruct, clState *pClState, + size_t typeSize, size_t vecWidth) +{ + size_t i; + cl_int targetSize = (cl_int)vecWidth; + cl_int *targetArr = (cl_int *)(pBufferStruct->m_pOut); + if (targetSize == 3) + { + targetSize = 4; // hack for 4-aligned vec3 types + } + for (i = 0; i < pClState->m_numThreads; ++i) + { + if (targetArr[i] != targetSize) + { + vlog_error("Error %ld (of %ld). Expected %d, got %d\n", i, + pClState->m_numThreads, targetSize, targetArr[i]); + return -1; + } + } + return 0; +} // vecSizeIdx indexes into g_arrVecAlignMasks, g_arrVecSizeNames // and g_arrVecSizes -int checkPackedCorrectness(bufferStruct * pBufferStruct, clState * pClState, +int checkPackedCorrectness(bufferStruct *pBufferStruct, clState *pClState, size_t totSize, size_t beforeSize) { size_t i; - cl_uint * targetArr = (cl_uint *)(pBufferStruct->m_pOut); - for(i = 0; i < pClState->m_numThreads; ++i) + cl_uint *targetArr = (cl_uint *)(pBufferStruct->m_pOut); + for (i = 0; i < pClState->m_numThreads; ++i) { - if((targetArr[i]-beforeSize)%totSize != (cl_uint)0) + if ((targetArr[i] - beforeSize) % totSize != (cl_uint)0) { - vlog_error("Error %d (of %d). Expected %d more than a multple of %d, got %d \n", - i, pClState->m_numThreads, beforeSize, - totSize, + vlog_error("Error %d (of %d). Expected %d more than a multple of " + "%d, got %d \n", + i, pClState->m_numThreads, beforeSize, totSize, targetArr[i]); return -1; } diff --git a/test_conformance/vectors/structs.h b/test_conformance/vectors/structs.h new file mode 100644 index 00000000..c6a1725e --- /dev/null +++ b/test_conformance/vectors/structs.h @@ -0,0 +1,75 @@ +// +// 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 "testBase.h" + + +#include "harness/conversions.h" +#include "harness/typeWrappers.h" + +typedef struct _clState +{ + cl_device_id m_device; + cl_context m_context; + cl_command_queue m_queue; + + cl_program m_program; + cl_kernel m_kernel; + size_t m_numThreads; +} clState; + +clState* newClState(cl_device_id device, cl_context context, + cl_command_queue queue); +clState* destroyClState(clState* pState); + +int clStateMakeProgram(clState* pState, const char* prog, + const char* kernelName); +void clStateDestroyProgramAndKernel(clState* pState); + +int runKernel(clState* pState, size_t numThreads); + +typedef struct _bufferStruct +{ + void* m_pIn; + void* m_pOut; + + cl_mem m_outBuffer; + cl_mem m_inBuffer; + + size_t m_bufSizeIn, m_bufSizeOut; + + int m_bufferUploaded; +} bufferStruct; + + +bufferStruct* newBufferStruct(size_t inSize, size_t outSize, clState* pClState); + +bufferStruct* destroyBufferStruct(bufferStruct* destroyMe, clState* pClState); + +void initContents(bufferStruct* pBufferStruct, clState* pClState, + size_t typeSize, size_t vecWidth); + +int pushArgs(bufferStruct* pBufferStruct, clState* pClState); +int retrieveResults(bufferStruct* pBufferStruct, clState* pClState); + +int checkCorrectnessStep(bufferStruct* pBufferStruct, clState* pClState, + size_t typeSize, size_t vecWidth); +// vecSizeIdx indexes into g_arrVecAlignMasks, g_arrVecSizeNames +// and g_arrVecSizes +int checkCorrectnessAlign(bufferStruct* pBufferStruct, clState* pClState, + size_t minAlign); + +int checkPackedCorrectness(bufferStruct* pBufferStruct, clState* pClState, + size_t totSize, size_t beforeSize); diff --git a/test_conformance/vec_align/testBase.h b/test_conformance/vectors/testBase.h similarity index 99% rename from test_conformance/vec_align/testBase.h rename to test_conformance/vectors/testBase.h index bd72e841..63086d7e 100644 --- a/test_conformance/vec_align/testBase.h +++ b/test_conformance/vectors/testBase.h @@ -1,6 +1,6 @@ // // 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 diff --git a/test_conformance/vec_step/test_step.cpp b/test_conformance/vectors/test_step.cpp similarity index 65% rename from test_conformance/vec_step/test_step.cpp rename to test_conformance/vectors/test_step.cpp index a2c57c39..2f6ad187 100644 --- a/test_conformance/vec_step/test_step.cpp +++ b/test_conformance/vectors/test_step.cpp @@ -1,6 +1,6 @@ // // 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 @@ -35,94 +35,99 @@ */ - -int test_step_internal(cl_device_id deviceID, cl_context context, cl_command_queue queue, const char * pattern, const char * testName) +int test_step_internal(cl_device_id deviceID, cl_context context, + cl_command_queue queue, const char* pattern, + const char* testName) { int err; int typeIdx, vecSizeIdx; char tempBuffer[2048]; - clState * pClState = newClState(deviceID, context, queue); - bufferStruct * pBuffers = - newBufferStruct(BUFFER_SIZE, BUFFER_SIZE, pClState); + clState* pClState = newClState(deviceID, context, queue); + bufferStruct* pBuffers = + newBufferStruct(BUFFER_SIZE, BUFFER_SIZE, pClState); - if(pBuffers == NULL) { + if (pBuffers == NULL) + { destroyClState(pClState); vlog_error("%s : Could not create buffer\n", testName); return -1; } - //detect whether profile of the device is embedded + // detect whether profile of the device is embedded char profile[1024] = ""; - err = clGetDeviceInfo(deviceID, CL_DEVICE_PROFILE, sizeof(profile), profile, NULL); + err = clGetDeviceInfo(deviceID, CL_DEVICE_PROFILE, sizeof(profile), profile, + NULL); if (err) { - print_error(err, "clGetDeviceInfo for CL_DEVICE_PROFILE failed\n" ); + print_error(err, "clGetDeviceInfo for CL_DEVICE_PROFILE failed\n"); return -1; } gIsEmbedded = NULL != strstr(profile, "EMBEDDED_PROFILE"); - for(typeIdx = 0; types[typeIdx] != kNumExplicitTypes; ++typeIdx) + for (typeIdx = 0; types[typeIdx] != kNumExplicitTypes; ++typeIdx) { - if( types[ typeIdx ] == kDouble ) + if (types[typeIdx] == kDouble) { // If we're testing doubles, we need to check for support first - if( !is_extension_available( deviceID, "cl_khr_fp64" ) ) + if (!is_extension_available(deviceID, "cl_khr_fp64")) { - log_info( "Not testing doubles (unsupported on this device)\n" ); + log_info("Not testing doubles (unsupported on this device)\n"); continue; } } - if( types[ typeIdx ] == kLong || types[ typeIdx ] == kULong ) - { - // If we're testing long/ulong, we need to check for embedded support - if( gIsEmbedded && !is_extension_available( deviceID, "cles_khr_int64") ) + if (types[typeIdx] == kLong || types[typeIdx] == kULong) { - log_info( "Not testing longs (unsupported on this embedded device)\n" ); - continue; + // If we're testing long/ulong, we need to check for embedded + // support + if (gIsEmbedded + && !is_extension_available(deviceID, "cles_khr_int64")) + { + log_info("Not testing longs (unsupported on this embedded " + "device)\n"); + continue; + } } - } char srcBuffer[2048]; - doSingleReplace(tempBuffer, 2048, pattern, - ".EXTENSIONS.", types[typeIdx] == kDouble + doSingleReplace(tempBuffer, 2048, pattern, ".EXTENSIONS.", + types[typeIdx] == kDouble ? "#pragma OPENCL EXTENSION cl_khr_fp64 : enable" : ""); - for(vecSizeIdx = 0; vecSizeIdx < NUM_VECTOR_SIZES; ++vecSizeIdx) + for (vecSizeIdx = 0; vecSizeIdx < NUM_VECTOR_SIZES; ++vecSizeIdx) { - doReplace(srcBuffer, 2048, tempBuffer, - ".TYPE.", g_arrTypeNames[typeIdx], - ".NUM.", g_arrVecSizeNames[vecSizeIdx]); + doReplace(srcBuffer, 2048, tempBuffer, ".TYPE.", + g_arrTypeNames[typeIdx], ".NUM.", + g_arrVecSizeNames[vecSizeIdx]); - if(srcBuffer[0] == '\0') { + if (srcBuffer[0] == '\0') + { vlog_error("%s: failed to fill source buf for type %s%s\n", - testName, - g_arrTypeNames[typeIdx], + testName, g_arrTypeNames[typeIdx], g_arrVecSizeNames[vecSizeIdx]); destroyBufferStruct(pBuffers, pClState); destroyClState(pClState); return -1; } - err = clStateMakeProgram(pClState, srcBuffer, testName ); + err = clStateMakeProgram(pClState, srcBuffer, testName); if (err) { - vlog_error("%s: Error compiling \"\n%s\n\"", - testName, srcBuffer); + vlog_error("%s: Error compiling \"\n%s\n\"", testName, + srcBuffer); destroyBufferStruct(pBuffers, pClState); destroyClState(pClState); return -1; } err = pushArgs(pBuffers, pClState); - if(err != 0) + if (err != 0) { - vlog_error("%s: failed to push args %s%s\n", - testName, + vlog_error("%s: failed to push args %s%s\n", testName, g_arrTypeNames[typeIdx], g_arrVecSizeNames[vecSizeIdx]); destroyBufferStruct(pBuffers, pClState); @@ -132,11 +137,10 @@ int test_step_internal(cl_device_id deviceID, cl_context context, cl_command_que // now we run the kernel err = runKernel(pClState, 1024); - if(err != 0) + if (err != 0) { - vlog_error("%s: runKernel fail (%ld threads) %s%s\n", - testName, pClState->m_numThreads, - g_arrTypeNames[typeIdx], + vlog_error("%s: runKernel fail (%ld threads) %s%s\n", testName, + pClState->m_numThreads, g_arrTypeNames[typeIdx], g_arrVecSizeNames[vecSizeIdx]); destroyBufferStruct(pBuffers, pClState); destroyClState(pClState); @@ -144,10 +148,9 @@ int test_step_internal(cl_device_id deviceID, cl_context context, cl_command_que } err = retrieveResults(pBuffers, pClState); - if(err != 0) + if (err != 0) { - vlog_error("%s: failed to retrieve results %s%s\n", - testName, + vlog_error("%s: failed to retrieve results %s%s\n", testName, g_arrTypeNames[typeIdx], g_arrVecSizeNames[vecSizeIdx]); destroyBufferStruct(pBuffers, pClState); @@ -155,24 +158,21 @@ int test_step_internal(cl_device_id deviceID, cl_context context, cl_command_que return -1; } - err = checkCorrectness(pBuffers, pClState, - g_arrTypeSizes[typeIdx], - g_arrVecSizes[vecSizeIdx]); + err = checkCorrectnessStep(pBuffers, pClState, + g_arrTypeSizes[typeIdx], + g_arrVecSizes[vecSizeIdx]); - if(err != 0) + if (err != 0) { - vlog_error("%s: incorrect results %s%s\n", - testName, + vlog_error("%s: incorrect results %s%s\n", testName, g_arrTypeNames[typeIdx], g_arrVecSizeNames[vecSizeIdx]); - vlog_error("%s: Source was \"\n%s\n\"", - testName, srcBuffer); + vlog_error("%s: Source was \"\n%s\n\"", testName, srcBuffer); destroyBufferStruct(pBuffers, pClState); destroyClState(pClState); return -1; } } - } destroyBufferStruct(pBuffers, pClState); @@ -184,9 +184,10 @@ int test_step_internal(cl_device_id deviceID, cl_context context, cl_command_que return 0; // -1; // fails on account of not being written. } -const char * patterns[] = { +static const char* patterns[] = { ".EXTENSIONS.\n" - "__kernel void test_step_type(__global .TYPE..NUM. *source, __global int *dest)\n" + "__kernel void test_step_type(__global .TYPE..NUM. *source, __global int " + "*dest)\n" "{\n" " int tid = get_global_id(0);\n" " dest[tid] = vec_step(.TYPE..NUM.);\n" @@ -194,7 +195,8 @@ const char * patterns[] = { "}\n", ".EXTENSIONS.\n" - "__kernel void test_step_var(__global .TYPE..NUM. *source, __global int *dest)\n" + "__kernel void test_step_var(__global .TYPE..NUM. *source, __global int " + "*dest)\n" "{\n" " int tid = get_global_id(0);\n" " dest[tid] = vec_step(source[tid]);\n" @@ -203,7 +205,8 @@ const char * patterns[] = { ".EXTENSIONS.\n" " typedef .TYPE..NUM. TypeToTest;\n" - "__kernel void test_step_typedef_type(__global TypeToTest *source, __global int *dest)\n" + "__kernel void test_step_typedef_type(__global TypeToTest *source, " + "__global int *dest)\n" "{\n" " int tid = get_global_id(0);\n" " dest[tid] = vec_step(TypeToTest);\n" @@ -212,7 +215,8 @@ const char * patterns[] = { ".EXTENSIONS.\n" " typedef .TYPE..NUM. TypeToTest;\n" - "__kernel void test_step_typedef_var(__global TypeToTest *source, __global int *dest)\n" + "__kernel void test_step_typedef_var(__global TypeToTest *source, __global " + "int *dest)\n" "{\n" " int tid = get_global_id(0);\n" " dest[tid] = vec_step(source[tid]);\n" @@ -227,25 +231,29 @@ const char * patterns[] = { test_step_typedef_var, */ -int test_step_type(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) +int test_step_type(cl_device_id deviceID, cl_context context, + cl_command_queue queue, int num_elements) { return test_step_internal(deviceID, context, queue, patterns[0], "test_step_type"); } -int test_step_var(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) +int test_step_var(cl_device_id deviceID, cl_context context, + cl_command_queue queue, int num_elements) { return test_step_internal(deviceID, context, queue, patterns[1], "test_step_var"); } -int test_step_typedef_type(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) +int test_step_typedef_type(cl_device_id deviceID, cl_context context, + cl_command_queue queue, int num_elements) { return test_step_internal(deviceID, context, queue, patterns[2], "test_step_typedef_type"); } -int test_step_typedef_var(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) +int test_step_typedef_var(cl_device_id deviceID, cl_context context, + cl_command_queue queue, int num_elements) { return test_step_internal(deviceID, context, queue, patterns[3], "test_step_typedef_var"); diff --git a/test_conformance/vec_align/test_vec_align.cpp b/test_conformance/vectors/test_vec_align.cpp similarity index 51% rename from test_conformance/vec_align/test_vec_align.cpp rename to test_conformance/vectors/test_vec_align.cpp index bc666a48..2f392f54 100644 --- a/test_conformance/vec_align/test_vec_align.cpp +++ b/test_conformance/vectors/test_vec_align.cpp @@ -1,6 +1,6 @@ // // 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 @@ -29,14 +29,15 @@ size_t get_align(size_t vecSize) { - if(vecSize == 3) + if (vecSize == 3) { return 4; } return vecSize; } -/* // Lots of conditionals means this is not gonna be an optimal min on intel. */ +/* // Lots of conditionals means this is not gonna be an optimal min on intel. + */ /* // That's okay, make sure we only call a few times per test, not for every */ /* // element */ /* size_t min_of_nonzero(size_t a, size_t b) */ @@ -53,7 +54,8 @@ size_t get_align(size_t vecSize) /* } */ -/* size_t get_min_packed_alignment(size_t preSize, size_t typeMultiplePreSize, */ +/* size_t get_min_packed_alignment(size_t preSize, size_t typeMultiplePreSize, + */ /* size_t postSize, size_t typeMultiplePostSize, */ /* ExplicitType kType, size_t vecSize) */ /* { */ @@ -71,12 +73,11 @@ size_t get_align(size_t vecSize) /* } */ - int test_vec_internal(cl_device_id deviceID, cl_context context, - cl_command_queue queue, const char * pattern, - const char * testName, size_t bufSize, - size_t preSize, size_t typeMultiplePreSize, - size_t postSize, size_t typeMultiplePostSize) + cl_command_queue queue, const char* pattern, + const char* testName, size_t bufSize, size_t preSize, + size_t typeMultiplePreSize, size_t postSize, + size_t typeMultiplePostSize) { int err; int typeIdx, vecSizeIdx; @@ -86,58 +87,64 @@ int test_vec_internal(cl_device_id deviceID, cl_context context, size_t preSizeBytes, postSizeBytes, typeSize, totSize; - clState * pClState = newClState(deviceID, context, queue); - bufferStruct * pBuffers = - newBufferStruct(bufSize, bufSize*sizeof(cl_uint)/sizeof(cl_char), pClState); + clState* pClState = newClState(deviceID, context, queue); + bufferStruct* pBuffers = newBufferStruct( + bufSize, bufSize * sizeof(cl_uint) / sizeof(cl_char), pClState); - if(pBuffers == NULL) { + if (pBuffers == NULL) + { destroyClState(pClState); vlog_error("%s : Could not create buffer\n", testName); return -1; } - for(typeIdx = 0; types[typeIdx] != kNumExplicitTypes; ++typeIdx) + for (typeIdx = 0; types[typeIdx] != kNumExplicitTypes; ++typeIdx) { // Skip doubles if it is not supported otherwise enable pragma - if (types[typeIdx] == kDouble) { - if (!is_extension_available(deviceID, "cl_khr_fp64")) { + if (types[typeIdx] == kDouble) + { + if (!is_extension_available(deviceID, "cl_khr_fp64")) + { continue; - } else { - doReplace(tmpBuffer, 2048, pattern, - ".PRAGMA.", "#pragma OPENCL EXTENSION cl_khr_fp64: ", - ".STATE.", "enable"); } - } else { - if (types[typeIdx] == kLong || types[typeIdx] == kULong) { - if (gIsEmbedded) - continue; + else + { + doReplace(tmpBuffer, 2048, pattern, ".PRAGMA.", + "#pragma OPENCL EXTENSION cl_khr_fp64: ", ".STATE.", + "enable"); + } + } + else + { + if (types[typeIdx] == kLong || types[typeIdx] == kULong) + { + if (gIsEmbedded) continue; } - doReplace(tmpBuffer, 2048, pattern, - ".PRAGMA.", " ", - ".STATE.", " "); + doReplace(tmpBuffer, 2048, pattern, ".PRAGMA.", " ", ".STATE.", + " "); } typeSize = get_explicit_type_size(types[typeIdx]); - preSizeBytes = preSize + typeSize*typeMultiplePreSize; - postSizeBytes = postSize + typeSize*typeMultiplePostSize; + preSizeBytes = preSize + typeSize * typeMultiplePreSize; + postSizeBytes = postSize + typeSize * typeMultiplePostSize; + for (vecSizeIdx = 1; vecSizeIdx < NUM_VECTOR_SIZES; ++vecSizeIdx) + { - for(vecSizeIdx = 1; vecSizeIdx < NUM_VECTOR_SIZES; ++vecSizeIdx) { + totSize = preSizeBytes + postSizeBytes + + typeSize * get_align(g_arrVecSizes[vecSizeIdx]); - totSize = preSizeBytes + postSizeBytes + - typeSize*get_align(g_arrVecSizes[vecSizeIdx]); + doReplace(srcBuffer, 2048, tmpBuffer, ".TYPE.", + g_arrTypeNames[typeIdx], ".NUM.", + g_arrVecSizeNames[vecSizeIdx]); - doReplace(srcBuffer, 2048, tmpBuffer, - ".TYPE.", g_arrTypeNames[typeIdx], - ".NUM.", g_arrVecSizeNames[vecSizeIdx]); - - if(srcBuffer[0] == '\0') { + if (srcBuffer[0] == '\0') + { vlog_error("%s: failed to fill source buf for type %s%s\n", - testName, - g_arrTypeNames[typeIdx], + testName, g_arrTypeNames[typeIdx], g_arrVecSizeNames[vecSizeIdx]); destroyBufferStruct(pBuffers, pClState); destroyClState(pClState); @@ -147,19 +154,20 @@ int test_vec_internal(cl_device_id deviceID, cl_context context, // log_info("Buffer is \"\n%s\n\"\n", srcBuffer); // fflush(stdout); - err = clStateMakeProgram(pClState, srcBuffer, testName ); - if (err) { - vlog_error("%s: Error compiling \"\n%s\n\"", - testName, srcBuffer); + err = clStateMakeProgram(pClState, srcBuffer, testName); + if (err) + { + vlog_error("%s: Error compiling \"\n%s\n\"", testName, + srcBuffer); destroyBufferStruct(pBuffers, pClState); destroyClState(pClState); return -1; } err = pushArgs(pBuffers, pClState); - if(err != 0) { - vlog_error("%s: failed to push args %s%s\n", - testName, + if (err != 0) + { + vlog_error("%s: failed to push args %s%s\n", testName, g_arrTypeNames[typeIdx], g_arrVecSizeNames[vecSizeIdx]); destroyBufferStruct(pBuffers, pClState); @@ -169,12 +177,14 @@ int test_vec_internal(cl_device_id deviceID, cl_context context, // log_info("About to Run kernel\n"); fflush(stdout); // now we run the kernel - err = runKernel(pClState, - bufSize/(g_arrVecSizes[vecSizeIdx]* g_arrTypeSizes[typeIdx])); - if(err != 0) { - vlog_error("%s: runKernel fail (%ld threads) %s%s\n", - testName, pClState->m_numThreads, - g_arrTypeNames[typeIdx], + err = runKernel( + pClState, + bufSize + / (g_arrVecSizes[vecSizeIdx] * g_arrTypeSizes[typeIdx])); + if (err != 0) + { + vlog_error("%s: runKernel fail (%ld threads) %s%s\n", testName, + pClState->m_numThreads, g_arrTypeNames[typeIdx], g_arrVecSizeNames[vecSizeIdx]); destroyBufferStruct(pBuffers, pClState); destroyClState(pClState); @@ -183,9 +193,9 @@ int test_vec_internal(cl_device_id deviceID, cl_context context, // log_info("About to retrieve results\n"); fflush(stdout); err = retrieveResults(pBuffers, pClState); - if(err != 0) { - vlog_error("%s: failed to retrieve results %s%s\n", - testName, + if (err != 0) + { + vlog_error("%s: failed to retrieve results %s%s\n", testName, g_arrTypeNames[typeIdx], g_arrVecSizeNames[vecSizeIdx]); destroyBufferStruct(pBuffers, pClState); @@ -194,13 +204,12 @@ int test_vec_internal(cl_device_id deviceID, cl_context context, } - - if(preSizeBytes+postSizeBytes == 0) + if (preSizeBytes + postSizeBytes == 0) { // log_info("About to Check Correctness\n"); fflush(stdout); - err = checkCorrectness(pBuffers, pClState, - get_align(g_arrVecSizes[vecSizeIdx])* - typeSize); + err = checkCorrectnessAlign(pBuffers, pClState, + get_align(g_arrVecSizes[vecSizeIdx]) + * typeSize); } else { @@ -209,20 +218,18 @@ int test_vec_internal(cl_device_id deviceID, cl_context context, preSizeBytes); } - if(err != 0) { - vlog_error("%s: incorrect results %s%s\n", - testName, + if (err != 0) + { + vlog_error("%s: incorrect results %s%s\n", testName, g_arrTypeNames[typeIdx], g_arrVecSizeNames[vecSizeIdx]); - vlog_error("%s: Source was \"\n%s\n\"", - testName, srcBuffer); + vlog_error("%s: Source was \"\n%s\n\"", testName, srcBuffer); destroyBufferStruct(pBuffers, pClState); destroyClState(pClState); return -1; } clStateDestroyProgramAndKernel(pClState); - } } @@ -236,10 +243,10 @@ int test_vec_internal(cl_device_id deviceID, cl_context context, } - -const char * patterns[] = { +static const char* patterns[] = { ".PRAGMA..STATE.\n" - "__kernel void test_vec_align_array(.SRC_SCOPE. .TYPE..NUM. *source, .DST_SCOPE. uint *dest)\n" + "__kernel void test_vec_align_array(.SRC_SCOPE. .TYPE..NUM. *source, " + ".DST_SCOPE. uint *dest)\n" "{\n" " int tid = get_global_id(0);\n" " dest[tid] = (uint)((.SRC_SCOPE. uchar *)(source+tid));\n" @@ -250,7 +257,8 @@ const char * patterns[] = { " .TYPE..NUM. vec;\n" ".POST." "} testStruct;\n" - "__kernel void test_vec_align_struct(__constant .TYPE..NUM. *source, .DST_SCOPE. uint *dest)\n" + "__kernel void test_vec_align_struct(__constant .TYPE..NUM. *source, " + ".DST_SCOPE. uint *dest)\n" "{\n" " .SRC_SCOPE. testStruct test;\n" " int tid = get_global_id(0);\n" @@ -262,11 +270,13 @@ const char * patterns[] = { " .TYPE..NUM. vec;\n" ".POST." "} testStruct;\n" - "__kernel void test_vec_align_packed_struct(__constant .TYPE..NUM. *source, .DST_SCOPE. uint *dest)\n" + "__kernel void test_vec_align_packed_struct(__constant .TYPE..NUM. " + "*source, .DST_SCOPE. uint *dest)\n" "{\n" " .SRC_SCOPE. testStruct test;\n" " int tid = get_global_id(0);\n" - " dest[tid] = (uint)((.SRC_SCOPE. uchar *)&(test.vec) - (.SRC_SCOPE. uchar *)&test);\n" + " dest[tid] = (uint)((.SRC_SCOPE. uchar *)&(test.vec) - (.SRC_SCOPE. " + "uchar *)&test);\n" "}\n", ".PRAGMA..STATE.\n" "typedef struct myStruct { \n" @@ -274,7 +284,8 @@ const char * patterns[] = { " .TYPE..NUM. vec;\n" ".POST." "} testStruct;\n" - "__kernel void test_vec_align_struct_arr(.SRC_SCOPE. testStruct *source, .DST_SCOPE. uint *dest)\n" + "__kernel void test_vec_align_struct_arr(.SRC_SCOPE. testStruct *source, " + ".DST_SCOPE. uint *dest)\n" "{\n" " int tid = get_global_id(0);\n" " dest[tid] = (uint)((.SRC_SCOPE. uchar *)&(source[tid].vec));\n" @@ -285,142 +296,123 @@ const char * patterns[] = { " .TYPE..NUM. vec;\n" ".POST." "} testStruct;\n" - "__kernel void test_vec_align_packed_struct_arr(.SRC_SCOPE. testStruct *source, .DST_SCOPE. uint *dest)\n" + "__kernel void test_vec_align_packed_struct_arr(.SRC_SCOPE. testStruct " + "*source, .DST_SCOPE. uint *dest)\n" "{\n" " int tid = get_global_id(0);\n" - " dest[tid] = (uint)((.SRC_SCOPE. uchar *)&(source[tid].vec) - (.SRC_SCOPE. uchar *)&(source[0]));\n" + " dest[tid] = (uint)((.SRC_SCOPE. uchar *)&(source[tid].vec) - " + "(.SRC_SCOPE. uchar *)&(source[0]));\n" "}\n", // __attribute__ ((packed)) }; - -const char * pre_substitution_arr[] = { - "", - "char c;\n", - "short3 s;", - ".TYPE.3 tPre;\n", - ".TYPE. arrPre[5];\n", - ".TYPE. arrPre[12];\n", - NULL -}; +const char* pre_substitution_arr[] = { "", + "char c;\n", + "short3 s;", + ".TYPE.3 tPre;\n", + ".TYPE. arrPre[5];\n", + ".TYPE. arrPre[12];\n", + NULL }; // alignments of everything in pre_substitution_arr as raw alignments // 0 if such a thing is meaningless -size_t pre_align_arr[] = { - 0, - sizeof(cl_char), - 4*sizeof(cl_short), - 0, // taken care of in type_multiple_pre_align_arr - 0, - 0 -}; +size_t pre_align_arr[] = { 0, + sizeof(cl_char), + 4 * sizeof(cl_short), + 0, // taken care of in type_multiple_pre_align_arr + 0, + 0 }; // alignments of everything in pre_substitution_arr as multiples of // sizeof(.TYPE.) // 0 if such a thing is meaningless -size_t type_multiple_pre_align_arr[] = { - 0, - 0, - 0, - 4, - 5, - 12 -}; +size_t type_multiple_pre_align_arr[] = { 0, 0, 0, 4, 5, 12 }; -const char * post_substitution_arr[] = { - "", - "char cPost;\n", - ".TYPE. arrPost[3];\n", - ".TYPE. arrPost[5];\n", - ".TYPE.3 arrPost;\n", - ".TYPE. arrPost[12];\n", - NULL -}; +const char* post_substitution_arr[] = { "", + "char cPost;\n", + ".TYPE. arrPost[3];\n", + ".TYPE. arrPost[5];\n", + ".TYPE.3 arrPost;\n", + ".TYPE. arrPost[12];\n", + NULL }; // alignments of everything in post_substitution_arr as raw alignments // 0 if such a thing is meaningless -size_t post_align_arr[] = { - 0, - sizeof(cl_char), - 0, // taken care of in type_multiple_post_align_arr - 0, - 0, - 0 -}; +size_t post_align_arr[] = { 0, sizeof(cl_char), + 0, // taken care of in type_multiple_post_align_arr + 0, 0, + 0 }; // alignments of everything in post_substitution_arr as multiples of // sizeof(.TYPE.) // 0 if such a thing is meaningless -size_t type_multiple_post_align_arr[] = { - 0, - 0, - 3, - 5, - 4, - 12 -}; +size_t type_multiple_post_align_arr[] = { 0, 0, 3, 5, 4, 12 }; // there hsould be a packed version of this? -int test_vec_align_array(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) +int test_vec_align_array(cl_device_id deviceID, cl_context context, + cl_command_queue queue, int num_elements) { char tmp[2048]; int result; log_info("Testing global\n"); - doReplace(tmp, (size_t)2048, patterns[0], - ".SRC_SCOPE.", "__global", + doReplace(tmp, (size_t)2048, patterns[0], ".SRC_SCOPE.", "__global", ".DST_SCOPE.", "__global"); // result = test_vec_internal(deviceID, context, queue, tmp, - "test_vec_align_array", - BUFFER_SIZE, 0, 0, 0, 0); + "test_vec_align_array", BUFFER_SIZE, 0, 0, 0, 0); return result; } -int test_vec_align_struct(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) +int test_vec_align_struct(cl_device_id deviceID, cl_context context, + cl_command_queue queue, int num_elements) { char tmp1[2048], tmp2[2048]; int result = 0; int preIdx, postIdx; log_info("testing __private\n"); - doReplace(tmp2, (size_t)2048, patterns[1], - ".SRC_SCOPE.", "__private", + doReplace(tmp2, (size_t)2048, patterns[1], ".SRC_SCOPE.", "__private", ".DST_SCOPE.", "__global"); // - for(preIdx = 0; pre_substitution_arr[preIdx] != NULL; ++preIdx) { - for(postIdx = 0; post_substitution_arr[postIdx] != NULL; ++postIdx) { - doReplace(tmp1, (size_t)2048, tmp2, - ".PRE.", pre_substitution_arr[preIdx], - ".POST.", post_substitution_arr[postIdx]); + for (preIdx = 0; pre_substitution_arr[preIdx] != NULL; ++preIdx) + { + for (postIdx = 0; post_substitution_arr[postIdx] != NULL; ++postIdx) + { + doReplace(tmp1, (size_t)2048, tmp2, ".PRE.", + pre_substitution_arr[preIdx], ".POST.", + post_substitution_arr[postIdx]); - result = test_vec_internal(deviceID, context, queue, tmp1, - "test_vec_align_struct", - 512, 0, 0, 0, 0); - if (result != 0) { + result = + test_vec_internal(deviceID, context, queue, tmp1, + "test_vec_align_struct", 512, 0, 0, 0, 0); + if (result != 0) + { return result; } } } log_info("testing __local\n"); - doReplace(tmp2, (size_t)2048, patterns[1], - ".SRC_SCOPE.", "__local", + doReplace(tmp2, (size_t)2048, patterns[1], ".SRC_SCOPE.", "__local", ".DST_SCOPE.", "__global"); // - for(preIdx = 0; pre_substitution_arr[preIdx] != NULL; ++preIdx) { - for(postIdx = 0; post_substitution_arr[postIdx] != NULL; ++postIdx) { - doReplace(tmp1, (size_t)2048, tmp2, - ".PRE.", pre_substitution_arr[preIdx], - ".POST.", post_substitution_arr[postIdx]); + for (preIdx = 0; pre_substitution_arr[preIdx] != NULL; ++preIdx) + { + for (postIdx = 0; post_substitution_arr[postIdx] != NULL; ++postIdx) + { + doReplace(tmp1, (size_t)2048, tmp2, ".PRE.", + pre_substitution_arr[preIdx], ".POST.", + post_substitution_arr[postIdx]); - result = test_vec_internal(deviceID, context, queue, tmp1, - "test_vec_align_struct", - 512, 0, 0, 0, 0); - if(result != 0) { + result = + test_vec_internal(deviceID, context, queue, tmp1, + "test_vec_align_struct", 512, 0, 0, 0, 0); + if (result != 0) + { return result; } } @@ -428,7 +420,8 @@ int test_vec_align_struct(cl_device_id deviceID, cl_context context, cl_command_ return 0; } -int test_vec_align_packed_struct(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) +int test_vec_align_packed_struct(cl_device_id deviceID, cl_context context, + cl_command_queue queue, int num_elements) { char tmp1[2048], tmp2[2048]; int result = 0; @@ -436,46 +429,46 @@ int test_vec_align_packed_struct(cl_device_id deviceID, cl_context context, cl_c log_info("Testing __private\n"); - doReplace(tmp2, (size_t)2048, patterns[2], - ".SRC_SCOPE.", "__private", + doReplace(tmp2, (size_t)2048, patterns[2], ".SRC_SCOPE.", "__private", ".DST_SCOPE.", "__global"); // - for(preIdx = 0; pre_substitution_arr[preIdx] != NULL; ++preIdx) { - for(postIdx = 0; post_substitution_arr[postIdx] != NULL; ++postIdx) { - doReplace(tmp1, (size_t)2048, tmp2, - ".PRE.", pre_substitution_arr[preIdx], - ".POST.", post_substitution_arr[postIdx]); + for (preIdx = 0; pre_substitution_arr[preIdx] != NULL; ++preIdx) + { + for (postIdx = 0; post_substitution_arr[postIdx] != NULL; ++postIdx) + { + doReplace(tmp1, (size_t)2048, tmp2, ".PRE.", + pre_substitution_arr[preIdx], ".POST.", + post_substitution_arr[postIdx]); - result = test_vec_internal(deviceID, context, queue, tmp1, - "test_vec_align_packed_struct", - 512, pre_align_arr[preIdx], - type_multiple_pre_align_arr[preIdx], - post_align_arr[postIdx], - type_multiple_post_align_arr[postIdx]); - if(result != 0) { + result = test_vec_internal( + deviceID, context, queue, tmp1, "test_vec_align_packed_struct", + 512, pre_align_arr[preIdx], type_multiple_pre_align_arr[preIdx], + post_align_arr[postIdx], type_multiple_post_align_arr[postIdx]); + if (result != 0) + { return result; } } } log_info("testing __local\n"); - doReplace(tmp2, (size_t)2048, patterns[2], - ".SRC_SCOPE.", "__local", + doReplace(tmp2, (size_t)2048, patterns[2], ".SRC_SCOPE.", "__local", ".DST_SCOPE.", "__global"); // - for(preIdx = 0; pre_substitution_arr[preIdx] != NULL; ++preIdx) { - for(postIdx = 0; post_substitution_arr[postIdx] != NULL; ++postIdx) { - doReplace(tmp1, (size_t)2048, tmp2, - ".PRE.", pre_substitution_arr[preIdx], - ".POST.", post_substitution_arr[postIdx]); + for (preIdx = 0; pre_substitution_arr[preIdx] != NULL; ++preIdx) + { + for (postIdx = 0; post_substitution_arr[postIdx] != NULL; ++postIdx) + { + doReplace(tmp1, (size_t)2048, tmp2, ".PRE.", + pre_substitution_arr[preIdx], ".POST.", + post_substitution_arr[postIdx]); - result = test_vec_internal(deviceID, context, queue, tmp1, - "test_vec_align_packed_struct", - 512, pre_align_arr[preIdx], - type_multiple_pre_align_arr[preIdx], - post_align_arr[postIdx], - type_multiple_post_align_arr[postIdx]); - if (result != 0) { + result = test_vec_internal( + deviceID, context, queue, tmp1, "test_vec_align_packed_struct", + 512, pre_align_arr[preIdx], type_multiple_pre_align_arr[preIdx], + post_align_arr[postIdx], type_multiple_post_align_arr[postIdx]); + if (result != 0) + { return result; } } @@ -483,7 +476,8 @@ int test_vec_align_packed_struct(cl_device_id deviceID, cl_context context, cl_c return 0; } -int test_vec_align_struct_arr(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) +int test_vec_align_struct_arr(cl_device_id deviceID, cl_context context, + cl_command_queue queue, int num_elements) { char tmp1[2048], tmp2[2048]; int result = 0; @@ -491,20 +485,22 @@ int test_vec_align_struct_arr(cl_device_id deviceID, cl_context context, cl_comm log_info("testing __global\n"); - doReplace(tmp2, (size_t)2048, patterns[3], - ".SRC_SCOPE.", "__global", + doReplace(tmp2, (size_t)2048, patterns[3], ".SRC_SCOPE.", "__global", ".DST_SCOPE.", "__global"); // - for(preIdx = 0; pre_substitution_arr[preIdx] != NULL; ++preIdx) { - for(postIdx = 0; post_substitution_arr[postIdx] != NULL; ++postIdx) { - doReplace(tmp1, (size_t)2048, tmp2, - ".PRE.", pre_substitution_arr[preIdx], - ".POST.", post_substitution_arr[postIdx]); + for (preIdx = 0; pre_substitution_arr[preIdx] != NULL; ++preIdx) + { + for (postIdx = 0; post_substitution_arr[postIdx] != NULL; ++postIdx) + { + doReplace(tmp1, (size_t)2048, tmp2, ".PRE.", + pre_substitution_arr[preIdx], ".POST.", + post_substitution_arr[postIdx]); result = test_vec_internal(deviceID, context, queue, tmp1, - "test_vec_align_struct_arr", - BUFFER_SIZE, 0, 0, 0, 0); - if(result != 0) { + "test_vec_align_struct_arr", BUFFER_SIZE, + 0, 0, 0, 0); + if (result != 0) + { return result; } } @@ -512,7 +508,8 @@ int test_vec_align_struct_arr(cl_device_id deviceID, cl_context context, cl_comm return 0; } -int test_vec_align_packed_struct_arr(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) +int test_vec_align_packed_struct_arr(cl_device_id deviceID, cl_context context, + cl_command_queue queue, int num_elements) { char tmp1[2048], tmp2[2048]; int result = 0; @@ -520,26 +517,24 @@ int test_vec_align_packed_struct_arr(cl_device_id deviceID, cl_context context, log_info("Testing __global\n"); - doReplace(tmp2, (size_t)2048, patterns[4], - ".SRC_SCOPE.", "__global", + doReplace(tmp2, (size_t)2048, patterns[4], ".SRC_SCOPE.", "__global", ".DST_SCOPE.", "__global"); // - for(preIdx = 0; pre_substitution_arr[preIdx] != NULL; ++preIdx) { - for(postIdx = 0; post_substitution_arr[postIdx] != NULL; ++postIdx) { - doReplace(tmp1, (size_t)2048, tmp2, - ".PRE.", pre_substitution_arr[preIdx], - ".POST.", post_substitution_arr[postIdx]); + for (preIdx = 0; pre_substitution_arr[preIdx] != NULL; ++preIdx) + { + for (postIdx = 0; post_substitution_arr[postIdx] != NULL; ++postIdx) + { + doReplace(tmp1, (size_t)2048, tmp2, ".PRE.", + pre_substitution_arr[preIdx], ".POST.", + post_substitution_arr[postIdx]); - result = test_vec_internal(deviceID, context, queue, tmp1, - "test_vec_align_packed_struct_arr", - BUFFER_SIZE, pre_align_arr[preIdx], - type_multiple_pre_align_arr[preIdx], - post_align_arr[postIdx], - type_multiple_post_align_arr[postIdx]); - if(result != 0) - return result; + result = test_vec_internal( + deviceID, context, queue, tmp1, + "test_vec_align_packed_struct_arr", BUFFER_SIZE, + pre_align_arr[preIdx], type_multiple_pre_align_arr[preIdx], + post_align_arr[postIdx], type_multiple_post_align_arr[postIdx]); + if (result != 0) return result; } } return 0; } - diff --git a/test_conformance/vec_align/type_replacer.cpp b/test_conformance/vectors/type_replacer.cpp similarity index 62% rename from test_conformance/vec_align/type_replacer.cpp rename to test_conformance/vectors/type_replacer.cpp index 74967b2c..39c6194e 100644 --- a/test_conformance/vec_align/type_replacer.cpp +++ b/test_conformance/vectors/type_replacer.cpp @@ -1,6 +1,6 @@ // // 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 @@ -18,15 +18,15 @@ #include #endif // !_MSC_VER -size_t doReplace(char * dest, size_t destLength, const char * source, - const char * stringToReplace1, const char * replaceWith1, - const char * stringToReplace2, const char * replaceWith2) +size_t doReplace(char* dest, size_t destLength, const char* source, + const char* stringToReplace1, const char* replaceWith1, + const char* stringToReplace2, const char* replaceWith2) { size_t copyCount = 0; - const char * sourcePtr = source; - char * destPtr = dest; - const char * ptr1; - const char * ptr2; + const char* sourcePtr = source; + char* destPtr = dest; + const char* ptr1; + const char* ptr2; size_t nJump; size_t len1, len2; size_t lenReplace1, lenReplace2; @@ -34,14 +34,18 @@ size_t doReplace(char * dest, size_t destLength, const char * source, len2 = strlen(stringToReplace2); lenReplace1 = strlen(replaceWith1); lenReplace2 = strlen(replaceWith2); - for(;copyCount < destLength && *sourcePtr; ) + for (; copyCount < destLength && *sourcePtr;) { ptr1 = strstr(sourcePtr, stringToReplace1); ptr2 = strstr(sourcePtr, stringToReplace2); - if(ptr1 != NULL && (ptr2 == NULL || ptr2 > ptr1)) + if (ptr1 != NULL && (ptr2 == NULL || ptr2 > ptr1)) { - nJump = ptr1-sourcePtr; - if(((uintptr_t)ptr1-(uintptr_t)sourcePtr) > destLength-copyCount) { return -1; } + nJump = ptr1 - sourcePtr; + if (((uintptr_t)ptr1 - (uintptr_t)sourcePtr) + > destLength - copyCount) + { + return -1; + } copyCount += nJump; strncpy(destPtr, sourcePtr, nJump); destPtr += nJump; @@ -49,10 +53,13 @@ size_t doReplace(char * dest, size_t destLength, const char * source, strcpy(destPtr, replaceWith1); destPtr += lenReplace1; } - else if(ptr2 != NULL && (ptr1 == NULL || ptr1 >= ptr2)) + else if (ptr2 != NULL && (ptr1 == NULL || ptr1 >= ptr2)) { - nJump = ptr2-sourcePtr; - if(nJump > destLength-copyCount) { return -2; } + nJump = ptr2 - sourcePtr; + if (nJump > destLength - copyCount) + { + return -2; + } copyCount += nJump; strncpy(destPtr, sourcePtr, nJump); destPtr += nJump; @@ -63,7 +70,10 @@ size_t doReplace(char * dest, size_t destLength, const char * source, else { nJump = strlen(sourcePtr); - if(nJump > destLength-copyCount) { return -3; } + if (nJump > destLength - copyCount) + { + return -3; + } copyCount += nJump; strcpy(destPtr, sourcePtr); destPtr += nJump; @@ -74,25 +84,29 @@ size_t doReplace(char * dest, size_t destLength, const char * source, return copyCount; } -size_t doSingleReplace(char * dest, size_t destLength, const char * source, - const char * stringToReplace, const char * replaceWith) +size_t doSingleReplace(char* dest, size_t destLength, const char* source, + const char* stringToReplace, const char* replaceWith) { size_t copyCount = 0; - const char * sourcePtr = source; - char * destPtr = dest; - const char * ptr; + const char* sourcePtr = source; + char* destPtr = dest; + const char* ptr; size_t nJump; size_t len; size_t lenReplace; len = strlen(stringToReplace); lenReplace = strlen(replaceWith); - for(;copyCount < destLength && *sourcePtr; ) + for (; copyCount < destLength && *sourcePtr;) { ptr = strstr(sourcePtr, stringToReplace); - if(ptr != NULL) + if (ptr != NULL) { - nJump = ptr-sourcePtr; - if(((uintptr_t)ptr-(uintptr_t)sourcePtr) > destLength-copyCount) { return -1; } + nJump = ptr - sourcePtr; + if (((uintptr_t)ptr - (uintptr_t)sourcePtr) + > destLength - copyCount) + { + return -1; + } copyCount += nJump; strncpy(destPtr, sourcePtr, nJump); destPtr += nJump; @@ -103,7 +117,10 @@ size_t doSingleReplace(char * dest, size_t destLength, const char * source, else { nJump = strlen(sourcePtr); - if(nJump > destLength-copyCount) { return -3; } + if (nJump > destLength - copyCount) + { + return -3; + } copyCount += nJump; strcpy(destPtr, sourcePtr); destPtr += nJump; diff --git a/test_conformance/vec_step/type_replacer.h b/test_conformance/vectors/type_replacer.h similarity index 62% rename from test_conformance/vec_step/type_replacer.h rename to test_conformance/vectors/type_replacer.h index f50b08d7..d7eb7efc 100644 --- a/test_conformance/vec_step/type_replacer.h +++ b/test_conformance/vectors/type_replacer.h @@ -1,6 +1,6 @@ // // 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 @@ -15,9 +15,9 @@ // #include -size_t doReplace(char * dest, size_t destLength, const char * source, - const char * stringToReplace1, const char * replaceWith1, - const char * stringToReplace2, const char * replaceWith2); +size_t doReplace(char* dest, size_t destLength, const char* source, + const char* stringToReplace1, const char* replaceWith1, + const char* stringToReplace2, const char* replaceWith2); -size_t doSingleReplace(char * dest, size_t destLength, const char * source, - const char * stringToReplace, const char * replaceWith); +size_t doSingleReplace(char* dest, size_t destLength, const char* source, + const char* stringToReplace, const char* replaceWith);