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 <kevin.petit@arm.com>

* format fixes

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 )

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
 

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

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
 

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

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 .
 

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
 

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
 

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
 

test_conformance/vec_align/globals.cpp

@@ -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
-    };
-

test_conformance/vec_align/main.cpp

@@ -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 <stdio.h>
-#include <string.h>
-#include "procs.h"
-#include "harness/testHarness.h"
-
-#if !defined(_WIN32)
-#include <unistd.h>
-#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 );
-}
-

test_conformance/vec_align/structs.h

@@ -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);

test_conformance/vec_align/type_replacer.h

@@ -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 <stdlib.h>
-
-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);

test_conformance/vec_step/CMakeLists.txt

@@ -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)

test_conformance/vec_step/defines.h

@@ -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[];

test_conformance/vec_step/globals.cpp

@@ -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
-    };
-

test_conformance/vec_step/procs.h

@@ -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);

test_conformance/vec_step/structs.cpp

@@ -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;
-}

test_conformance/vec_step/structs.h

@@ -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);

test_conformance/vec_step/testBase.h

@@ -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 <stdio.h>
-#include <string.h>
-#include <sys/types.h>
-#include <sys/stat.h>
-
-#include "procs.h"
-
-#endif // _testBase_h

test_conformance/vec_step/type_replacer.cpp

@@ -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 <string.h>
-#if !defined(_MSC_VER)
-#include <stdint.h>
-#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;
-}

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
 )

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 };

test_conformance/vectors/main.cpp

@@ -29,6 +29,11 @@ test_definition test_list[] = {
     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);
@@ -37,4 +42,3 @@ int main(int argc, const char *argv[])
 {
     return runTestHarness(argc, argv, test_num, test_list, false, false, 0);
 }
-

test_conformance/vectors/procs.h

@@ -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);

test_conformance/vectors/structs.cpp

@@ -33,7 +33,8 @@
  */
 
 
-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));
 #if DEBUG_MEM_ALLOC
@@ -44,7 +45,8 @@ 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;
 }
 
@@ -64,26 +66,24 @@ int clStateMakeProgram(clState * pState, const char * prog,
 {
     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);
+    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;
@@ -95,17 +95,20 @@ void clStateDestroyProgramAndKernel(clState * pState)
 #if DEBUG_MEM_ALLOC
     log_info("destroy program and kernel\n");
 #endif
-    if(pState->m_kernel != NULL) {
+    if (pState->m_kernel != NULL)
+    {
         clReleaseKernel(pState->m_kernel);
         pState->m_kernel = NULL;
     }
-    if(pState->m_program != NULL) {
+    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));
 #if DEBUG_MEM_ALLOC
@@ -133,11 +136,8 @@ bufferStruct * newBufferStruct(size_t inSize, size_t outSize, clState * pClState
     log_info("clCreateBuffer %x\n", pResult->m_inBuffer);
 #endif
 
-    pResult->m_outBuffer = clCreateBuffer( pClState->m_context,
-                                          CL_MEM_WRITE_ONLY,
-                                          outSize,
-                                          NULL,
-                                          &error );
+    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);
@@ -152,31 +152,36 @@ bufferStruct * newBufferStruct(size_t inSize, size_t outSize, clState * pClState
     return pResult;
 }
 
-bufferStruct * destroyBufferStruct(bufferStruct * destroyMe, clState * pClState) {
+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
@@ -193,8 +198,7 @@ bufferStruct * destroyBufferStruct(bufferStruct * destroyMe, clState * pClState)
 }
 
 void initContents(bufferStruct *pBufferStruct, clState *pClState,
-                  size_t typeSize,
-                  size_t countIn, size_t countOut )
+                  size_t typeSize, size_t countIn, size_t countOut)
 {
     size_t i;
 
@@ -219,18 +223,22 @@ void initContents(bufferStruct * pBufferStruct, clState * pClState,
             break;
         }
         case 4: {
-            if (!g_wimpyMode) {
+            if (!g_wimpyMode)
+            {
                 uint32_t *ui = (uint32_t *)(pBufferStruct->m_pIn);
-                for (i=0; i < countIn; ++i) {
+                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
+            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) {
+                for (i = 0; i < countIn; ++i, ++start)
+                {
                     ui[i] = (int32_t)start * sign;
                     sign = sign * -1;
                 }
@@ -242,7 +250,8 @@ void initContents(bufferStruct * pBufferStruct, clState * pClState,
             // 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) {
+            for (i = 0; i < countIn; ++i, ++start)
+            {
                 ll[i] = start * sign;
                 sign = sign * -1;
             }
@@ -275,7 +284,8 @@ int pushArgs(bufferStruct * pBufferStruct, clState * pClState)
         pBufferStruct->m_bufferUploaded = true;
     }
 
-    err = clSetKernelArg(pClState->m_kernel, 0,
+    err = clSetKernelArg(
+        pClState->m_kernel, 0,
         sizeof(pBufferStruct->m_inBuffer), // pBufferStruct->m_bufSizeIn,
         &(pBufferStruct->m_inBuffer));
 #if DEBUG_MEM_ALLOC
@@ -287,7 +297,8 @@ int pushArgs(bufferStruct * pBufferStruct, clState * pClState)
         return -1;
     }
 
-    err = clSetKernelArg(pClState->m_kernel, 1,
+    err = clSetKernelArg(
+        pClState->m_kernel, 1,
         sizeof(pBufferStruct->m_outBuffer), // pBufferStruct->m_bufSizeOut,
         &(pBufferStruct->m_outBuffer));
     if (err != CL_SUCCESS)
@@ -319,7 +330,7 @@ int retrieveResults(bufferStruct * pBufferStruct, clState * pClState)
 
 // vecSizeIdx indexes into g_arrVecAlignMasks, g_arrVecSizeNames
 // and g_arrVecSizes
-int checkCorrectness(bufferStruct * pBufferStruct, clState * pClState,
+int checkCorrectnessAlign(bufferStruct *pBufferStruct, clState *pClState,
                           size_t minAlign)
 {
     size_t i;
@@ -329,9 +340,7 @@ int checkCorrectness(bufferStruct * pBufferStruct, clState * pClState,
         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,6 +354,27 @@ 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
@@ -357,9 +387,9 @@ int checkPackedCorrectness(bufferStruct * pBufferStruct, clState * pClState,
     {
         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;
         }

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);

test_conformance/vectors/test_step.cpp

@@ -35,8 +35,9 @@
  */
 
 
-
-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;
@@ -47,7 +48,8 @@ int test_step_internal(cl_device_id deviceID, cl_context context, cl_command_que
     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;
@@ -55,7 +57,8 @@ int test_step_internal(cl_device_id deviceID, cl_context context, cl_command_que
 
     // 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");
@@ -77,31 +80,34 @@ int test_step_internal(cl_device_id deviceID, cl_context context, cl_command_que
 
         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 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" );
+                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)
         {
-            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);
@@ -111,8 +117,8 @@ int test_step_internal(cl_device_id deviceID, cl_context context, cl_command_que
             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;
@@ -121,8 +127,7 @@ int test_step_internal(cl_device_id deviceID, cl_context context, cl_command_que
             err = pushArgs(pBuffers, pClState);
             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);
@@ -134,9 +139,8 @@ int test_step_internal(cl_device_id deviceID, cl_context context, cl_command_que
             err = runKernel(pClState, 1024);
             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);
@@ -146,8 +150,7 @@ int test_step_internal(cl_device_id deviceID, cl_context context, cl_command_que
             err = retrieveResults(pBuffers, pClState);
             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,
+            err = checkCorrectnessStep(pBuffers, pClState,
                                        g_arrTypeSizes[typeIdx],
                                        g_arrVecSizes[vecSizeIdx]);
 
             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");

test_conformance/vectors/test_vec_align.cpp

@@ -36,7 +36,8 @@ size_t get_align(size_t vecSize)
     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)
+                      const char* testName, size_t bufSize, size_t preSize,
+                      size_t typeMultiplePreSize, size_t postSize,
+                      size_t typeMultiplePostSize)
 {
     int err;
     int typeIdx, vecSizeIdx;
@@ -87,10 +88,11 @@ 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);
+    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;
@@ -100,23 +102,28 @@ int test_vec_internal(cl_device_id deviceID, cl_context context,
     {
 
         // 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]);
@@ -124,20 +131,20 @@ int test_vec_internal(cl_device_id deviceID, cl_context context,
         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);
@@ -148,18 +155,19 @@ int test_vec_internal(cl_device_id deviceID, cl_context context,
             // fflush(stdout);
 
             err = clStateMakeProgram(pClState, srcBuffer, testName);
-            if (err) {
-                vlog_error("%s: Error compiling \"\n%s\n\"",
-                           testName, srcBuffer);
+            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)
             {
                 // 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[] = {
-    "",
+const char* pre_substitution_arr[] = { "",
                                        "char c;\n",
                                        "short3 s;",
                                        ".TYPE.3 tPre;\n",
                                        ".TYPE. arrPre[5];\n",
                                        ".TYPE. arrPre[12];\n",
-    NULL
-};
+                                       NULL };
 
 
 // alignments of everything in pre_substitution_arr as raw alignments
 // 0 if such a thing is meaningless
-size_t pre_align_arr[] = {
-    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
-};
+                           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[] = {
-    "",
+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
-};
+                                        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),
+size_t post_align_arr[] = { 0, sizeof(cl_char),
                             0, // taken care of in type_multiple_post_align_arr
-    0,
-    0,
-    0
-};
+                            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;
 }
-

test_conformance/vectors/type_replacer.cpp

@@ -41,7 +41,11 @@ size_t doReplace(char * dest, size_t destLength, const char * source,
         if (ptr1 != NULL && (ptr2 == NULL || ptr2 > ptr1))
         {
             nJump = ptr1 - sourcePtr;
-            if(((uintptr_t)ptr1-(uintptr_t)sourcePtr) > destLength-copyCount) { return -1; }
+            if (((uintptr_t)ptr1 - (uintptr_t)sourcePtr)
+                > destLength - copyCount)
+            {
+                return -1;
+            }
             copyCount += nJump;
             strncpy(destPtr, sourcePtr, nJump);
             destPtr += nJump;
@@ -52,7 +56,10 @@ size_t doReplace(char * dest, size_t destLength, const char * source,
         else if (ptr2 != NULL && (ptr1 == NULL || ptr1 >= ptr2))
         {
             nJump = ptr2 - sourcePtr;
-            if(nJump > destLength-copyCount) { return -2; }
+            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;
@@ -92,7 +102,11 @@ size_t doSingleReplace(char * dest, size_t destLength, const char * source,
         if (ptr != NULL)
         {
             nJump = ptr - sourcePtr;
-            if(((uintptr_t)ptr-(uintptr_t)sourcePtr) > destLength-copyCount) { return -1; }
+            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;