ahmed/OpenCL-CTS
1
0
Fork 0
mirror of https://github.com/KhronosGroup/OpenCL-CTS.git synced 2026-03-19 06:09:01 +00:00
Code Issues Packages Projects Releases Wiki Activity

Merge branch 'main' into cl_khr_unified_svm

Browse Source
This commit is contained in:
Ben Ashbaugh
2025-03-18 11:57:53 -07:00
parent 3f15fcab9e e88e5be93e
commit 09199a79eb
98 changed files with 1400 additions and 2298 deletions
Download Patch File Download Diff File Expand all files Collapse all files

2
.github/workflows/presubmit.yml vendored
View File

@@ -73,7 +73,7 @@ jobs:
- name: Install Vulkan SDK
uses: humbletim/install-vulkan-sdk@main
with:
version: 1.3.275.0
version: 1.4.309.0
cache: true
- name: Install Android NDK
if: ${{ matrix.arch == 'android-arm' || matrix.arch == 'android-aarch64' }}

7
test_conformance/api/test_sub_group_dispatch.cpp
View File

@@ -102,10 +102,13 @@ REGISTER_TEST_VERSION(sub_group_dispatch, Version(2, 1))
out = clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR, sizeof(size_t), NULL, &error);
test_error(error, "clCreateBuffer failed");
error = clGetDeviceInfo(device, CL_DEVICE_MAX_WORK_GROUP_SIZE,
sizeof(size_t), &max_local, NULL);
size_t max_work_item_sizes[3] = {};
error = clGetDeviceInfo(device, CL_DEVICE_MAX_WORK_ITEM_SIZES,
sizeof(max_work_item_sizes), &max_work_item_sizes,
nullptr);
test_error(error, "clGetDeviceInfo failed");
max_local = max_work_item_sizes[0];
error = clGetDeviceInfo(device, CL_DEVICE_PLATFORM, sizeof(platform),
(void *)&platform, NULL);

151
test_conformance/basic/main.cpp
View File

@@ -14,151 +14,12 @@
// limitations under the License.
//
#include "harness/compat.h"
#if !defined(_WIN32)
#include <unistd.h>
#endif
#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#include "harness/deviceInfo.h"
#include "harness/kernelHelpers.h"
#include "harness/testHarness.h"
#include <CL/cl_half.h>
#include "harness/testHarness.h"
#include "procs.h"
test_definition test_list[] = {
ADD_TEST(hostptr),
ADD_TEST(fpmath),
ADD_TEST(intmath_int),
ADD_TEST(intmath_int2),
ADD_TEST(intmath_int4),
ADD_TEST(intmath_long),
ADD_TEST(intmath_long2),
ADD_TEST(intmath_long4),
ADD_TEST(hiloeo),
ADD_TEST(if),
ADD_TEST(sizeof),
ADD_TEST(loop),
ADD_TEST(pointer_cast),
ADD_TEST(local_arg_def),
ADD_TEST(local_kernel_def),
ADD_TEST(local_kernel_scope),
ADD_TEST(constant),
ADD_TEST(constant_source),
ADD_TEST(readimage),
ADD_TEST(readimage_int16),
ADD_TEST(readimage_fp32),
ADD_TEST(writeimage),
ADD_TEST(writeimage_int16),
ADD_TEST(writeimage_fp32),
ADD_TEST(mri_one),
ADD_TEST(mri_multiple),
ADD_TEST(image_r8),
ADD_TEST(barrier),
ADD_TEST_VERSION(wg_barrier, Version(2, 0)),
ADD_TEST(int2fp),
ADD_TEST(fp2int),
ADD_TEST(imagereadwrite),
ADD_TEST(imagereadwrite3d),
ADD_TEST(readimage3d),
ADD_TEST(readimage3d_int16),
ADD_TEST(readimage3d_fp32),
ADD_TEST(bufferreadwriterect),
ADD_TEST(arrayreadwrite),
ADD_TEST(arraycopy),
ADD_TEST(imagearraycopy),
ADD_TEST(imagearraycopy3d),
ADD_TEST(imagecopy),
ADD_TEST(imagecopy3d),
ADD_TEST(imagerandomcopy),
ADD_TEST(arrayimagecopy),
ADD_TEST(arrayimagecopy3d),
ADD_TEST(imagenpot),
ADD_TEST(vload_global),
ADD_TEST(vload_local),
ADD_TEST(vload_constant),
ADD_TEST(vload_private),
ADD_TEST(vstore_global),
ADD_TEST(vstore_local),
ADD_TEST(vstore_private),
ADD_TEST(createkernelsinprogram),
ADD_TEST(imagedim_pow2),
ADD_TEST(imagedim_non_pow2),
ADD_TEST(image_param),
ADD_TEST(image_multipass_integer_coord),
ADD_TEST(image_multipass_float_coord),
ADD_TEST(explicit_s2v),
ADD_TEST(enqueue_map_buffer),
ADD_TEST(enqueue_map_image),
ADD_TEST(work_item_functions),
ADD_TEST(work_item_functions_out_of_range),
ADD_TEST(work_item_functions_out_of_range_hardcoded),
ADD_TEST(astype),
ADD_TEST(async_copy_global_to_local),
ADD_TEST(async_copy_local_to_global),
ADD_TEST(async_strided_copy_global_to_local),
ADD_TEST(async_strided_copy_local_to_global),
ADD_TEST(async_copy_global_to_local2D),
ADD_TEST(async_copy_local_to_global2D),
ADD_TEST(async_copy_global_to_local3D),
ADD_TEST(async_copy_local_to_global3D),
ADD_TEST(async_work_group_copy_fence_import_after_export_aliased_local),
ADD_TEST(async_work_group_copy_fence_import_after_export_aliased_global),
ADD_TEST(
async_work_group_copy_fence_import_after_export_aliased_global_and_local),
ADD_TEST(async_work_group_copy_fence_export_after_import_aliased_local),
ADD_TEST(async_work_group_copy_fence_export_after_import_aliased_global),
ADD_TEST(
async_work_group_copy_fence_export_after_import_aliased_global_and_local),
ADD_TEST(prefetch),
ADD_TEST(kernel_call_kernel_function),
ADD_TEST(host_numeric_constants),
ADD_TEST(kernel_numeric_constants),
ADD_TEST(kernel_limit_constants),
ADD_TEST(kernel_preprocessor_macros),
ADD_TEST(parameter_types),
ADD_TEST(vector_creation),
ADD_TEST(vector_swizzle),
ADD_TEST(vec_type_hint),
ADD_TEST(kernel_memory_alignment_local),
ADD_TEST(kernel_memory_alignment_global),
ADD_TEST(kernel_memory_alignment_constant),
ADD_TEST(kernel_memory_alignment_private),
ADD_TEST_VERSION(progvar_prog_scope_misc, Version(2, 0)),
ADD_TEST_VERSION(progvar_prog_scope_uninit, Version(2, 0)),
ADD_TEST_VERSION(progvar_prog_scope_init, Version(2, 0)),
ADD_TEST_VERSION(progvar_func_scope, Version(2, 0)),
ADD_TEST(global_work_offsets),
ADD_TEST(get_global_offset),
ADD_TEST_VERSION(global_linear_id, Version(2, 0)),
ADD_TEST_VERSION(local_linear_id, Version(2, 0)),
ADD_TEST_VERSION(enqueued_local_size, Version(2, 0)),
ADD_TEST(simple_read_image_pitch),
ADD_TEST(simple_write_image_pitch),
#if defined(__APPLE__)
ADD_TEST(queue_priority),
#endif
ADD_TEST_VERSION(get_linear_ids, Version(2, 0)),
ADD_TEST_VERSION(rw_image_access_qualifier, Version(2, 0)),
};
const int test_num = ARRAY_SIZE( test_list );
cl_half_rounding_mode halfRoundingMode = CL_HALF_RTE;
test_status InitCL(cl_device_id device)
@@ -187,7 +48,7 @@ test_status InitCL(cl_device_id device)
int main(int argc, const char *argv[])
{
return runTestHarnessWithCheck(argc, argv, test_num, test_list, false, 0,
InitCL);
return runTestHarnessWithCheck(
argc, argv, test_registry::getInstance().num_tests(),
test_registry::getInstance().definitions(), false, 0, InitCL);
}

198
test_conformance/basic/procs.h
View File

@@ -1,198 +0,0 @@
//
// Copyright (c) 2023 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/kernelHelpers.h"
#include "harness/testHarness.h"
#include "harness/errorHelpers.h"
#include "harness/typeWrappers.h"
#include "harness/conversions.h"
#include "harness/rounding_mode.h"
extern int test_hostptr(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_fpmath(cl_device_id deviceID, cl_context context,
cl_command_queue queue, int num_elements);
extern int test_intmath_int(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_intmath_int2(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_intmath_int4(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_intmath_long(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_intmath_long2(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_intmath_long4(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_hiloeo(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_if(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_sizeof(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_loop(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_pointer_cast(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_local_arg_def(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_local_kernel_def(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_local_kernel_scope(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_constant(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_constant_source(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_readimage(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_readimage_int16(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_readimage_fp32(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_writeimage(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_writeimage_int16(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_writeimage_fp32(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_mri_one(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_mri_multiple(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_image_r8(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_simplebarrier(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_barrier(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_wg_barrier(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_int2fp(cl_device_id deviceID, cl_context context,
cl_command_queue queue, int num_elements);
extern int test_fp2int(cl_device_id deviceID, cl_context context,
cl_command_queue queue, int num_elements);
extern int test_imagearraycopy(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_imagearraycopy3d(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_imagereadwrite(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_imagereadwrite3d(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_readimage3d(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_readimage3d_int16(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_readimage3d_fp32(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_bufferreadwriterect(cl_device_id device, cl_context context, cl_command_queue queue_, int num_elements);
extern int test_imagecopy(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_imagecopy3d(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_imagerandomcopy(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_arraycopy(cl_device_id device, cl_context context, cl_command_queue queue, int n_elems);
extern int test_arrayimagecopy(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_arrayimagecopy3d(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_imagenpot(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_sampler_float(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_sampler_int(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_createkernelsinprogram(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_single_large_allocation(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_multiple_max_allocation(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_arrayreadwrite(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_imagedim_pow2(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_imagedim_non_pow2(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_image_param(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_image_multipass_integer_coord(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_image_multipass_float_coord(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_vload_global(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_vload_local(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_vload_constant(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_vload_private(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_vstore_global(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_vstore_local(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_vstore_private(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_explicit_s2v(cl_device_id deviceID, cl_context context,
cl_command_queue queue, int num_elements);
extern int test_enqueue_map_buffer(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_enqueue_map_image(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_work_item_functions(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_work_item_functions_out_of_range(cl_device_id deviceID,
cl_context context,
cl_command_queue queue,
int num_elements);
extern int test_work_item_functions_out_of_range_hardcoded(
cl_device_id deviceID, cl_context context, cl_command_queue queue,
int num_elements);
extern int test_astype(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_native_kernel(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_async_copy_global_to_local(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_async_copy_local_to_global(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_async_strided_copy_global_to_local(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_async_strided_copy_local_to_global(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_async_copy_global_to_local2D(cl_device_id deviceID,
cl_context context,
cl_command_queue queue,
int num_elements);
extern int test_async_copy_local_to_global2D(cl_device_id deviceID,
cl_context context,
cl_command_queue queue,
int num_elements);
extern int test_async_copy_global_to_local3D(cl_device_id deviceID,
cl_context context,
cl_command_queue queue,
int num_elements);
extern int test_async_copy_local_to_global3D(cl_device_id deviceID,
cl_context context,
cl_command_queue queue,
int num_elements);
extern int test_async_work_group_copy_fence_import_after_export_aliased_local(
cl_device_id deviceID, cl_context context, cl_command_queue queue,
int num_elements);
extern int test_async_work_group_copy_fence_import_after_export_aliased_global(
cl_device_id deviceID, cl_context context, cl_command_queue queue,
int num_elements);
extern int
test_async_work_group_copy_fence_import_after_export_aliased_global_and_local(
cl_device_id deviceID, cl_context context, cl_command_queue queue,
int num_elements);
extern int test_async_work_group_copy_fence_export_after_import_aliased_local(
cl_device_id deviceID, cl_context context, cl_command_queue queue,
int num_elements);
extern int test_async_work_group_copy_fence_export_after_import_aliased_global(
cl_device_id deviceID, cl_context context, cl_command_queue queue,
int num_elements);
extern int
test_async_work_group_copy_fence_export_after_import_aliased_global_and_local(
cl_device_id deviceID, cl_context context, cl_command_queue queue,
int num_elements);
extern int test_prefetch(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_host_numeric_constants(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_kernel_numeric_constants(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_kernel_limit_constants(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_kernel_preprocessor_macros(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_kernel_call_kernel_function(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_parameter_types(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements);
extern int test_vector_creation(cl_device_id deviceID, cl_context context,
cl_command_queue queue, int num_elements);
extern int test_vector_swizzle(cl_device_id deviceID, cl_context context,
cl_command_queue queue, int num_elements);
extern int test_vec_type_hint(cl_device_id deviceID, cl_context context,
cl_command_queue queue, int num_elements);
extern int test_kernel_memory_alignment_local(cl_device_id device, cl_context context, cl_command_queue queue, int n_elems );
extern int test_kernel_memory_alignment_global(cl_device_id device, cl_context context, cl_command_queue queue, int n_elems );
extern int test_kernel_memory_alignment_constant(cl_device_id device, cl_context context, cl_command_queue queue, int n_elems );
extern int test_kernel_memory_alignment_private(cl_device_id device, cl_context context, cl_command_queue queue, int n_elems );
extern int test_progvar_prog_scope_misc(cl_device_id device, cl_context context, cl_command_queue queue, int n_elems );
extern int test_progvar_prog_scope_uninit(cl_device_id device, cl_context context, cl_command_queue queue, int n_elems );
extern int test_progvar_prog_scope_init(cl_device_id device, cl_context context, cl_command_queue queue, int n_elems );
extern int test_progvar_func_scope(cl_device_id device, cl_context context, cl_command_queue queue, int n_elems );
extern int test_global_work_offsets(cl_device_id device, cl_context context, cl_command_queue queue, int n_elems );
extern int test_get_global_offset(cl_device_id device, cl_context context, cl_command_queue queue, int n_elems );
extern int test_global_linear_id(cl_device_id device, cl_context context, cl_command_queue queue, int n_elems );
extern int test_local_linear_id(cl_device_id device, cl_context context, cl_command_queue queue, int n_elems );
extern int test_enqueued_local_size(cl_device_id device, cl_context context, cl_command_queue queue, int n_elems );
extern int test_simple_read_image_pitch(cl_device_id device, cl_context cl_context_, cl_command_queue q, int num_elements);
extern int test_simple_write_image_pitch(cl_device_id device, cl_context cl_context_, cl_command_queue q, int num_elements);
#if defined( __APPLE__ )
extern int test_queue_priority(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements);
#endif
extern int test_get_linear_ids(cl_device_id device, cl_context cl_context_, cl_command_queue q, int num_elements);
extern int test_rw_image_access_qualifier(cl_device_id device_id, cl_context context, cl_command_queue commands, int num_elements);

26
test_conformance/basic/testBase.h Normal file
View File

@@ -0,0 +1,26 @@
//
// Copyright (c) 2025 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 <CL/cl.h>
#include "harness/conversions.h"
#include "harness/testHarness.h"
#include "harness/typeWrappers.h"
#include "harness/rounding_mode.h"
#endif // _testBase_h

11
test_conformance/basic/test_arraycopy.cpp
View File

@@ -20,8 +20,7 @@
#include <sys/types.h>
#include <sys/stat.h>
#include "procs.h"
#include "testBase.h"
const char *copy_kernel_code =
"__kernel void test_copy(__global unsigned int *src, __global unsigned int *dst)\n"
@@ -31,14 +30,13 @@ const char *copy_kernel_code =
" dst[tid] = src[tid];\n"
"}\n";
int
test_arraycopy(cl_device_id device, cl_context context, cl_command_queue queue, int n_elems)
REGISTER_TEST(arraycopy)
{
cl_uint *input_ptr, *output_ptr;
cl_mem streams[4], results;
cl_program program;
cl_kernel kernel;
unsigned num_elements = 128 * 1024;
num_elements = 128 * 1024;
cl_uint num_copies = 1;
size_t delta_offset;
unsigned i;
@@ -202,6 +200,3 @@ test_arraycopy(cl_device_id device, cl_context context, cl_command_queue queue,
return err;
}

30
test_conformance/basic/test_arrayimagecopy.cpp
View File

@@ -23,19 +23,14 @@
#include <sys/stat.h>
#include <vector>
#include <memory>
#include "procs.h"
using test_function_t = int (*)(cl_device_id, cl_context, cl_command_queue,
cl_mem_flags, cl_mem_flags, cl_mem_object_type,
const cl_image_format *);
int test_arrayimagecopy_single_format(cl_device_id device, cl_context context,
cl_command_queue queue,
cl_mem_flags buffer_flags,
cl_mem_flags image_flags,
cl_mem_object_type image_type,
const cl_image_format *format)
static int test_arrayimagecopy_single_format(
cl_device_id device, cl_context context, cl_command_queue queue,
cl_mem_flags buffer_flags, cl_mem_flags image_flags,
cl_mem_object_type image_type, const cl_image_format *format)
{
std::unique_ptr<cl_uchar, decltype(&free)> bufptr{ nullptr, free },
imgptr{ nullptr, free };
@@ -153,11 +148,12 @@ int test_arrayimagecopy_single_format(cl_device_id device, cl_context context,
}
int test_arrayimagecommon(cl_device_id device, cl_context context,
cl_command_queue queue, cl_mem_flags buffer_flags,
cl_mem_flags image_flags,
cl_mem_object_type image_type,
test_function_t test_function)
static int test_arrayimagecommon(cl_device_id device, cl_context context,
cl_command_queue queue,
cl_mem_flags buffer_flags,
cl_mem_flags image_flags,
cl_mem_object_type image_type,
test_function_t test_function)
{
cl_int err;
cl_uint num_formats;
@@ -188,8 +184,7 @@ int test_arrayimagecommon(cl_device_id device, cl_context context,
return err;
}
int test_arrayimagecopy(cl_device_id device, cl_context context,
cl_command_queue queue, int num_elements)
REGISTER_TEST(arrayimagecopy)
{
PASSIVE_REQUIRE_IMAGE_SUPPORT(device)
@@ -199,8 +194,7 @@ int test_arrayimagecopy(cl_device_id device, cl_context context,
}
int test_arrayimagecopy3d(cl_device_id device, cl_context context,
cl_command_queue queue, int num_elements)
REGISTER_TEST(arrayimagecopy3d)
{
PASSIVE_REQUIRE_3D_IMAGE_SUPPORT(device)

10
test_conformance/basic/test_arrayreadwrite.cpp
View File

@@ -21,12 +21,9 @@
#include <sys/types.h>
#include <sys/stat.h>
#include "testBase.h"
#include "procs.h"
int
test_arrayreadwrite(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
REGISTER_TEST(arrayreadwrite)
{
cl_uint *inptr, *outptr;
cl_mem streams[1];
@@ -91,6 +88,3 @@ test_arrayreadwrite(cl_device_id device, cl_context context, cl_command_queue qu
return err;
}

27
test_conformance/basic/test_astype.cpp
View File

@@ -24,9 +24,6 @@
#include <sys/types.h>
#include <sys/stat.h>
#include <vector>
#include "procs.h"
// clang-format off
static char extension[128] = { 0 };
@@ -48,9 +45,10 @@ extension,
// clang-format on
int test_astype_set( cl_device_id device, cl_context context, cl_command_queue queue, ExplicitType inVecType, ExplicitType outVecType,
unsigned int vecSize, unsigned int outVecSize,
int numElements )
static int test_astype_set(cl_device_id device, cl_context context,
cl_command_queue queue, ExplicitType inVecType,
ExplicitType outVecType, unsigned int vecSize,
unsigned int outVecSize, int numElements)
{
int error;
@@ -179,7 +177,7 @@ int test_astype_set( cl_device_id device, cl_context context, cl_command_queue q
return 0;
}
int test_astype(cl_device_id device, cl_context context, cl_command_queue queue, int n_elems )
REGISTER_TEST(astype)
{
// Note: although casting to different vector element sizes that match the same size (i.e. short2 -> char4) is
// legal in OpenCL 1.0, the result is dependent on the device it runs on, which means there's no actual way
@@ -236,19 +234,24 @@ int test_astype(cl_device_id device, cl_context context, cl_command_queue queue,
{
continue;
}
error += test_astype_set( device, context, queue, vecTypes[ inTypeIdx ], vecTypes[ outTypeIdx ], vecSizes[ sizeIdx ], vecSizes[outSizeIdx], n_elems );
error += test_astype_set(
device, context, queue, vecTypes[inTypeIdx],
vecTypes[outTypeIdx], vecSizes[sizeIdx],
vecSizes[outSizeIdx], num_elements);
}
}
if(get_explicit_type_size(vecTypes[inTypeIdx]) ==
get_explicit_type_size(vecTypes[outTypeIdx])) {
// as_type3(vec4) allowed, as_type4(vec3) not allowed
error += test_astype_set( device, context, queue, vecTypes[ inTypeIdx ], vecTypes[ outTypeIdx ], 3, 4, n_elems );
error += test_astype_set( device, context, queue, vecTypes[ inTypeIdx ], vecTypes[ outTypeIdx ], 4, 3, n_elems );
error +=
test_astype_set(device, context, queue, vecTypes[inTypeIdx],
vecTypes[outTypeIdx], 3, 4, num_elements);
error +=
test_astype_set(device, context, queue, vecTypes[inTypeIdx],
vecTypes[outTypeIdx], 4, 3, num_elements);
}
}
}
return error;
}

33
test_conformance/basic/test_async_copy.cpp
View File

@@ -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
@@ -13,8 +13,6 @@
// See the License for the specific language governing permissions and
// limitations under the License.
//
#include "harness/compat.h"
#include <stdio.h>
#include <stdlib.h>
#include <string.h>
@@ -22,8 +20,7 @@
#include <sys/stat.h>
#include <vector>
#include "procs.h"
#include "harness/conversions.h"
#include "testBase.h"
static const char *async_global_to_local_kernel =
"%s\n" // optional pragma string
@@ -75,10 +72,9 @@ static const char *prefetch_kernel =
"}\n" ;
int test_copy(cl_device_id deviceID, cl_context context, cl_command_queue queue, const char *kernelCode,
ExplicitType vecType, int vecSize
)
static int test_copy(cl_device_id deviceID, cl_context context,
cl_command_queue queue, const char *kernelCode,
ExplicitType vecType, int vecSize)
{
int error;
clProgramWrapper program;
@@ -238,7 +234,9 @@ int test_copy(cl_device_id deviceID, cl_context context, cl_command_queue queue,
return failuresPrinted ? -1 : 0;
}
int test_copy_all_types(cl_device_id deviceID, cl_context context, cl_command_queue queue, const char *kernelCode) {
static int test_copy_all_types(cl_device_id deviceID, cl_context context,
cl_command_queue queue, const char *kernelCode)
{
const std::vector<ExplicitType> vecType = { kChar, kUChar, kShort, kUShort,
kInt, kUInt, kLong, kULong,
kFloat, kHalf, kDouble };
@@ -271,18 +269,19 @@ int test_copy_all_types(cl_device_id deviceID, cl_context context, cl_command_qu
return 0;
}
int test_async_copy_global_to_local(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
REGISTER_TEST(async_copy_global_to_local)
{
return test_copy_all_types( deviceID, context, queue, async_global_to_local_kernel );
return test_copy_all_types(device, context, queue,
async_global_to_local_kernel);
}
int test_async_copy_local_to_global(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
REGISTER_TEST(async_copy_local_to_global)
{
return test_copy_all_types( deviceID, context, queue, async_local_to_global_kernel );
return test_copy_all_types(device, context, queue,
async_local_to_global_kernel);
}
int test_prefetch(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
REGISTER_TEST(prefetch)
{
return test_copy_all_types( deviceID, context, queue, prefetch_kernel );
return test_copy_all_types(device, context, queue, prefetch_kernel);
}

30
test_conformance/basic/test_async_copy2D.cpp
View File

@@ -13,8 +13,6 @@
// See the License for the specific language governing permissions and
// limitations under the License.
//
#include "../../test_common/harness/compat.h"
#include <algorithm>
#include <stdio.h>
#include <stdlib.h>
@@ -22,8 +20,7 @@
#include <sys/stat.h>
#include <sys/types.h>
#include "../../test_common/harness/conversions.h"
#include "procs.h"
#include "testBase.h"
static const char *async_global_to_local_kernel2D = R"OpenCLC(
#pragma OPENCL EXTENSION cl_khr_extended_async_copies : enable
@@ -113,10 +110,11 @@ __kernel void test_fn(const __global VarSizeStruct *src, __global VarSizeStruct
};
)OpenCLC";
int test_copy2D(const cl_device_id deviceID, const cl_context context,
const cl_command_queue queue, const char *const kernelCode,
const size_t elementSize, const int srcMargin,
const int dstMargin, const bool localIsDst)
static int test_copy2D(const cl_device_id deviceID, const cl_context context,
const cl_command_queue queue,
const char *const kernelCode, const size_t elementSize,
const int srcMargin, const int dstMargin,
const bool localIsDst)
{
int error;
@@ -388,9 +386,9 @@ int test_copy2D(const cl_device_id deviceID, const cl_context context,
return failuresPrinted ? -1 : 0;
}
int test_copy2D_all_types(cl_device_id deviceID, cl_context context,
cl_command_queue queue, const char *kernelCode,
bool localIsDst)
static int test_copy2D_all_types(cl_device_id deviceID, cl_context context,
cl_command_queue queue, const char *kernelCode,
bool localIsDst)
{
const unsigned int elemSizes[] = { 1, 2, 3, 4, 5, 6, 7,
8, 13, 16, 32, 47, 64 };
@@ -430,16 +428,14 @@ int test_copy2D_all_types(cl_device_id deviceID, cl_context context,
return errors ? -1 : 0;
}
int test_async_copy_global_to_local2D(cl_device_id deviceID, cl_context context,
cl_command_queue queue, int num_elements)
REGISTER_TEST(async_copy_global_to_local2D)
{
return test_copy2D_all_types(deviceID, context, queue,
return test_copy2D_all_types(device, context, queue,
async_global_to_local_kernel2D, true);
}
int test_async_copy_local_to_global2D(cl_device_id deviceID, cl_context context,
cl_command_queue queue, int num_elements)
REGISTER_TEST(async_copy_local_to_global2D)
{
return test_copy2D_all_types(deviceID, context, queue,
return test_copy2D_all_types(device, context, queue,
async_local_to_global_kernel2D, false);
}

32
test_conformance/basic/test_async_copy3D.cpp
View File

@@ -13,8 +13,6 @@
// See the License for the specific language governing permissions and
// limitations under the License.
//
#include "../../test_common/harness/compat.h"
#include <algorithm>
#include <stdio.h>
#include <stdlib.h>
@@ -22,8 +20,7 @@
#include <sys/stat.h>
#include <sys/types.h>
#include "../../test_common/harness/conversions.h"
#include "procs.h"
#include "testBase.h"
static const char *async_global_to_local_kernel3D = R"OpenCLC(
#pragma OPENCL EXTENSION cl_khr_extended_async_copies : enable
@@ -124,11 +121,12 @@ __kernel void test_fn(const __global VarSizeStruct *src, __global VarSizeStruct
}
)OpenCLC";
int test_copy3D(const cl_device_id deviceID, const cl_context context,
const cl_command_queue queue, const char *const kernelCode,
const size_t elementSize, const int srcLineMargin,
const int dstLineMargin, const int srcPlaneMargin,
const int dstPlaneMargin, const bool localIsDst)
static int test_copy3D(const cl_device_id deviceID, const cl_context context,
const cl_command_queue queue,
const char *const kernelCode, const size_t elementSize,
const int srcLineMargin, const int dstLineMargin,
const int srcPlaneMargin, const int dstPlaneMargin,
const bool localIsDst)
{
int error;
@@ -445,9 +443,9 @@ int test_copy3D(const cl_device_id deviceID, const cl_context context,
return failuresPrinted ? -1 : 0;
}
int test_copy3D_all_types(cl_device_id deviceID, cl_context context,
cl_command_queue queue, const char *kernelCode,
bool localIsDst)
static int test_copy3D_all_types(cl_device_id deviceID, cl_context context,
cl_command_queue queue, const char *kernelCode,
bool localIsDst)
{
const unsigned int elemSizes[] = { 1, 2, 3, 4, 5, 6, 7,
8, 13, 16, 32, 47, 64 };
@@ -500,16 +498,14 @@ int test_copy3D_all_types(cl_device_id deviceID, cl_context context,
return 0;
}
int test_async_copy_global_to_local3D(cl_device_id deviceID, cl_context context,
cl_command_queue queue, int num_elements)
REGISTER_TEST(async_copy_global_to_local3D)
{
return test_copy3D_all_types(deviceID, context, queue,
return test_copy3D_all_types(device, context, queue,
async_global_to_local_kernel3D, true);
}
int test_async_copy_local_to_global3D(cl_device_id deviceID, cl_context context,
cl_command_queue queue, int num_elements)
REGISTER_TEST(async_copy_local_to_global3D)
{
return test_copy3D_all_types(deviceID, context, queue,
return test_copy3D_all_types(device, context, queue,
async_local_to_global_kernel3D, false);
}

62
test_conformance/basic/test_async_copy_fence.cpp
View File

@@ -13,16 +13,13 @@
// See the License for the specific language governing permissions and
// limitations under the License.
//
#include "../../test_common/harness/compat.h"
#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#include <sys/stat.h>
#include <sys/types.h>
#include "../../test_common/harness/conversions.h"
#include "procs.h"
#include "testBase.h"
static const char *import_after_export_aliased_local_kernel =
"#pragma OPENCL EXTENSION cl_khr_async_work_group_copy_fence : enable\n"
@@ -331,10 +328,11 @@ static const char *export_after_import_aliased_global_and_local_kernel =
" }\n"
"}\n";
int test_copy_fence(cl_device_id deviceID, cl_context context,
cl_command_queue queue, const char *kernelCode,
ExplicitType vecType, int vecSize, bool export_after_import,
bool aliased_local_mem, bool aliased_global_mem)
static int test_copy_fence(cl_device_id deviceID, cl_context context,
cl_command_queue queue, const char *kernelCode,
ExplicitType vecType, int vecSize,
bool export_after_import, bool aliased_local_mem,
bool aliased_global_mem)
{
int error;
clProgramWrapper program;
@@ -710,10 +708,12 @@ int test_copy_fence(cl_device_id deviceID, cl_context context,
return failuresPrinted ? -1 : 0;
}
int test_copy_fence_all_types(cl_device_id deviceID, cl_context context,
cl_command_queue queue, const char *kernelCode,
bool export_after_import, bool aliased_local_mem,
bool aliased_global_mem)
static int test_copy_fence_all_types(cl_device_id deviceID, cl_context context,
cl_command_queue queue,
const char *kernelCode,
bool export_after_import,
bool aliased_local_mem,
bool aliased_global_mem)
{
ExplicitType vecType[] = {
kChar, kUChar, kShort, kUShort, kInt, kUInt, kLong,
@@ -757,56 +757,46 @@ int test_copy_fence_all_types(cl_device_id deviceID, cl_context context,
return 0;
}
int test_async_work_group_copy_fence_import_after_export_aliased_local(
cl_device_id deviceID, cl_context context, cl_command_queue queue,
int num_elements)
REGISTER_TEST(async_work_group_copy_fence_import_after_export_aliased_local)
{
return test_copy_fence_all_types(deviceID, context, queue,
return test_copy_fence_all_types(device, context, queue,
import_after_export_aliased_local_kernel,
false, true, false);
}
int test_async_work_group_copy_fence_import_after_export_aliased_global(
cl_device_id deviceID, cl_context context, cl_command_queue queue,
int num_elements)
REGISTER_TEST(async_work_group_copy_fence_import_after_export_aliased_global)
{
return test_copy_fence_all_types(deviceID, context, queue,
return test_copy_fence_all_types(device, context, queue,
import_after_export_aliased_global_kernel,
false, false, true);
}
int test_async_work_group_copy_fence_import_after_export_aliased_global_and_local(
cl_device_id deviceID, cl_context context, cl_command_queue queue,
int num_elements)
REGISTER_TEST(
async_work_group_copy_fence_import_after_export_aliased_global_and_local)
{
return test_copy_fence_all_types(
deviceID, context, queue,
device, context, queue,
import_after_export_aliased_global_and_local_kernel, false, true, true);
}
int test_async_work_group_copy_fence_export_after_import_aliased_local(
cl_device_id deviceID, cl_context context, cl_command_queue queue,
int num_elements)
REGISTER_TEST(async_work_group_copy_fence_export_after_import_aliased_local)
{
return test_copy_fence_all_types(deviceID, context, queue,
return test_copy_fence_all_types(device, context, queue,
export_after_import_aliased_local_kernel,
true, true, false);
}
int test_async_work_group_copy_fence_export_after_import_aliased_global(
cl_device_id deviceID, cl_context context, cl_command_queue queue,
int num_elements)
REGISTER_TEST(async_work_group_copy_fence_export_after_import_aliased_global)
{
return test_copy_fence_all_types(deviceID, context, queue,
return test_copy_fence_all_types(device, context, queue,
export_after_import_aliased_global_kernel,
true, false, true);
}
int test_async_work_group_copy_fence_export_after_import_aliased_global_and_local(
cl_device_id deviceID, cl_context context, cl_command_queue queue,
int num_elements)
REGISTER_TEST(
async_work_group_copy_fence_export_after_import_aliased_global_and_local)
{
return test_copy_fence_all_types(
deviceID, context, queue,
device, context, queue,
export_after_import_aliased_global_and_local_kernel, true, true, true);
}

25
test_conformance/basic/test_async_strided_copy.cpp
View File

@@ -13,8 +13,6 @@
// See the License for the specific language governing permissions and
// limitations under the License.
//
#include "harness/compat.h"
#include <stdio.h>
#include <stdlib.h>
#include <string.h>
@@ -22,8 +20,7 @@
#include <sys/stat.h>
#include <vector>
#include "procs.h"
#include "harness/conversions.h"
#include "testBase.h"
// clang-format off
@@ -66,7 +63,9 @@ static const char *async_strided_local_to_global_kernel =
// clang-format on
int test_strided_copy(cl_device_id deviceID, cl_context context, cl_command_queue queue, const char *kernelCode, ExplicitType vecType, int vecSize, int stride)
static int test_strided_copy(cl_device_id deviceID, cl_context context,
cl_command_queue queue, const char *kernelCode,
ExplicitType vecType, int vecSize, int stride)
{
int error;
clProgramWrapper program;
@@ -233,7 +232,10 @@ int test_strided_copy(cl_device_id deviceID, cl_context context, cl_command_queu
return 0;
}
int test_strided_copy_all_types(cl_device_id deviceID, cl_context context, cl_command_queue queue, const char *kernelCode)
static int test_strided_copy_all_types(cl_device_id deviceID,
cl_context context,
cl_command_queue queue,
const char *kernelCode)
{
const std::vector<ExplicitType> vecType = { kChar, kUChar, kShort, kUShort,
kInt, kUInt, kLong, kULong,
@@ -272,13 +274,14 @@ int test_strided_copy_all_types(cl_device_id deviceID, cl_context context, cl_co
return 0;
}
int test_async_strided_copy_global_to_local(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
REGISTER_TEST(async_strided_copy_global_to_local)
{
return test_strided_copy_all_types( deviceID, context, queue, async_strided_global_to_local_kernel );
return test_strided_copy_all_types(device, context, queue,
async_strided_global_to_local_kernel);
}
int test_async_strided_copy_local_to_global(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
REGISTER_TEST(async_strided_copy_local_to_global)
{
return test_strided_copy_all_types( deviceID, context, queue, async_strided_local_to_global_kernel );
return test_strided_copy_all_types(device, context, queue,
async_strided_local_to_global_kernel);
}

14
test_conformance/basic/test_barrier.cpp
View File

@@ -25,7 +25,7 @@
#include <numeric>
#include <vector>
#include "procs.h"
#include "testBase.h"
namespace {
const char *barrier_kernel_code = R"(
@@ -66,9 +66,9 @@ void generate_random_inputs(std::vector<cl_int> &v)
std::generate(v.begin(), v.end(), random_generator);
}
int test_barrier_common(cl_device_id device, cl_context context,
cl_command_queue queue, int num_elements,
std::string barrier_str)
static int test_barrier_common(cl_device_id device, cl_context context,
cl_command_queue queue, int num_elements,
std::string barrier_str)
{
clMemWrapper streams[3];
clProgramWrapper program;
@@ -142,14 +142,12 @@ int test_barrier_common(cl_device_id device, cl_context context,
}
}
int test_barrier(cl_device_id device, cl_context context,
cl_command_queue queue, int num_elements)
REGISTER_TEST(barrier)
{
return test_barrier_common(device, context, queue, num_elements, "barrier");
}
int test_wg_barrier(cl_device_id device, cl_context context,
cl_command_queue queue, int num_elements)
REGISTER_TEST_VERSION(wg_barrier, Version(2, 0))
{
return test_barrier_common(device, context, queue, num_elements,
"work_group_barrier");

9
test_conformance/basic/test_basic_parameter_types.cpp
View File

@@ -21,7 +21,7 @@
#include <sys/types.h>
#include <sys/stat.h>
#include "procs.h"
#include "testBase.h"
const char *kernel_code = R"(
__kernel void test_kernel(
@@ -46,8 +46,8 @@ __global float%s *result)
result[1] = %s(ul);
})";
int test_parameter_types_long(cl_device_id device, cl_context context,
cl_command_queue queue, int num_elements)
static int test_parameter_types_long(cl_device_id device, cl_context context,
cl_command_queue queue, int num_elements)
{
clMemWrapper results;
int error;
@@ -196,8 +196,7 @@ int test_parameter_types_long(cl_device_id device, cl_context context,
return total_errors;
}
int test_parameter_types(cl_device_id device, cl_context context,
cl_command_queue queue, int num_elements)
REGISTER_TEST(parameter_types)
{
clMemWrapper results;
int error;

64
test_conformance/basic/test_bufferreadwriterect.cpp
View File

@@ -21,7 +21,7 @@
#include <sys/types.h>
#include <sys/stat.h>
#include "procs.h"
#include "testBase.h"
#define CL_EXIT_ERROR(cmd,format,...) \
{ \
@@ -36,7 +36,7 @@ log_error("\n"); \
typedef unsigned char BufferType;
// Globals for test
cl_command_queue queue;
cl_command_queue gQueue;
// Width and height of each pair of images.
enum { TotalImages = 8 };
@@ -150,13 +150,12 @@ int copy_region(size_t src, size_t soffset[3], size_t sregion[3], size_t dst, si
log_info( "Copy overlap reported, skipping copy buffer rect\n" );
return CL_SUCCESS;
} else {
if ((err = clEnqueueCopyBufferRect(queue,
buffer[src],buffer[dst],
soffset, doffset,
sregion,/*dregion,*/
width[src], src_slice_pitch,
width[dst], dst_slice_pitch,
0, NULL, NULL)) != CL_SUCCESS)
if ((err = clEnqueueCopyBufferRect(
gQueue, buffer[src], buffer[dst], soffset, doffset,
sregion, /*dregion,*/
width[src], src_slice_pitch, width[dst], dst_slice_pitch, 0,
NULL, NULL))
!= CL_SUCCESS)
{
CL_EXIT_ERROR(err, "clEnqueueCopyBufferRect failed between %u and %u",(unsigned)src,(unsigned)dst);
}
@@ -253,15 +252,12 @@ int read_verify_region(size_t src, size_t soffset[3], size_t sregion[3], size_t
size_t dst_slice_pitch = (width[dst]*height[dst] != 1) ? width[dst]*height[dst] : 0;
// Copy the source region of the cl buffer, to the destination region of the temporary buffer.
CL_EXIT_ERROR(clEnqueueReadBufferRect(queue,
buffer[src],
CL_TRUE,
soffset,doffset,
sregion,
width[src], src_slice_pitch,
width[dst], dst_slice_pitch,
tmp_buffer,
0, NULL, NULL), "clEnqueueCopyBufferRect failed between %u and %u",(unsigned)src,(unsigned)dst);
CL_EXIT_ERROR(clEnqueueReadBufferRect(
gQueue, buffer[src], CL_TRUE, soffset, doffset, sregion,
width[src], src_slice_pitch, width[dst], dst_slice_pitch,
tmp_buffer, 0, NULL, NULL),
"clEnqueueCopyBufferRect failed between %u and %u",
(unsigned)src, (unsigned)dst);
return verify_region(tmp_buffer,src,soffset,sregion,dst,doffset);
}
@@ -276,7 +272,9 @@ int map_verify_region(size_t src) {
// Copy the source region of the cl buffer, to the destination region of the temporary buffer.
cl_int err;
BufferType* mapped = (BufferType*)clEnqueueMapBuffer(queue,buffer[src],CL_TRUE,CL_MAP_READ,0,size_bytes,0,NULL,NULL,&err);
BufferType* mapped = (BufferType*)clEnqueueMapBuffer(
gQueue, buffer[src], CL_TRUE, CL_MAP_READ, 0, size_bytes, 0, NULL, NULL,
&err);
CL_EXIT_ERROR(err, "clEnqueueMapBuffer failed for buffer %u",(unsigned)src);
size_t soffset[] = { 0, 0, 0 };
@@ -284,8 +282,9 @@ int map_verify_region(size_t src) {
int ret = verify_region(mapped,src,soffset,sregion,src,soffset);
CL_EXIT_ERROR(clEnqueueUnmapMemObject(queue,buffer[src],mapped,0,NULL,NULL),
"clEnqueueUnmapMemObject failed for buffer %u",(unsigned)src);
CL_EXIT_ERROR(
clEnqueueUnmapMemObject(gQueue, buffer[src], mapped, 0, NULL, NULL),
"clEnqueueUnmapMemObject failed for buffer %u", (unsigned)src);
return ret;
}
@@ -301,15 +300,12 @@ int write_region(size_t src, size_t soffset[3], size_t sregion[3], size_t dst, s
size_t dst_slice_pitch = (width[dst]*height[dst] != 1) ? width[dst]*height[dst] : 0;
// Copy the source region of the cl buffer, to the destination region of the temporary buffer.
CL_EXIT_ERROR(clEnqueueWriteBufferRect(queue,
buffer[dst],
CL_TRUE,
doffset,soffset,
/*sregion,*/dregion,
width[dst], dst_slice_pitch,
width[src], src_slice_pitch,
tmp_buffer,
0, NULL, NULL), "clEnqueueWriteBufferRect failed between %u and %u",(unsigned)src,(unsigned)dst);
CL_EXIT_ERROR(clEnqueueWriteBufferRect(
gQueue, buffer[dst], CL_TRUE, doffset, soffset,
/*sregion,*/ dregion, width[dst], dst_slice_pitch,
width[src], src_slice_pitch, tmp_buffer, 0, NULL, NULL),
"clEnqueueWriteBufferRect failed between %u and %u",
(unsigned)src, (unsigned)dst);
// Copy from the temporary buffer to the host buffer.
size_t spitch = width[src];
@@ -345,10 +341,9 @@ void CL_CALLBACK mem_obj_destructor_callback( cl_mem, void *data )
}
// This is the main test function for the conformance test.
int
test_bufferreadwriterect(cl_device_id device, cl_context context, cl_command_queue queue_, int num_elements)
REGISTER_TEST(bufferreadwriterect)
{
queue = queue_;
gQueue = queue;
cl_int err;
// Initialize the random number generator.
@@ -564,6 +559,3 @@ test_bufferreadwriterect(cl_device_id device, cl_context context, cl_command_que
return err;
}

5
test_conformance/basic/test_constant.cpp
View File

@@ -24,7 +24,7 @@
#include <algorithm>
#include <vector>
#include "procs.h"
#include "testBase.h"
namespace {
const char* constant_kernel_code = R"(
@@ -99,8 +99,7 @@ template <typename T> void generate_random_inputs(std::vector<T>& v)
}
}
int test_constant(cl_device_id device, cl_context context,
cl_command_queue queue, int num_elements)
REGISTER_TEST(constant)
{
clMemWrapper streams[3];
clProgramWrapper program;

10
test_conformance/basic/test_constant_source.cpp
View File

@@ -21,8 +21,7 @@
#include <sys/types.h>
#include <sys/stat.h>
#include "procs.h"
#include "testBase.h"
const char *constant_source_kernel_code[] = {
"__constant int outVal = 42;\n"
@@ -44,7 +43,7 @@ const char *constant_source_kernel_code[] = {
" }\n"
"}\n" };
int test_constant_source(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
REGISTER_TEST(constant_source)
{
clProgramWrapper program;
clKernelWrapper kernel;
@@ -94,8 +93,3 @@ int test_constant_source(cl_device_id device, cl_context context, cl_command_que
return 0;
}

11
test_conformance/basic/test_createkernelsinprogram.cpp
View File

@@ -21,8 +21,7 @@
#include <sys/types.h>
#include <sys/stat.h>
#include "procs.h"
#include "testBase.h"
const char *sample_single_kernel = {
"__kernel void sample_test(__global float *src, __global int *dst)\n"
@@ -50,8 +49,7 @@ const char *sample_double_kernel = {
"}\n"};
int
test_createkernelsinprogram(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
REGISTER_TEST(createkernelsinprogram)
{
cl_program program;
cl_kernel kernel[2];
@@ -98,8 +96,3 @@ test_createkernelsinprogram(cl_device_id device, cl_context context, cl_command_
return err;
}

11
test_conformance/basic/test_enqueue_map.cpp
View File

@@ -20,9 +20,6 @@
#include <string.h>
#include <sys/types.h>
#include <sys/stat.h>
#include "procs.h"
#include "harness/conversions.h"
#include "harness/typeWrappers.h"
@@ -44,8 +41,7 @@ const char *flag_set_names[] = {
};
// clang-format on
int test_enqueue_map_buffer(cl_device_id deviceID, cl_context context,
cl_command_queue queue, int num_elements)
REGISTER_TEST(enqueue_map_buffer)
{
int error;
const size_t bufferSize = 256 * 256;
@@ -141,15 +137,14 @@ int test_enqueue_map_buffer(cl_device_id deviceID, cl_context context,
return 0;
}
int test_enqueue_map_image(cl_device_id deviceID, cl_context context,
cl_command_queue queue, int num_elements)
REGISTER_TEST(enqueue_map_image)
{
int error;
cl_image_format format = { CL_RGBA, CL_UNSIGNED_INT32 };
const size_t imageSize = 256;
const size_t imageDataSize = imageSize * imageSize * 4 * sizeof(cl_uint);
PASSIVE_REQUIRE_IMAGE_SUPPORT(deviceID)
PASSIVE_REQUIRE_IMAGE_SUPPORT(device)
BufferOwningPtr<cl_uint> hostPtrData{ malloc(imageDataSize) };
BufferOwningPtr<cl_uint> referenceData{ malloc(imageDataSize) };

5
test_conformance/basic/test_enqueued_local_size.cpp
View File

@@ -24,7 +24,7 @@
#include <algorithm>
#include "procs.h"
#include "testBase.h"
static const char *enqueued_local_size_2d_code = R"(
__kernel void test_enqueued_local_size_2d(global int *dst)
@@ -65,8 +65,7 @@ static int verify_enqueued_local_size(int *result, size_t *expected, int n)
}
int test_enqueued_local_size(cl_device_id device, cl_context context,
cl_command_queue queue, int num_elements)
REGISTER_TEST_VERSION(enqueued_local_size, Version(2, 0))
{
clMemWrapper stream;
clProgramWrapper program[2];

16
test_conformance/basic/test_explicit_s2v.cpp
View File

@@ -25,8 +25,6 @@ using std::isnan;
#include <vector>
#include <CL/cl_half.h>
#include "procs.h"
#include "harness/conversions.h"
#include "harness/typeWrappers.h"
@@ -114,10 +112,11 @@ bool IsHalfNaN(cl_half v)
return (h_exp == 0x1F && h_mant != 0);
}
int test_explicit_s2v_function(cl_context context, cl_command_queue queue,
cl_kernel kernel, ExplicitType srcType,
unsigned int count, ExplicitType destType,
unsigned int vecSize, void *inputData)
static int test_explicit_s2v_function(cl_context context,
cl_command_queue queue, cl_kernel kernel,
ExplicitType srcType, unsigned int count,
ExplicitType destType,
unsigned int vecSize, void *inputData)
{
int error;
clMemWrapper streams[2];
@@ -387,12 +386,11 @@ protected:
} // anonymous namespace
int test_explicit_s2v(cl_device_id deviceID, cl_context context,
cl_command_queue queue, int num_elements)
REGISTER_TEST(explicit_s2v)
{
try
{
TypesIterator(deviceID, context, queue);
TypesIterator(device, context, queue);
} catch (const std::runtime_error &e)
{
log_error("%s", e.what());

5
test_conformance/basic/test_fpmath.cpp
View File

@@ -31,7 +31,7 @@
#include <string>
#include <vector>
#include "procs.h"
#include "testBase.h"
extern cl_half_rounding_mode halfRoundingMode;
@@ -380,8 +380,7 @@ protected:
} // anonymous namespace
int test_fpmath(cl_device_id device, cl_context context, cl_command_queue queue,
int num_elements)
REGISTER_TEST(fpmath)
{
try
{

9
test_conformance/basic/test_get_linear_ids.cpp
View File

@@ -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
@@ -13,9 +13,10 @@
// See the License for the specific language governing permissions and
// limitations under the License.
//
#include "procs.h"
#include <ctype.h>
#include "testBase.h"
static const char *linear_ids_source[1] = {
"__kernel void test_linear_ids(__global int2 *out)\n"
"{\n"
@@ -46,8 +47,7 @@ static const char *linear_ids_source[1] = {
#define MAX_3D 16
#define MAX_OFFSET 100000
int
test_get_linear_ids(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
REGISTER_TEST_VERSION(get_linear_ids, Version(2, 0))
{
clProgramWrapper program;
clKernelWrapper kernel;
@@ -193,4 +193,3 @@ test_get_linear_ids(cl_device_id device, cl_context context, cl_command_queue qu
free_mtdata(seed);
return 0;
}

5
test_conformance/basic/test_global_linear_id.cpp
View File

@@ -25,7 +25,7 @@
#include <algorithm>
#include <vector>
#include "procs.h"
#include "testBase.h"
namespace {
const char *global_linear_id_2d_code = R"(
@@ -64,8 +64,7 @@ int verify_global_linear_id(std::vector<cl_int> &result, int n)
}
}
int test_global_linear_id(cl_device_id device, cl_context context,
cl_command_queue queue, int num_elements)
REGISTER_TEST_VERSION(global_linear_id, Version(2, 0))
{
clProgramWrapper program[2];
clKernelWrapper kernel[2];

9
test_conformance/basic/test_global_work_offsets.cpp
View File

@@ -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
@@ -13,9 +13,9 @@
// See the License for the specific language governing permissions and
// limitations under the License.
//
#include "procs.h"
#include <ctype.h>
#include "testBase.h"
const char *work_offset_test[] = {
"__kernel void test( __global int * outputID_A, \n"
@@ -122,7 +122,7 @@ int check_results( size_t threads[], size_t offsets[], cl_int outputA[], cl_int
return ( missed | multiple | errored | corrected );
}
int test_global_work_offsets(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
REGISTER_TEST(global_work_offsets)
{
clProgramWrapper program;
clKernelWrapper kernel;
@@ -222,7 +222,7 @@ const char *get_offset_test[] = {
"}\n"
};
int test_get_global_offset(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
REGISTER_TEST(get_global_offset)
{
clProgramWrapper program;
clKernelWrapper kernel;
@@ -300,4 +300,3 @@ int test_get_global_offset(cl_device_id deviceID, cl_context context, cl_command
// All done!
return 0;
}

12
test_conformance/basic/test_hiloeo.cpp
View File

@@ -21,7 +21,7 @@
#include <sys/stat.h>
#include <vector>
#include "procs.h"
#include "testBase.h"
int hi_offset( int index, int vectorSize) { return index + vectorSize / 2; }
int lo_offset( int index, int vectorSize) { return index; }
@@ -48,7 +48,7 @@ static const char *vector_size_names[] = { "", "2", "3", "4", "8", "16"};
static const size_t kSizes[] = { 1, 1, 2, 2, 4, 4, 8, 8, 2, 4, 8 };
static int CheckResults( void *in, void *out, size_t elementCount, int type, int vectorSize, int operatorToUse );
int test_hiloeo(cl_device_id device, cl_context context, cl_command_queue queue, int n_elems)
REGISTER_TEST(hiloeo)
{
int err;
int hasDouble = is_extension_available( device, "cl_khr_fp64" );
@@ -60,12 +60,12 @@ int test_hiloeo(cl_device_id device, cl_context context, cl_command_queue queue,
int expressionMode;
int numExpressionModes = 2;
size_t length = sizeof(cl_int) * 4 * n_elems;
size_t length = sizeof(cl_int) * 4 * num_elements;
std::vector<cl_int> input_ptr(4 * n_elems);
std::vector<cl_int> output_ptr(4 * n_elems);
std::vector<cl_int> input_ptr(4 * num_elements);
std::vector<cl_int> output_ptr(4 * num_elements);
for (cl_uint i = 0; i < 4 * (cl_uint)n_elems; i++)
for (cl_uint i = 0; i < 4 * (cl_uint)num_elements; i++)
input_ptr[i] = genrand_int32(d);
for( type = 0; type < sizeof( test_str_names ) / sizeof( test_str_names[0] ); type++ )

11
test_conformance/basic/test_hostptr.cpp
View File

@@ -21,8 +21,7 @@
#include <sys/types.h>
#include <sys/stat.h>
#include "procs.h"
#include "testBase.h"
const char *hostptr_kernel_code =
"__kernel void test_hostptr(__global float *srcA, __global float *srcB, __global float *dst)\n"
@@ -92,8 +91,7 @@ verify_rgba8_image(unsigned char *image, unsigned char *outptr, int w, int h)
return 0;
}
int
test_hostptr(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
REGISTER_TEST(hostptr)
{
cl_float *input_ptr[2], *output_ptr;
cl_program program;
@@ -282,8 +280,3 @@ test_hostptr(cl_device_id device, cl_context context, cl_command_queue queue, in
return err;
}

5
test_conformance/basic/test_if.cpp
View File

@@ -24,7 +24,7 @@
#include <algorithm>
#include <vector>
#include "procs.h"
#include "testBase.h"
namespace {
const char *conditional_kernel_code = R"(
@@ -88,8 +88,7 @@ void generate_random_inputs(std::vector<cl_int> &v)
std::generate(v.begin(), v.end(), random_generator);
}
}
int test_if(cl_device_id device, cl_context context, cl_command_queue queue,
int num_elements)
REGISTER_TEST(if)
{
clMemWrapper streams[2];
clProgramWrapper program;

14
test_conformance/basic/test_image_multipass.cpp
View File

@@ -21,8 +21,7 @@
#include <sys/types.h>
#include <sys/stat.h>
#include "procs.h"
#include "testBase.h"
static const char *image_to_image_kernel_integer_coord_code =
"\n"
@@ -143,8 +142,7 @@ verify_byte_image(unsigned char *image, unsigned char *outptr, int w, int h, int
return 0;
}
int
test_image_multipass_integer_coord(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
REGISTER_TEST(image_multipass_integer_coord)
{
int img_width = 512;
int img_height = 512;
@@ -397,8 +395,7 @@ test_image_multipass_integer_coord(cl_device_id device, cl_context context, cl_c
return err;
}
int
test_image_multipass_float_coord(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
REGISTER_TEST(image_multipass_float_coord)
{
int img_width = 512;
int img_height = 512;
@@ -637,8 +634,3 @@ test_image_multipass_float_coord(cl_device_id device, cl_context context, cl_com
return err;
}

4
test_conformance/basic/test_image_param.cpp
View File

@@ -21,8 +21,6 @@
#include <sys/types.h>
#include <sys/stat.h>
#include "procs.h"
#include "harness/typeWrappers.h"
#include "harness/imageHelpers.h"
#include "harness/conversions.h"
@@ -161,7 +159,7 @@ int validate_results( size_t width, size_t height, cl_image_format &format, char
return 0;
}
int test_image_param(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
REGISTER_TEST(image_param)
{
size_t sizes[] = { 64, 100, 128, 250, 512 };
cl_image_format formats[] = { { CL_RGBA, CL_UNORM_INT8 }, { CL_RGBA, CL_UNORM_INT16 }, { CL_RGBA, CL_FLOAT }, { CL_BGRA, CL_UNORM_INT8 } };

5
test_conformance/basic/test_image_r8.cpp
View File

@@ -24,7 +24,7 @@
#include <algorithm>
#include <vector>
#include "procs.h"
#include "testBase.h"
namespace {
const char *r_uint8_kernel_code = R"(
@@ -52,8 +52,7 @@ void generate_random_inputs(std::vector<cl_uchar> &v)
}
}
int test_image_r8(cl_device_id device, cl_context context,
cl_command_queue queue, int num_elements)
REGISTER_TEST(image_r8)
{
clMemWrapper streams[2];
clProgramWrapper program;

30
test_conformance/basic/test_imagearraycopy.cpp
View File

@@ -23,19 +23,14 @@
#include <sys/stat.h>
#include <vector>
#include <memory>
#include "procs.h"
using test_function_t = int (*)(cl_device_id, cl_context, cl_command_queue,
cl_mem_flags, cl_mem_flags, cl_mem_object_type,
const cl_image_format *);
int test_imagearraycopy_single_format(cl_device_id device, cl_context context,
cl_command_queue queue,
cl_mem_flags image_flags,
cl_mem_flags buffer_flags,
cl_mem_object_type image_type,
const cl_image_format *format)
static int test_imagearraycopy_single_format(
cl_device_id device, cl_context context, cl_command_queue queue,
cl_mem_flags image_flags, cl_mem_flags buffer_flags,
cl_mem_object_type image_type, const cl_image_format *format)
{
std::unique_ptr<cl_uchar, decltype(&free)> bufptr{ nullptr, free },
imgptr{ nullptr, free };
@@ -153,11 +148,12 @@ int test_imagearraycopy_single_format(cl_device_id device, cl_context context,
return err;
}
int test_imagearraycommon(cl_device_id device, cl_context context,
cl_command_queue queue, cl_mem_flags image_flags,
cl_mem_flags buffer_flags,
cl_mem_object_type image_type,
test_function_t test_function)
static int test_imagearraycommon(cl_device_id device, cl_context context,
cl_command_queue queue,
cl_mem_flags image_flags,
cl_mem_flags buffer_flags,
cl_mem_object_type image_type,
test_function_t test_function)
{
cl_int err;
cl_uint num_formats;
@@ -188,8 +184,7 @@ int test_imagearraycommon(cl_device_id device, cl_context context,
return err;
}
int test_imagearraycopy(cl_device_id device, cl_context context,
cl_command_queue queue, int num_elements)
REGISTER_TEST(imagearraycopy)
{
PASSIVE_REQUIRE_IMAGE_SUPPORT(device)
@@ -199,8 +194,7 @@ int test_imagearraycopy(cl_device_id device, cl_context context,
}
int test_imagearraycopy3d(cl_device_id device, cl_context context,
cl_command_queue queue, int num_elements)
REGISTER_TEST(imagearraycopy3d)
{
PASSIVE_REQUIRE_3D_IMAGE_SUPPORT(device)

9
test_conformance/basic/test_imagecopy.cpp
View File

@@ -21,8 +21,7 @@
#include <sys/types.h>
#include <sys/stat.h>
#include "procs.h"
#include "testBase.h"
static unsigned char *
generate_rgba8_image(int w, int h, MTdata d)
@@ -105,8 +104,7 @@ verify_rgbafp_image(float *image, float *outptr, int w, int h)
}
int
test_imagecopy(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
REGISTER_TEST(imagecopy)
{
cl_image_format img_format;
unsigned char *rgba8_inptr, *rgba8_outptr;
@@ -236,6 +234,3 @@ test_imagecopy(cl_device_id device, cl_context context, cl_command_queue queue,
return err;
}

38
test_conformance/basic/test_imagecopy3d.cpp
View File

@@ -21,8 +21,7 @@
#include <sys/types.h>
#include <sys/stat.h>
#include "procs.h"
#include "testBase.h"
static unsigned char *
generate_uint8_image(unsigned num_elements, MTdata d)
@@ -105,8 +104,7 @@ verify_float_image(float *image, float *outptr, unsigned num_elements)
}
int
test_imagecopy3d(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements_ignored)
REGISTER_TEST(imagecopy3d)
{
cl_image_format img_format;
unsigned char *rgba8_inptr, *rgba8_outptr;
@@ -118,20 +116,21 @@ test_imagecopy3d(cl_device_id device, cl_context context, cl_command_queue queue
int img_depth = 64;
int i;
cl_int err;
unsigned num_elements = img_width * img_height * img_depth * 4;
unsigned num_elems = img_width * img_height * img_depth * 4;
MTdata d;
PASSIVE_REQUIRE_3D_IMAGE_SUPPORT( device )
d = init_genrand( gRandomSeed );
rgba8_inptr = (unsigned char *)generate_uint8_image(num_elements, d);
rgba16_inptr = (unsigned short *)generate_uint16_image(num_elements, d);
rgbafp_inptr = (float *)generate_float_image(num_elements, d);
rgba8_inptr = (unsigned char *)generate_uint8_image(num_elems, d);
rgba16_inptr = (unsigned short *)generate_uint16_image(num_elems, d);
rgbafp_inptr = (float *)generate_float_image(num_elems, d);
free_mtdata(d); d = NULL;
rgba8_outptr = (unsigned char*)malloc(sizeof(unsigned char) * num_elements);
rgba16_outptr = (unsigned short*)malloc(sizeof(unsigned short) * num_elements);
rgbafp_outptr = (float*)malloc(sizeof(float) * num_elements);
rgba8_outptr = (unsigned char *)malloc(sizeof(unsigned char) * num_elems);
rgba16_outptr =
(unsigned short *)malloc(sizeof(unsigned short) * num_elems);
rgbafp_outptr = (float *)malloc(sizeof(float) * num_elems);
img_format.image_channel_order = CL_RGBA;
img_format.image_channel_data_type = CL_UNORM_INT8;
@@ -202,16 +201,18 @@ test_imagecopy3d(cl_device_id device, cl_context context, cl_command_queue queue
switch (i)
{
case 0:
err = verify_uint8_image(rgba8_inptr, rgba8_outptr, num_elements);
if (err) log_error("Failed uint8\n");
err = verify_uint8_image(rgba8_inptr, rgba8_outptr, num_elems);
if (err) log_error("Failed uint8\n");
break;
case 1:
err = verify_uint16_image(rgba16_inptr, rgba16_outptr, num_elements);
if (err) log_error("Failed uint16\n");
err =
verify_uint16_image(rgba16_inptr, rgba16_outptr, num_elems);
if (err) log_error("Failed uint16\n");
break;
case 2:
err = verify_float_image(rgbafp_inptr, rgbafp_outptr, num_elements);
if (err) log_error("Failed float\n");
err =
verify_float_image(rgbafp_inptr, rgbafp_outptr, num_elems);
if (err) log_error("Failed float\n");
break;
}
@@ -233,6 +234,3 @@ test_imagecopy3d(cl_device_id device, cl_context context, cl_command_queue queue
return err;
}

14
test_conformance/basic/test_imagedim.cpp
View File

@@ -24,7 +24,7 @@
#include <algorithm>
#include <vector>
#include "procs.h"
#include "testBase.h"
namespace {
const char *image_dim_kernel_code = R"(
@@ -105,9 +105,9 @@ int get_max_image_dimensions(cl_device_id device, size_t &max_img_width,
return err;
}
int test_imagedim_common(cl_context context, cl_command_queue queue,
cl_kernel kernel, size_t *local_threads,
size_t img_width, size_t img_height)
static int test_imagedim_common(cl_context context, cl_command_queue queue,
cl_kernel kernel, size_t *local_threads,
size_t img_width, size_t img_height)
{
int err;
@@ -173,8 +173,7 @@ int test_imagedim_common(cl_context context, cl_command_queue queue,
}
}
int test_imagedim_pow2(cl_device_id device, cl_context context,
cl_command_queue queue, int n_elems)
REGISTER_TEST(imagedim_pow2)
{
clProgramWrapper program;
clKernelWrapper kernel;
@@ -211,8 +210,7 @@ int test_imagedim_pow2(cl_device_id device, cl_context context,
}
int test_imagedim_non_pow2(cl_device_id device, cl_context context,
cl_command_queue queue, int n_elems)
REGISTER_TEST(imagedim_non_pow2)
{
clProgramWrapper program;
clKernelWrapper kernel;

24
test_conformance/basic/test_imagenpot.cpp
View File

@@ -21,8 +21,7 @@
#include <sys/types.h>
#include <sys/stat.h>
#include "procs.h"
#include "testBase.h"
static const char *rgba8888_kernel_code =
"\n"
@@ -74,8 +73,7 @@ verify_rgba8888_image(unsigned char *src, unsigned char *dst, int w, int h)
int img_width_selection[] = { 97, 111, 322, 479 };
int img_height_selection[] = { 149, 222, 754, 385 };
int
test_imagenpot(cl_device_id device_id, cl_context context, cl_command_queue queue, int num_elements)
REGISTER_TEST(imagenpot)
{
cl_mem streams[2];
cl_image_format img_format;
@@ -91,10 +89,11 @@ test_imagenpot(cl_device_id device_id, cl_context context, cl_command_queue queu
size_t max_local_workgroup_size[3];
MTdata d;
PASSIVE_REQUIRE_IMAGE_SUPPORT( device_id )
PASSIVE_REQUIRE_IMAGE_SUPPORT(device)
cl_device_type device_type;
err = clGetDeviceInfo(device_id, CL_DEVICE_TYPE, sizeof(device_type), &device_type, NULL);
err = clGetDeviceInfo(device, CL_DEVICE_TYPE, sizeof(device_type),
&device_type, NULL);
if (err) {
log_error("Failed to get device type: %d\n",err);
return -1;
@@ -163,10 +162,14 @@ test_imagenpot(cl_device_id device_id, cl_context context, cl_command_queue queu
return -1;
}
err = clGetKernelWorkGroupInfo(kernel, device_id, CL_KERNEL_WORK_GROUP_SIZE, sizeof(local_workgroup_size), &local_workgroup_size, NULL);
err = clGetKernelWorkGroupInfo(
kernel, device, CL_KERNEL_WORK_GROUP_SIZE,
sizeof(local_workgroup_size), &local_workgroup_size, NULL);
test_error(err, "clGetKernelWorkGroupInfo for CL_KERNEL_WORK_GROUP_SIZE failed");
err = clGetDeviceInfo(device_id, CL_DEVICE_MAX_WORK_ITEM_SIZES, sizeof(max_local_workgroup_size), max_local_workgroup_size, NULL);
err = clGetDeviceInfo(device, CL_DEVICE_MAX_WORK_ITEM_SIZES,
sizeof(max_local_workgroup_size),
max_local_workgroup_size, NULL);
test_error(err, "clGetDeviceInfo failed for CL_DEVICE_MAX_WORK_ITEM_SIZES");
// Pick the minimum of the device and the kernel
@@ -214,8 +217,3 @@ test_imagenpot(cl_device_id device_id, cl_context context, cl_command_queue queu
return err;
}

9
test_conformance/basic/test_imagerandomcopy.cpp
View File

@@ -21,8 +21,7 @@
#include <sys/types.h>
#include <sys/stat.h>
#include "procs.h"
#include "testBase.h"
static unsigned char *
generate_rgba8_image(int w, int h, MTdata d)
@@ -117,8 +116,7 @@ verify_rgbafp_image(float *image, float *outptr, int x, int y, int w, int h, int
#define NUM_COPIES 10
static const char *test_str_names[] = { "CL_RGBA CL_UNORM_INT8", "CL_RGBA CL_UNORM_INT16", "CL_RGBA CL_FLOAT" };
int
test_imagerandomcopy(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
REGISTER_TEST(imagerandomcopy)
{
cl_image_format img_format;
unsigned char *rgba8_inptr, *rgba8_outptr;
@@ -271,6 +269,3 @@ test_imagerandomcopy(cl_device_id device, cl_context context, cl_command_queue q
return err;
}

9
test_conformance/basic/test_imagereadwrite.cpp
View File

@@ -21,8 +21,7 @@
#include <sys/types.h>
#include <sys/stat.h>
#include "procs.h"
#include "testBase.h"
static unsigned char *
generate_rgba8_image(int w, int h, MTdata d)
@@ -188,8 +187,7 @@ verify_rgbafp_image(float *image, float *outptr, int w, int h)
}
int
test_imagereadwrite(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
REGISTER_TEST(imagereadwrite)
{
cl_image_format img_format;
unsigned char *rgba8_inptr, *rgba8_outptr;
@@ -422,6 +420,3 @@ test_imagereadwrite(cl_device_id device, cl_context context, cl_command_queue qu
return err;
}

9
test_conformance/basic/test_imagereadwrite3d.cpp
View File

@@ -21,8 +21,7 @@
#include <sys/types.h>
#include <sys/stat.h>
#include "procs.h"
#include "testBase.h"
static unsigned char *
generate_rgba8_image(int w, int h, int d, MTdata mtData)
@@ -196,8 +195,7 @@ verify_rgbafp_image(float *image, float *outptr, int w, int h, int d)
}
int
test_imagereadwrite3d(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
REGISTER_TEST(imagereadwrite3d)
{
cl_image_format img_format;
unsigned char *rgba8_inptr, *rgba8_outptr;
@@ -419,6 +417,3 @@ test_imagereadwrite3d(cl_device_id device, cl_context context, cl_command_queue
return err;
}

8
test_conformance/basic/test_int2fp.cpp
View File

@@ -29,7 +29,7 @@
#include <map>
#include <vector>
#include "procs.h"
#include "testBase.h"
extern cl_half_rounding_mode halfRoundingMode;
@@ -294,8 +294,7 @@ protected:
}
int test_int2fp(cl_device_id device, cl_context context, cl_command_queue queue,
int num_elements)
REGISTER_TEST(int2fp)
{
try
{
@@ -309,8 +308,7 @@ int test_int2fp(cl_device_id device, cl_context context, cl_command_queue queue,
return TEST_PASS;
}
int test_fp2int(cl_device_id device, cl_context context, cl_command_queue queue,
int num_elements)
REGISTER_TEST(fp2int)
{
try
{

25
test_conformance/basic/test_intmath.cpp
View File

@@ -18,7 +18,7 @@
#include <string>
#include <vector>
#include "procs.h"
#include "testBase.h"
template <typename T> struct TestDef
{
@@ -28,8 +28,9 @@ template <typename T> struct TestDef
};
template <typename T, unsigned N>
int test_intmath(cl_device_id device, cl_context context,
cl_command_queue queue, int num_elements, std::string typestr)
static int test_intmath(cl_device_id device, cl_context context,
cl_command_queue queue, int num_elements,
std::string typestr)
{
TestDef<T> tests[] = {
// Test addition
@@ -197,43 +198,37 @@ int test_intmath(cl_device_id device, cl_context context,
return TEST_PASS;
}
int test_intmath_int(cl_device_id device, cl_context context,
cl_command_queue queue, int num_elements)
REGISTER_TEST(intmath_int)
{
return test_intmath<cl_uint, 1>(device, context, queue, num_elements,
"uint");
}
int test_intmath_int2(cl_device_id device, cl_context context,
cl_command_queue queue, int num_elements)
REGISTER_TEST(intmath_int2)
{
return test_intmath<cl_uint, 2>(device, context, queue, num_elements,
"uint2");
}
int test_intmath_int4(cl_device_id device, cl_context context,
cl_command_queue queue, int num_elements)
REGISTER_TEST(intmath_int4)
{
return test_intmath<cl_uint, 4>(device, context, queue, num_elements,
"uint4");
}
int test_intmath_long(cl_device_id device, cl_context context,
cl_command_queue queue, int num_elements)
REGISTER_TEST(intmath_long)
{
return test_intmath<cl_ulong, 1>(device, context, queue, num_elements,
"ulong");
}
int test_intmath_long2(cl_device_id device, cl_context context,
cl_command_queue queue, int num_elements)
REGISTER_TEST(intmath_long2)
{
return test_intmath<cl_ulong, 2>(device, context, queue, num_elements,
"ulong2");
}
int test_intmath_long4(cl_device_id device, cl_context context,
cl_command_queue queue, int num_elements)
REGISTER_TEST(intmath_long4)
{
return test_intmath<cl_ulong, 4>(device, context, queue, num_elements,
"ulong4");

9
test_conformance/basic/test_kernel_call_kernel_function.cpp
View File

@@ -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
@@ -13,7 +13,7 @@
// See the License for the specific language governing permissions and
// limitations under the License.
//
#include "procs.h"
#include "testBase.h"
const char *kernel_call_kernel_code[] = {
"void test_function_to_call(__global int *output, __global int *input, int where);\n"
@@ -57,8 +57,7 @@ const char *kernel_call_kernel_code[] = {
};
int test_kernel_call_kernel_function(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
REGISTER_TEST(kernel_call_kernel_function)
{
num_elements = 256;
@@ -249,5 +248,3 @@ int test_kernel_call_kernel_function(cl_device_id deviceID, cl_context context,
return errors;
}

27
test_conformance/basic/test_kernel_memory_alignment.cpp
View File

@@ -17,7 +17,6 @@
#include <unistd.h>
#endif
#include "procs.h"
#include "harness/conversions.h"
#include "harness/typeWrappers.h"
#include "harness/errorHelpers.h"
@@ -148,7 +147,9 @@ const char * get_explicit_address_name( AddressSpaces address )
}
int test_kernel_memory_alignment(cl_device_id device, cl_context context, cl_command_queue queue, int n_elems, AddressSpaces address )
static int test_kernel_memory_alignment(cl_device_id device, cl_context context,
cl_command_queue queue, int n_elems,
AddressSpaces address)
{
const char *constant_kernel;
const char *parameter_kernel;
@@ -495,17 +496,19 @@ int test_kernel_memory_alignment(cl_device_id device, cl_context context, cl_com
}
int test_kernel_memory_alignment_local(cl_device_id device, cl_context context, cl_command_queue queue, int n_elems )
REGISTER_TEST(kernel_memory_alignment_local)
{
return test_kernel_memory_alignment( device, context, queue, n_elems, kLocal );
return test_kernel_memory_alignment(device, context, queue, num_elements,
kLocal);
}
int test_kernel_memory_alignment_global(cl_device_id device, cl_context context, cl_command_queue queue, int n_elems )
REGISTER_TEST(kernel_memory_alignment_global)
{
return test_kernel_memory_alignment( device, context, queue, n_elems, kGlobal );
return test_kernel_memory_alignment(device, context, queue, num_elements,
kGlobal);
}
int test_kernel_memory_alignment_constant(cl_device_id device, cl_context context, cl_command_queue queue, int n_elems )
REGISTER_TEST(kernel_memory_alignment_constant)
{
// There is a class of approved OpenCL 1.0 conformant devices out there that in some circumstances
// are unable to meaningfully take (or more precisely use) the address of constant data by virtue
@@ -559,12 +562,12 @@ int test_kernel_memory_alignment_constant(cl_device_id device, cl_context contex
free(version_string);
// Everyone else is to be ground mercilessly under the wheels of progress
return test_kernel_memory_alignment( device, context, queue, n_elems, kConstant );
return test_kernel_memory_alignment(device, context, queue, num_elements,
kConstant);
}
int test_kernel_memory_alignment_private(cl_device_id device, cl_context context, cl_command_queue queue, int n_elems )
REGISTER_TEST(kernel_memory_alignment_private)
{
return test_kernel_memory_alignment( device, context, queue, n_elems, kPrivate );
return test_kernel_memory_alignment(device, context, queue, num_elements,
kPrivate);
}

10
test_conformance/basic/test_local.cpp
View File

@@ -21,8 +21,7 @@
#include <sys/types.h>
#include <sys/stat.h>
#include "procs.h"
#include "testBase.h"
const char *barrier_with_localmem_kernel_code[] = {
"__kernel void compute_sum_with_localmem(__global int *a, int n, __local int *tmp_sum, __global int *sum)\n"
@@ -115,7 +114,7 @@ verify_sum(int *inptr, int *outptr, int n)
return 0;
}
int test_local_arg_def(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
REGISTER_TEST(local_arg_def)
{
cl_mem streams[2];
cl_program program;
@@ -231,7 +230,7 @@ int test_local_arg_def(cl_device_id device, cl_context context, cl_command_queue
return err;
}
int test_local_kernel_def(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
REGISTER_TEST(local_kernel_def)
{
cl_mem streams[2];
cl_program program;
@@ -363,6 +362,3 @@ int test_local_kernel_def(cl_device_id device, cl_context context, cl_command_qu
return err;
}

7
test_conformance/basic/test_local_kernel_scope.cpp
View File

@@ -21,8 +21,7 @@
#include <sys/types.h>
#include <sys/stat.h>
#include "procs.h"
#include "testBase.h"
#define MAX_LOCAL_STORAGE_SIZE 256
#define MAX_LOCAL_STORAGE_SIZE_STRING "256"
@@ -56,7 +55,7 @@ const char *kernelSource[] = {
"}\n"
};
int test_local_kernel_scope(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
REGISTER_TEST(local_kernel_scope)
{
cl_int error;
clProgramWrapper program;
@@ -138,5 +137,3 @@ int test_local_kernel_scope(cl_device_id device, cl_context context, cl_command_
free(outputData);
return 0;
}

7
test_conformance/basic/test_local_linear_id.cpp
View File

@@ -20,9 +20,9 @@
#include <string.h>
#include <sys/types.h>
#include <sys/stat.h>
#include "harness/rounding_mode.h"
#include "procs.h"
#include "testBase.h"
#include "harness/rounding_mode.h"
static const char *local_linear_id_1d_code =
"__kernel void test_local_linear_id_1d(global int *dst)\n"
@@ -63,8 +63,7 @@ verify_local_linear_id(int *result, int n)
}
int
test_local_linear_id(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
REGISTER_TEST_VERSION(local_linear_id, Version(2, 0))
{
cl_mem streams;
cl_program program[2];

5
test_conformance/basic/test_loop.cpp
View File

@@ -23,7 +23,7 @@
#include <vector>
#include "procs.h"
#include "testBase.h"
namespace {
const char *loop_kernel_code = R"(
@@ -68,8 +68,7 @@ int verify_loop(std::vector<cl_int> inptr, std::vector<cl_int> loopindx,
return 0;
}
}
int test_loop(cl_device_id device, cl_context context, cl_command_queue queue,
int num_elements)
REGISTER_TEST(loop)
{
clMemWrapper streams[4];
clProgramWrapper program;

11
test_conformance/basic/test_multireadimagemultifmt.cpp
View File

@@ -21,8 +21,7 @@
#include <sys/types.h>
#include <sys/stat.h>
#include "procs.h"
#include "testBase.h"
static const char *multireadimage_kernel_code =
"__kernel void test_multireadimage(read_only image2d_t img0, read_only image2d_t img1, \n"
@@ -110,8 +109,7 @@ verify_multireadimage(void *image[], float *outptr, int w, int h)
}
int
test_mri_multiple(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
REGISTER_TEST(mri_multiple)
{
cl_mem streams[4];
cl_image_format img_format;
@@ -229,8 +227,3 @@ test_mri_multiple(cl_device_id device, cl_context context, cl_command_queue queu
return err;
}

10
test_conformance/basic/test_multireadimageonefmt.cpp
View File

@@ -21,8 +21,7 @@
#include <sys/types.h>
#include <sys/stat.h>
#include "procs.h"
#include "testBase.h"
static const char *multireadimage_kernel_code =
"__kernel void test_multireadimage(int n, int m, sampler_t sampler, \n"
@@ -93,7 +92,7 @@ verify_multireadimage(void *image[], int num_images, float *outptr, int w, int h
}
int test_mri_one(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
REGISTER_TEST(mri_one)
{
cl_mem streams[8];
cl_image_format img_format;
@@ -190,8 +189,3 @@ int test_mri_one(cl_device_id device, cl_context context, cl_command_queue queue
return err;
}

28
test_conformance/basic/test_numeric_constants.cpp
View File

@@ -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
@@ -13,7 +13,7 @@
// See the License for the specific language governing permissions and
// limitations under the License.
//
#include "procs.h"
#include "testBase.h"
#define TEST_VALUE_POSITIVE( string_name, name, value ) \
{ \
@@ -55,7 +55,7 @@ log_info("\t" string_name ": " #name " = %a (%17.21g)\n", value, value); \
} \
}
int test_host_numeric_constants(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
REGISTER_TEST(host_numeric_constants)
{
int errors = 0;
TEST_VALUE_EQUAL_LITERAL( "CL_CHAR_BIT", CL_CHAR_BIT, 8)
@@ -215,7 +215,7 @@ const char *kernel_double[] = {
};
int test_kernel_numeric_constants(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
REGISTER_TEST(kernel_numeric_constants)
{
int error, errors = 0;
// clProgramWrapper program;
@@ -307,7 +307,7 @@ int test_kernel_numeric_constants(cl_device_id deviceID, cl_context context, cl_
TEST_VALUE_EQUAL( "M_SQRT1_2_F", float_out[15], CL_M_SQRT1_2_F )
// We need to check these values against what we know is supported on the device
if( checkForImageSupport( deviceID ) == 0 )
if (checkForImageSupport(device) == 0)
{ // has images
// If images are supported, the constant should have been defined to the value 1
if( int_out[18] == 0xf00baa )
@@ -383,7 +383,8 @@ int test_kernel_numeric_constants(cl_device_id deviceID, cl_context context, cl_
/** DOUBLEs **/
if(!is_extension_available(deviceID, "cl_khr_fp64")) {
if (!is_extension_available(device, "cl_khr_fp64"))
{
log_info("Extension cl_khr_fp64 not supported; skipping double tests.\n");
}
else
@@ -529,7 +530,7 @@ const char *kernel_constant_double_limits[] = {
#define TEST_FLOAT_ASSERTION( a, msg, f ) if( !( a ) ) { log_error( "ERROR: Float constant failed requirement: %s (bitwise value is 0x%8.8x)\n", msg, *( (uint32_t *)&f ) ); return -1; }
#define TEST_DOUBLE_ASSERTION( a, msg, f ) if( !( a ) ) { log_error( "ERROR: Double constant failed requirement: %s (bitwise value is 0x%16.16llx)\n", msg, *( (uint64_t *)&f ) ); return -1; }
int test_kernel_limit_constants(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
REGISTER_TEST(kernel_limit_constants)
{
int error;
size_t threads[] = {1,1,1};
@@ -583,7 +584,8 @@ int test_kernel_limit_constants(cl_device_id deviceID, cl_context context, cl_co
// Stage 2: INFINITY and NAN
char profileStr[128] = "";
error = clGetDeviceInfo( deviceID, CL_DEVICE_PROFILE, sizeof( profileStr ), &profileStr, NULL );
error = clGetDeviceInfo(device, CL_DEVICE_PROFILE, sizeof(profileStr),
&profileStr, NULL);
test_error( error, "Unable to run INFINITY/NAN tests (unable to get CL_DEVICE_PROFILE" );
bool testInfNan = true;
@@ -591,7 +593,8 @@ int test_kernel_limit_constants(cl_device_id deviceID, cl_context context, cl_co
{
// We test if we're not an embedded profile, OR if the inf/nan flag in the config is set
cl_device_fp_config single = 0;
error = clGetDeviceInfo( deviceID, CL_DEVICE_SINGLE_FP_CONFIG, sizeof( single ), &single, NULL );
error = clGetDeviceInfo(device, CL_DEVICE_SINGLE_FP_CONFIG,
sizeof(single), &single, NULL);
test_error( error, "Unable to run INFINITY/NAN tests (unable to get FP_CONFIG bits)" );
if( ( single & CL_FP_INF_NAN ) == 0 )
@@ -666,12 +669,13 @@ int test_kernel_limit_constants(cl_device_id deviceID, cl_context context, cl_co
}
// Stage 3: limits on HUGE_VAL (double)
if( !is_extension_available( deviceID, "cl_khr_fp64" ) )
if (!is_extension_available(device, "cl_khr_fp64"))
log_info( "Note: Skipping double HUGE_VAL tests (doubles unsupported on device)\n" );
else
{
cl_device_fp_config config = 0;
error = clGetDeviceInfo( deviceID, CL_DEVICE_DOUBLE_FP_CONFIG, sizeof( config ), &config, NULL );
error = clGetDeviceInfo(device, CL_DEVICE_DOUBLE_FP_CONFIG,
sizeof(config), &config, NULL);
test_error( error, "Unable to run INFINITY/NAN tests (unable to get double FP_CONFIG bits)" );
if( ( config & CL_FP_INF_NAN ) == 0 )
@@ -716,5 +720,3 @@ int test_kernel_limit_constants(cl_device_id deviceID, cl_context context, cl_co
return 0;
}

7
test_conformance/basic/test_pointercast.cpp
View File

@@ -21,8 +21,7 @@
#include <sys/types.h>
#include <sys/stat.h>
#include "procs.h"
#include "testBase.h"
static const char *pointer_cast_kernel_code =
"__kernel void test_pointer_cast(__global unsigned char *src, __global unsigned int *dst)\n"
@@ -57,7 +56,7 @@ verify_pointer_cast(unsigned char *inptr, unsigned int *outptr, int n)
return 0;
}
int test_pointer_cast(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
REGISTER_TEST(pointer_cast)
{
cl_mem streams[2];
unsigned char *input_ptr;
@@ -137,5 +136,3 @@ int test_pointer_cast(cl_device_id device, cl_context context, cl_command_queue
return err;
}

22
test_conformance/basic/test_preprocessors.cpp
View File

@@ -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
@@ -13,9 +13,10 @@
// See the License for the specific language governing permissions and
// limitations under the License.
//
#include "procs.h"
#include <ctype.h>
#include "testBase.h"
// Test __FILE__, __LINE__, __OPENCL_VERSION__, __OPENCL_C_VERSION__, __ENDIAN_LITTLE__, __ROUNDING_MODE__, __IMAGE_SUPPORT__, __FAST_RELAXED_MATH__
// __kernel_exec
@@ -83,7 +84,7 @@ const char *preprocessor_test = {
"}\n"
};
int test_kernel_preprocessor_macros(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
REGISTER_TEST(kernel_preprocessor_macros)
{
clProgramWrapper program;
clKernelWrapper kernel;
@@ -156,7 +157,7 @@ int test_kernel_preprocessor_macros(cl_device_id deviceID, cl_context context, c
/////// Check the integer results
// We need to check these values against what we know is supported on the device
if( checkForImageSupport( deviceID ) == 0 )
if (checkForImageSupport(device) == 0)
{
// If images are supported, the constant should have been defined to the value 1
if( results[ 0 ] == 0xf00baa )
@@ -182,7 +183,9 @@ int test_kernel_preprocessor_macros(cl_device_id deviceID, cl_context context, c
// __ENDIAN_LITTLE__ is similar to __IMAGE_SUPPORT__: 1 if it's true, undefined if it isn't
cl_bool deviceIsLittleEndian;
error = clGetDeviceInfo( deviceID, CL_DEVICE_ENDIAN_LITTLE, sizeof( deviceIsLittleEndian ), &deviceIsLittleEndian, NULL );
error = clGetDeviceInfo(device, CL_DEVICE_ENDIAN_LITTLE,
sizeof(deviceIsLittleEndian), &deviceIsLittleEndian,
NULL);
test_error( error, "Unable to get endian property of device to validate against" );
if( deviceIsLittleEndian )
@@ -216,7 +219,7 @@ int test_kernel_preprocessor_macros(cl_device_id deviceID, cl_context context, c
// The OpenCL version reported by the macro reports the feature level supported by the compiler. Since
// this doesn't directly match any property we can query, we just check to see if it's a sane value
auto device_cl_version = get_device_cl_version(deviceID);
auto device_cl_version = get_device_cl_version(device);
int device_cl_version_int = device_cl_version.to_uint() * 10;
if ((results[2] < 100) || (results[2] > device_cl_version_int))
{
@@ -241,11 +244,11 @@ int test_kernel_preprocessor_macros(cl_device_id deviceID, cl_context context, c
int cl_c_minor_version = (results[3] / 10) % 10;
if ((results[3] < 100)
|| (!device_supports_cl_c_version(
deviceID,
device,
Version{ (cl_uint)cl_c_major_version,
(cl_uint)cl_c_minor_version })))
{
auto device_version = get_device_cl_c_version(deviceID);
auto device_version = get_device_cl_c_version(device);
log_error(
"ERROR: Kernel preprocessor __OPENCL_C_VERSION__ does not make "
"sense w.r.t. device's version string! "
@@ -337,7 +340,7 @@ int test_kernel_preprocessor_macros(cl_device_id deviceID, cl_context context, c
test_error( error, "Unable to create test program" );
// Try compiling
error = clBuildProgram( programB, 1, &deviceID, "-cl-fast-relaxed-math", NULL, NULL );
error = clBuildProgram( programB, 1, &device, "-cl-fast-relaxed-math", NULL, NULL );
test_error( error, "Unable to build program" );
// Create a kernel again to run against
@@ -373,4 +376,3 @@ int test_kernel_preprocessor_macros(cl_device_id deviceID, cl_context context, c
return 0;
}

14
test_conformance/basic/test_progvar.cpp
View File

@@ -50,8 +50,6 @@
#include "harness/errorHelpers.h"
#include "harness/featureHelpers.h"
#include "harness/mt19937.h"
#include "procs.h"
////////////////////
// Device capabilities
@@ -2006,8 +2004,7 @@ static cl_int should_skip(cl_device_id device, cl_bool& skip)
// Test support for variables at program scope. Miscellaneous
int test_progvar_prog_scope_misc(cl_device_id device, cl_context context,
cl_command_queue queue, int num_elements)
REGISTER_TEST_VERSION(progvar_prog_scope_misc, Version(2, 0))
{
cl_bool skip{ CL_FALSE };
auto error = should_skip(device, skip);
@@ -2038,8 +2035,7 @@ int test_progvar_prog_scope_misc(cl_device_id device, cl_context context,
// Test support for variables at program scope. Unitialized data
int test_progvar_prog_scope_uninit(cl_device_id device, cl_context context,
cl_command_queue queue, int num_elements)
REGISTER_TEST_VERSION(progvar_prog_scope_uninit, Version(2, 0))
{
cl_bool skip{ CL_FALSE };
auto error = should_skip(device, skip);
@@ -2068,8 +2064,7 @@ int test_progvar_prog_scope_uninit(cl_device_id device, cl_context context,
}
// Test support for variables at program scope. Initialized data.
int test_progvar_prog_scope_init(cl_device_id device, cl_context context,
cl_command_queue queue, int num_elements)
REGISTER_TEST_VERSION(progvar_prog_scope_init, Version(2, 0))
{
cl_bool skip{ CL_FALSE };
auto error = should_skip(device, skip);
@@ -2098,8 +2093,7 @@ int test_progvar_prog_scope_init(cl_device_id device, cl_context context,
// A simple test for support of static variables inside a kernel.
int test_progvar_func_scope(cl_device_id device, cl_context context,
cl_command_queue queue, int num_elements)
REGISTER_TEST_VERSION(progvar_func_scope, Version(2, 0))
{
cl_bool skip{ CL_FALSE };
auto error = should_skip(device, skip);

7
test_conformance/basic/test_queue_priority.cpp
View File

@@ -20,9 +20,9 @@
#include <string.h>
#include <sys/types.h>
#include <sys/stat.h>
#include "harness/rounding_mode.h"
#include "procs.h"
#include "testBase.h"
#include "harness/rounding_mode.h"
static const char *fpadd_kernel_code =
"__kernel void test_fpadd(__global float *srcA, __global float *srcB, __global float *dst)\n"
@@ -131,7 +131,7 @@ verify_fpmul(float *inptrA, float *inptrB, float *outptr, int n, int fileNum)
#if defined( __APPLE__ )
int test_queue_priority(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
REGISTER_TEST(queue_priority)
{
int err;
int command_queue_priority = 0;
@@ -342,6 +342,5 @@ int test_queue_priority(cl_device_id device, cl_context context, cl_command_queu
}
#endif

25
test_conformance/basic/test_readimage.cpp
View File

@@ -26,7 +26,7 @@
#include <string>
#include <vector>
#include "procs.h"
#include "testBase.h"
#define TEST_IMAGE_WIDTH_2D (512)
#define TEST_IMAGE_HEIGHT_2D (512)
@@ -204,8 +204,9 @@ cl_mem create_image_xd(cl_context context, cl_mem_flags flags,
}
template <cl_mem_object_type IMG_TYPE, typename T>
int test_readimage(cl_device_id device, cl_context context,
cl_command_queue queue, const cl_image_format *img_format)
static int test_readimage(cl_device_id device, cl_context context,
cl_command_queue queue,
const cl_image_format *img_format)
{
clMemWrapper streams[2];
clProgramWrapper program;
@@ -301,8 +302,7 @@ bool check_format(cl_device_id device, cl_context context,
}
}
int test_readimage(cl_device_id device, cl_context context,
cl_command_queue queue, int num_elements)
REGISTER_TEST(readimage)
{
const cl_image_format format[] = { { CL_RGBA, CL_UNORM_INT8 },
{ CL_BGRA, CL_UNORM_INT8 } };
@@ -319,24 +319,21 @@ int test_readimage(cl_device_id device, cl_context context,
return err;
}
int test_readimage_int16(cl_device_id device, cl_context context,
cl_command_queue queue, int num_elements)
REGISTER_TEST(readimage_int16)
{
const cl_image_format format = { CL_RGBA, CL_UNORM_INT16 };
return test_readimage<CL_MEM_OBJECT_IMAGE2D, cl_ushort>(device, context,
queue, &format);
}
int test_readimage_fp32(cl_device_id device, cl_context context,
cl_command_queue queue, int num_elements)
REGISTER_TEST(readimage_fp32)
{
const cl_image_format format = { CL_RGBA, CL_FLOAT };
return test_readimage<CL_MEM_OBJECT_IMAGE2D, cl_float>(device, context,
queue, &format);
}
int test_readimage3d(cl_device_id device, cl_context context,
cl_command_queue queue, int num_elements)
REGISTER_TEST(readimage3d)
{
const cl_image_format format[] = { { CL_RGBA, CL_UNORM_INT8 },
{ CL_BGRA, CL_UNORM_INT8 } };
@@ -355,8 +352,7 @@ int test_readimage3d(cl_device_id device, cl_context context,
return err;
}
int test_readimage3d_int16(cl_device_id device, cl_context context,
cl_command_queue queue, int num_elements)
REGISTER_TEST(readimage3d_int16)
{
const cl_image_format format = { CL_RGBA, CL_UNORM_INT16 };
@@ -365,8 +361,7 @@ int test_readimage3d_int16(cl_device_id device, cl_context context,
return test_readimage<CL_MEM_OBJECT_IMAGE3D, cl_ushort>(device, context,
queue, &format);
}
int test_readimage3d_fp32(cl_device_id device, cl_context context,
cl_command_queue queue, int num_elements)
REGISTER_TEST(readimage3d_fp32)
{
const cl_image_format format = { CL_RGBA, CL_FLOAT };

17
test_conformance/basic/test_rw_image_access_qualifier.cpp
View File

@@ -18,7 +18,7 @@
#include <string.h>
#include <sys/stat.h>
#include "procs.h"
#include "testBase.h"
#include "harness/clImageHelper.h"
static const char* rw_kernel_code =
@@ -41,10 +41,10 @@ static const char* rw_kernel_code =
"}\n";
int test_rw_image_access_qualifier(cl_device_id device_id, cl_context context, cl_command_queue commands, int num_elements)
REGISTER_TEST_VERSION(rw_image_access_qualifier, Version(2, 0))
{
// This test should be skipped if images are not supported.
if (checkForImageSupport(device_id))
if (checkForImageSupport(device))
{
return TEST_SKIPPED_ITSELF;
}
@@ -53,11 +53,11 @@ int test_rw_image_access_qualifier(cl_device_id device_id, cl_context context, c
// or 2.X device if the device supports images. In OpenCL-3.0
// read-write images are optional. This test is already being skipped
// for 1.X devices.
if (get_device_cl_version(device_id) >= Version(3, 0))
if (get_device_cl_version(device) >= Version(3, 0))
{
cl_uint are_rw_images_supported{};
test_error(
clGetDeviceInfo(device_id, CL_DEVICE_MAX_READ_WRITE_IMAGE_ARGS,
clGetDeviceInfo(device, CL_DEVICE_MAX_READ_WRITE_IMAGE_ARGS,
sizeof(are_rw_images_supported),
&are_rw_images_supported, nullptr),
"clGetDeviceInfo failed for CL_DEVICE_MAX_READ_IMAGE_ARGS\n");
@@ -155,15 +155,14 @@ int test_rw_image_access_qualifier(cl_device_id device_id, cl_context context, c
err = CL_SUCCESS;
unsigned int num_iter = 1;
for(i = 0; i < num_iter; i++) {
err |= clEnqueueNDRangeKernel(commands, kernel, dim_count,
NULL, global_dim, local_dim,
0, NULL, NULL);
err |= clEnqueueNDRangeKernel(queue, kernel, dim_count, NULL,
global_dim, local_dim, 0, NULL, NULL);
}
/* Read back the results from the device to verify the output */
const size_t origin[3] = {0, 0, 0};
const size_t region[3] = {size_x, size_y, 1};
err |= clEnqueueReadImage(commands, src_image, CL_TRUE, origin, region, 0, 0,
err |= clEnqueueReadImage(queue, src_image, CL_TRUE, origin, region, 0, 0,
output, 0, NULL, NULL);
if (err != CL_SUCCESS) {
log_error("Error: clEnqueueReadBuffer failed\n");

26
test_conformance/basic/test_simple_image_pitch.cpp
View File

@@ -21,9 +21,9 @@
#include <sys/types.h>
#include <sys/stat.h>
#include "procs.h"
#include "testBase.h"
int test_simple_read_image_pitch(cl_device_id device, cl_context cl_context_, cl_command_queue q, int num_elements)
REGISTER_TEST(simple_read_image_pitch)
{
cl_int err = CL_SUCCESS;
@@ -51,7 +51,9 @@ int test_simple_read_image_pitch(cl_device_id device, cl_context cl_context_, cl
desc.image_width = imageW;
desc.image_height = imageH;
cl_mem image = clCreateImage(cl_context_, CL_MEM_COPY_HOST_PTR|CL_MEM_READ_WRITE, &fmt, &desc, host_image, &err);
cl_mem image =
clCreateImage(context, CL_MEM_COPY_HOST_PTR | CL_MEM_READ_WRITE, &fmt,
&desc, host_image, &err);
test_error(err,"clCreateImage");
char* host_buffer = (char*)malloc(buffer_bytes);
@@ -61,7 +63,8 @@ int test_simple_read_image_pitch(cl_device_id device, cl_context cl_context_, cl
size_t origin[] = { 0, 0, 0 };
size_t region[] = { imageW, imageH, 1 };
err = clEnqueueReadImage(q, image, CL_TRUE, origin, region, bufferW, 0, host_buffer, 0, NULL, NULL);
err = clEnqueueReadImage(queue, image, CL_TRUE, origin, region, bufferW, 0,
host_buffer, 0, NULL, NULL);
test_error(err,"clEnqueueReadImage");
size_t errors = 0;
@@ -88,7 +91,7 @@ int test_simple_read_image_pitch(cl_device_id device, cl_context cl_context_, cl
return errors == 0 ? TEST_PASS : TEST_FAIL;
}
int test_simple_write_image_pitch(cl_device_id device, cl_context cl_context_, cl_command_queue q, int num_elements)
REGISTER_TEST(simple_write_image_pitch)
{
cl_int err = CL_SUCCESS;
@@ -116,7 +119,9 @@ int test_simple_write_image_pitch(cl_device_id device, cl_context cl_context_, c
desc.image_width = imageW;
desc.image_height = imageH;
cl_mem image = clCreateImage(cl_context_, CL_MEM_COPY_HOST_PTR|CL_MEM_READ_WRITE, &fmt, &desc, host_image, &err);
cl_mem image =
clCreateImage(context, CL_MEM_COPY_HOST_PTR | CL_MEM_READ_WRITE, &fmt,
&desc, host_image, &err);
test_error(err,"clCreateImage");
char* host_buffer = (char*)malloc(buffer_bytes);
@@ -126,11 +131,14 @@ int test_simple_write_image_pitch(cl_device_id device, cl_context cl_context_, c
size_t origin[] = { 0, 0, 0 };
size_t region[] = { imageW, imageH, 1 };
err = clEnqueueWriteImage(q, image, CL_TRUE, origin, region, bufferW, 0, host_buffer, 0, NULL, NULL);
err = clEnqueueWriteImage(queue, image, CL_TRUE, origin, region, bufferW, 0,
host_buffer, 0, NULL, NULL);
test_error(err,"clEnqueueWriteImage");
size_t mapped_pitch = 0;
char* mapped_image = (char*)clEnqueueMapImage(q, image, CL_TRUE, CL_MAP_READ, origin, region, &mapped_pitch, NULL, 0, NULL, NULL, &err);
char* mapped_image = (char*)clEnqueueMapImage(
queue, image, CL_TRUE, CL_MAP_READ, origin, region, &mapped_pitch, NULL,
0, NULL, NULL, &err);
test_error(err,"clEnqueueMapImage");
size_t errors = 0;
@@ -145,7 +153,7 @@ int test_simple_write_image_pitch(cl_device_id device, cl_context cl_context_, c
}
}
err = clEnqueueUnmapMemObject(q, image, (void *)mapped_image, 0, 0, 0);
err = clEnqueueUnmapMemObject(queue, image, (void*)mapped_image, 0, 0, 0);
test_error(err,"clEnqueueUnmapMemObject");
test_error(clReleaseMemObject(image),"clReleaseMemObject");

7
test_conformance/basic/test_sizeof.cpp
View File

@@ -20,9 +20,8 @@
#include <string.h>
#include <sys/types.h>
#include <sys/stat.h>
#include "procs.h"
#include "testBase.h"
cl_int get_type_size( cl_context context, cl_command_queue queue, const char *type, cl_ulong *size, cl_device_id device )
{
@@ -127,7 +126,7 @@ const char *other_types[] =
static int IsPowerOfTwo( cl_ulong x ){ return 0 == (x & (x-1)); }
int test_sizeof(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
REGISTER_TEST(sizeof)
{
size_t i, j;
cl_ulong test;
@@ -376,5 +375,3 @@ int test_sizeof(cl_device_id device, cl_context context, cl_command_queue queue,
return err;
}

7
test_conformance/basic/test_vec_type_hint.cpp
View File

@@ -21,7 +21,6 @@
#include <sys/stat.h>
#include <vector>
#include "procs.h"
#include "harness/conversions.h"
#include "harness/typeWrappers.h"
@@ -36,7 +35,7 @@ static const char *sample_kernel = {
"}\n"
};
int test_vec_type_hint(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
REGISTER_TEST(vec_type_hint)
{
int error;
int vec_type_index, vec_size_index;
@@ -51,14 +50,14 @@ int test_vec_type_hint(cl_device_id deviceID, cl_context context, cl_command_que
{
if (vecType[vec_type_index] == kHalf
&& !is_extension_available(deviceID, "cl_khr_fp16"))
&& !is_extension_available(device, "cl_khr_fp16"))
{
log_info(
"Extension cl_khr_fp16 not supported; skipping half tests.\n");
continue;
}
else if (vecType[vec_type_index] == kDouble
&& !is_extension_available(deviceID, "cl_khr_fp64"))
&& !is_extension_available(device, "cl_khr_fp64"))
{
log_info(
"Extension cl_khr_fp64 not supported; skipping double tests.\n");

8
test_conformance/basic/test_vector_creation.cpp
View File

@@ -13,7 +13,6 @@
// See the License for the specific language governing permissions and
// limitations under the License.
//
#include "procs.h"
#include "harness/conversions.h"
#include "harness/typeWrappers.h"
#include "harness/errorHelpers.h"
@@ -244,8 +243,7 @@ int create_kernel(ExplicitType type, int output_size, char *program,
}
int test_vector_creation(cl_device_id deviceID, cl_context context,
cl_command_queue queue, int num_elements)
REGISTER_TEST(vector_creation)
{
const std::vector<ExplicitType> vecType = { kChar, kUChar, kShort, kUShort,
kInt, kUInt, kLong, kULong,
@@ -274,7 +272,7 @@ int test_vector_creation(cl_device_id deviceID, cl_context context,
}
else if (vecType[type_index] == kDouble)
{
if (!is_extension_available(deviceID, "cl_khr_fp64"))
if (!is_extension_available(device, "cl_khr_fp64"))
{
log_info("Extension cl_khr_fp64 not supported; skipping double "
"tests.\n");
@@ -285,7 +283,7 @@ int test_vector_creation(cl_device_id deviceID, cl_context context,
}
else if (vecType[type_index] == kHalf)
{
if (!is_extension_available(deviceID, "cl_khr_fp16"))
if (!is_extension_available(device, "cl_khr_fp16"))
{
log_info("Extension cl_khr_fp16 not supported; skipping half "
"tests.\n");

6
test_conformance/basic/test_vector_swizzle.cpp
View File

@@ -19,8 +19,7 @@
#include <string>
#include <vector>
#include "procs.h"
#include "harness/testHarness.h"
#include "testBase.h"
static std::string pragma_extension;
@@ -692,8 +691,7 @@ static int test_type(const char* type_name, cl_device_id device,
| test_vectype<T, 16>(type_name, device, context, queue);
}
int test_vector_swizzle(cl_device_id device, cl_context context,
cl_command_queue queue, int num_elements)
REGISTER_TEST(vector_swizzle)
{
int hasDouble = is_extension_available(device, "cl_khr_fp64");
int hasHalf = is_extension_available(device, "cl_khr_fp16");

46
test_conformance/basic/test_vloadstore.cpp
View File

@@ -23,8 +23,6 @@
#include <vector>
#include <CL/cl_half.h>
#include "procs.h"
#include "harness/conversions.h"
#include "harness/errorHelpers.h"
#include "harness/stringHelpers.h"
@@ -46,10 +44,10 @@ char load_str[128] = { 0 };
extern cl_half_rounding_mode halfRoundingMode;
// clang-format off
static const char *store_pattern= "results[ tid ] = tmp;\n";
static const char *store_patternV3 = "results[3*tid] = tmp.s0; results[3*tid+1] = tmp.s1; results[3*tid+2] = tmp.s2;\n";
static const char *load_pattern = "sSharedStorage[ i ] = src[ i ];\n";
static const char *load_patternV3 = "sSharedStorage[3*i] = src[ 3*i]; sSharedStorage[3*i+1] = src[3*i+1]; sSharedStorage[3*i+2] = src[3*i+2];\n";
static const char *const store_pattern= "results[ tid ] = tmp;\n";
static const char *const store_patternV3 = "results[3*tid] = tmp.s0; results[3*tid+1] = tmp.s1; results[3*tid+2] = tmp.s2;\n";
static const char *const load_pattern = "sSharedStorage[ i ] = src[ i ];\n";
static const char *const load_patternV3 = "sSharedStorage[3*i] = src[ 3*i]; sSharedStorage[3*i+1] = src[3*i+1]; sSharedStorage[3*i+2] = src[3*i+2];\n";
static const char *kernel_pattern[] = {
pragma_str,
"#define STYPE %s\n"
@@ -111,9 +109,10 @@ typedef void (*create_program_fn)(std::string &, size_t, ExplicitType, size_t,
typedef int (*test_fn)(cl_device_id, cl_context, cl_command_queue, ExplicitType,
unsigned int, create_program_fn, size_t);
int test_vload(cl_device_id device, cl_context context, cl_command_queue queue,
ExplicitType type, unsigned int vecSize,
create_program_fn createFn, size_t bufferSize)
static int test_vload(cl_device_id device, cl_context context,
cl_command_queue queue, ExplicitType type,
unsigned int vecSize, create_program_fn createFn,
size_t bufferSize)
{
clProgramWrapper program;
clKernelWrapper kernel;
@@ -278,8 +277,9 @@ int test_vload(cl_device_id device, cl_context context, cl_command_queue queue,
}
template <test_fn test_func_ptr>
int test_vset(cl_device_id device, cl_context context, cl_command_queue queue,
create_program_fn createFn, size_t bufferSize)
static int test_vset(cl_device_id device, cl_context context,
cl_command_queue queue, create_program_fn createFn,
size_t bufferSize)
{
std::vector<ExplicitType> vecType = { kChar, kUChar, kShort, kUShort,
kInt, kUInt, kLong, kULong,
@@ -342,7 +342,7 @@ void create_global_load_code(std::string &destBuffer, size_t inBufferSize,
typeName, (int)inVectorSize, (int)inVectorSize);
}
int test_vload_global(cl_device_id device, cl_context context, cl_command_queue queue, int n_elems )
REGISTER_TEST(vload_global)
{
return test_vset<test_vload>(device, context, queue,
create_global_load_code, 10240);
@@ -372,7 +372,7 @@ void create_local_load_code(std::string &destBuffer, size_t inBufferSize,
(int)inVectorSize, (int)inVectorSize, typeName);
}
int test_vload_local(cl_device_id device, cl_context context, cl_command_queue queue, int n_elems )
REGISTER_TEST(vload_local)
{
// Determine the max size of a local buffer that we can test against
cl_ulong localSize;
@@ -408,7 +408,7 @@ void create_constant_load_code(std::string &destBuffer, size_t inBufferSize,
typeName, (int)inVectorSize, (int)inVectorSize);
}
int test_vload_constant(cl_device_id device, cl_context context, cl_command_queue queue, int n_elems )
REGISTER_TEST(vload_constant)
{
// Determine the max size of a local buffer that we can test against
cl_ulong maxSize;
@@ -448,7 +448,7 @@ void create_private_load_code(std::string &destBuffer, size_t inBufferSize,
(int)inVectorSize, (int)inVectorSize, typeName);
}
int test_vload_private(cl_device_id device, cl_context context, cl_command_queue queue, int n_elems )
REGISTER_TEST(vload_private)
{
// We have no idea how much actual private storage is available, so just pick a reasonable value,
// which is that we can fit at least two 16-element long, which is 2*8 bytes * 16 = 256 bytes
@@ -459,9 +459,10 @@ int test_vload_private(cl_device_id device, cl_context context, cl_command_queue
///////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
#pragma mark -------------------- vstore harness --------------------------
int test_vstore(cl_device_id device, cl_context context, cl_command_queue queue,
ExplicitType type, unsigned int vecSize,
create_program_fn createFn, size_t bufferSize)
static int test_vstore(cl_device_id device, cl_context context,
cl_command_queue queue, ExplicitType type,
unsigned int vecSize, create_program_fn createFn,
size_t bufferSize)
{
clProgramWrapper program;
clKernelWrapper kernel;
@@ -716,7 +717,7 @@ void create_global_store_code(std::string &destBuffer, size_t inBufferSize,
}
}
int test_vstore_global(cl_device_id device, cl_context context, cl_command_queue queue, int n_elems )
REGISTER_TEST(vstore_global)
{
return test_vset<test_vstore>(device, context, queue,
create_global_store_code, 10240);
@@ -805,7 +806,7 @@ void create_local_store_code(std::string &destBuffer, size_t inBufferSize,
}
}
int test_vstore_local(cl_device_id device, cl_context context, cl_command_queue queue, int n_elems )
REGISTER_TEST(vstore_local)
{
// Determine the max size of a local buffer that we can test against
cl_ulong localSize;
@@ -890,13 +891,10 @@ void create_private_store_code(std::string &destBuffer, size_t inBufferSize,
}
}
int test_vstore_private(cl_device_id device, cl_context context, cl_command_queue queue, int n_elems )
REGISTER_TEST(vstore_private)
{
// We have no idea how much actual private storage is available, so just pick a reasonable value,
// which is that we can fit at least two 16-element long, which is 2*8 bytes * 16 = 256 bytes
return test_vset<test_vstore>(device, context, queue,
create_private_store_code, 256);
}

20
test_conformance/basic/test_work_item_functions.cpp
View File

@@ -23,7 +23,6 @@
#include <sys/stat.h>
#include <vector>
#include "procs.h"
#include "harness/conversions.h"
#include "harness/typeWrappers.h"
@@ -551,29 +550,22 @@ struct TestWorkItemFnsOutOfRange
} // anonymous namespace
int test_work_item_functions(cl_device_id deviceID, cl_context context,
cl_command_queue queue, int num_elements)
REGISTER_TEST(work_item_functions)
{
TestWorkItemFns fnct(deviceID, context, queue);
TestWorkItemFns fnct(device, context, queue);
return fnct.Run();
}
int test_work_item_functions_out_of_range(cl_device_id deviceID,
cl_context context,
cl_command_queue queue,
int num_elements)
REGISTER_TEST(work_item_functions_out_of_range)
{
TestWorkItemFnsOutOfRange fnct(deviceID, context, queue,
TestWorkItemFnsOutOfRange fnct(device, context, queue,
outOfRangeWorkItemKernelCode);
return fnct.Run();
}
int test_work_item_functions_out_of_range_hardcoded(cl_device_id deviceID,
cl_context context,
cl_command_queue queue,
int num_elements)
REGISTER_TEST(work_item_functions_out_of_range_hardcoded)
{
TestWorkItemFnsOutOfRange fnct(deviceID, context, queue,
TestWorkItemFnsOutOfRange fnct(device, context, queue,
outOfRangeWorkItemHardcodedKernelCode);
return fnct.Run();
}

20
test_conformance/basic/test_writeimage.cpp
View File

@@ -22,14 +22,11 @@
#include <sys/types.h>
#include <sys/stat.h>
#include "procs.h"
#include <algorithm>
#include <string>
#include <vector>
#include "procs.h"
#include "testBase.h"
namespace {
const char *kernel_source = R"(
@@ -120,9 +117,9 @@ const char *get_mem_flag_name(cl_mem_flags flags)
}
template <typename T>
int test_writeimage(cl_device_id device, cl_context context,
cl_command_queue queue, const cl_image_format *img_format,
cl_mem_flags img_flags)
static int
test_writeimage(cl_device_id device, cl_context context, cl_command_queue queue,
const cl_image_format *img_format, cl_mem_flags img_flags)
{
clMemWrapper streams[2];
clProgramWrapper program;
@@ -207,8 +204,7 @@ bool check_format(cl_device_id device, cl_context context,
&img_format);
}
}
int test_writeimage(cl_device_id device, cl_context context,
cl_command_queue queue, int num_elements)
REGISTER_TEST(writeimage)
{
int err = 0;
const cl_image_format format[] = { { CL_RGBA, CL_UNORM_INT8 },
@@ -230,8 +226,7 @@ int test_writeimage(cl_device_id device, cl_context context,
return err;
}
int test_writeimage_int16(cl_device_id device, cl_context context,
cl_command_queue queue, int num_elements)
REGISTER_TEST(writeimage_int16)
{
int err = 0;
const cl_image_format format = { CL_RGBA, CL_UNORM_INT16 };
@@ -245,8 +240,7 @@ int test_writeimage_int16(cl_device_id device, cl_context context,
return err;
}
int test_writeimage_fp32(cl_device_id device, cl_context context,
cl_command_queue queue, int num_elements)
REGISTER_TEST(writeimage_fp32)
{
int err = 0;
const cl_image_format format = { CL_RGBA, CL_FLOAT };

97
test_conformance/compiler/main.cpp
View File

@@ -13,103 +13,10 @@
// See the License for the specific language governing permissions and
// limitations under the License.
//
#include "harness/compat.h"
#include "harness/testHarness.h"
#include "procs.h"
#include <stdio.h>
#include <string.h>
#if !defined(_WIN32)
#include <unistd.h>
#endif
test_definition test_list[] = {
ADD_TEST(load_program_source),
ADD_TEST(load_multistring_source),
ADD_TEST(load_two_kernel_source),
ADD_TEST(load_null_terminated_source),
ADD_TEST(load_null_terminated_multi_line_source),
ADD_TEST(load_null_terminated_partial_multi_line_source),
ADD_TEST(load_discreet_length_source),
ADD_TEST(get_program_source),
ADD_TEST(get_program_build_info),
ADD_TEST(get_program_info),
ADD_TEST(get_program_info_kernel_names),
ADD_TEST(get_program_info_mult_devices),
ADD_TEST(large_compile),
ADD_TEST(async_build),
ADD_TEST(options_build_optimizations),
ADD_TEST(options_build_macro),
ADD_TEST(options_build_macro_existence),
ADD_TEST(options_include_directory),
ADD_TEST(options_denorm_cache),
ADD_TEST(preprocessor_define_udef),
ADD_TEST(preprocessor_include),
ADD_TEST(preprocessor_line_error),
ADD_TEST(preprocessor_pragma),
ADD_TEST(opencl_c_versions),
ADD_TEST(compiler_defines_for_extensions),
ADD_TEST(image_macro),
ADD_TEST(simple_compile_only),
ADD_TEST(simple_static_compile_only),
ADD_TEST(simple_extern_compile_only),
ADD_TEST(simple_compile_with_callback),
ADD_TEST(simple_embedded_header_compile),
ADD_TEST(simple_link_only),
ADD_TEST(two_file_regular_variable_access),
ADD_TEST(two_file_regular_struct_access),
ADD_TEST(two_file_regular_function_access),
ADD_TEST(simple_link_with_callback),
ADD_TEST(simple_embedded_header_link),
ADD_TEST(execute_after_simple_compile_and_link),
ADD_TEST(execute_after_simple_compile_and_link_no_device_info),
ADD_TEST(execute_after_simple_compile_and_link_with_defines),
ADD_TEST(execute_after_simple_compile_and_link_with_callbacks),
ADD_TEST(execute_after_simple_library_with_link),
ADD_TEST(execute_after_two_file_link),
ADD_TEST(execute_after_embedded_header_link),
ADD_TEST(execute_after_included_header_link),
ADD_TEST(execute_after_serialize_reload_object),
ADD_TEST(execute_after_serialize_reload_library),
ADD_TEST(simple_library_only),
ADD_TEST(simple_library_with_callback),
ADD_TEST(simple_library_with_link),
ADD_TEST(two_file_link),
ADD_TEST(multi_file_libraries),
ADD_TEST(multiple_files),
ADD_TEST(multiple_libraries),
ADD_TEST(multiple_files_multiple_libraries),
ADD_TEST(multiple_embedded_headers),
ADD_TEST(program_binary_type),
ADD_TEST(compile_and_link_status_options_log),
ADD_TEST_VERSION(pragma_unroll, Version(2, 0)),
ADD_TEST_VERSION(features_macro, Version(3, 0)),
ADD_TEST(features_macro_coupling),
ADD_TEST(unload_valid),
// ADD_TEST(unload_invalid), // disabling temporarily, see GitHub #977
ADD_TEST(unload_repeated),
ADD_TEST(unload_compile_unload_link),
ADD_TEST(unload_build_unload_create_kernel),
ADD_TEST(unload_link_different),
ADD_TEST(unload_build_threaded),
ADD_TEST(unload_build_info),
ADD_TEST(unload_program_binaries),
};
const int test_num = ARRAY_SIZE(test_list);
int main(int argc, const char *argv[])
{
return runTestHarness(argc, argv, test_num, test_list, false, 0);
return runTestHarness(argc, argv, test_registry::getInstance().num_tests(),
test_registry::getInstance().definitions(), false, 0);
}

272
test_conformance/compiler/procs.h
View File

@@ -1,272 +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/conversions.h"
#include "harness/errorHelpers.h"
#include "harness/kernelHelpers.h"
#include "harness/mt19937.h"
#include "harness/typeWrappers.h"
// This is a macro rather than a function to be able to use and act like the
// existing test_error macro.
//
// Not all compiler tests need to use this macro, only those that don't use the
// test harness compiler helpers.
#define check_compiler_available(DEVICE) \
{ \
cl_bool compilerAvailable = CL_FALSE; \
cl_int error = clGetDeviceInfo((DEVICE), CL_DEVICE_COMPILER_AVAILABLE, \
sizeof(compilerAvailable), \
&compilerAvailable, NULL); \
test_error(error, "Unable to query CL_DEVICE_COMPILER_AVAILABLE"); \
if (compilerAvailable == CL_FALSE) \
{ \
log_info("Skipping test - no compiler is available.\n"); \
return TEST_SKIPPED_ITSELF; \
} \
}
extern int test_load_program_source(cl_device_id deviceID, cl_context context,
cl_command_queue queue, int num_elements);
extern int test_load_multistring_source(cl_device_id deviceID,
cl_context context,
cl_command_queue queue,
int num_elements);
extern int test_load_two_kernel_source(cl_device_id deviceID,
cl_context context,
cl_command_queue queue,
int num_elements);
extern int test_load_null_terminated_source(cl_device_id deviceID,
cl_context context,
cl_command_queue queue,
int num_elements);
extern int test_load_null_terminated_multi_line_source(cl_device_id deviceID,
cl_context context,
cl_command_queue queue,
int num_elements);
extern int test_load_null_terminated_partial_multi_line_source(
cl_device_id deviceID, cl_context context, cl_command_queue queue,
int num_elements);
extern int test_load_discreet_length_source(cl_device_id deviceID,
cl_context context,
cl_command_queue queue,
int num_elements);
extern int test_get_program_source(cl_device_id deviceID, cl_context context,
cl_command_queue queue, int num_elements);
extern int test_get_program_build_info(cl_device_id deviceID,
cl_context context,
cl_command_queue queue,
int num_elements);
extern int test_get_program_info(cl_device_id deviceID, cl_context context,
cl_command_queue queue, int num_elements);
extern int test_get_program_info_kernel_names(cl_device_id deviceID,
cl_context context,
cl_command_queue queue,
int num_elements);
extern int test_get_program_info_mult_devices(cl_device_id deviceID,
cl_context context,
cl_command_queue queue,
int num_elements);
extern int test_large_compile(cl_device_id deviceID, cl_context context,
cl_command_queue queue, int num_elements);
extern int test_async_build(cl_device_id deviceID, cl_context context,
cl_command_queue queue, int num_elements);
extern int test_options_build_optimizations(cl_device_id deviceID,
cl_context context,
cl_command_queue queue,
int num_elements);
extern int test_options_build_macro(cl_device_id deviceID, cl_context context,
cl_command_queue queue, int num_elements);
extern int test_options_build_macro_existence(cl_device_id deviceID,
cl_context context,
cl_command_queue queue,
int num_elements);
extern int test_options_include_directory(cl_device_id deviceID,
cl_context context,
cl_command_queue queue,
int num_elements);
extern int test_options_denorm_cache(cl_device_id deviceID, cl_context context,
cl_command_queue queue, int num_elements);
extern int test_preprocessor_define_udef(cl_device_id deviceID,
cl_context context,
cl_command_queue queue,
int num_elements);
extern int test_preprocessor_include(cl_device_id deviceID, cl_context context,
cl_command_queue queue, int num_elements);
extern int test_preprocessor_line_error(cl_device_id deviceID,
cl_context context,
cl_command_queue queue,
int num_elements);
extern int test_preprocessor_pragma(cl_device_id deviceID, cl_context context,
cl_command_queue queue, int num_elements);
extern int test_opencl_c_versions(cl_device_id device, cl_context context,
cl_command_queue queue, int num_elements);
extern int test_compiler_defines_for_extensions(cl_device_id device,
cl_context context,
cl_command_queue queue,
int n_elems);
extern int test_image_macro(cl_device_id deviceID, cl_context context,
cl_command_queue queue, int num_elements);
extern int test_simple_compile_only(cl_device_id deviceID, cl_context context,
cl_command_queue queue, int num_elements);
extern int test_simple_static_compile_only(cl_device_id deviceID,
cl_context context,
cl_command_queue queue,
int num_elements);
extern int test_simple_extern_compile_only(cl_device_id deviceID,
cl_context context,
cl_command_queue queue,
int num_elements);
extern int test_simple_compile_with_callback(cl_device_id deviceID,
cl_context context,
cl_command_queue queue,
int num_elements);
extern int test_simple_embedded_header_compile(cl_device_id deviceID,
cl_context context,
cl_command_queue queue,
int num_elements);
extern int test_simple_link_only(cl_device_id deviceID, cl_context context,
cl_command_queue queue, int num_elements);
extern int test_two_file_regular_variable_access(cl_device_id deviceID,
cl_context context,
cl_command_queue queue,
int num_elements);
extern int test_two_file_regular_struct_access(cl_device_id deviceID,
cl_context context,
cl_command_queue queue,
int num_elements);
extern int test_two_file_regular_function_access(cl_device_id deviceID,
cl_context context,
cl_command_queue queue,
int num_elements);
extern int test_simple_link_with_callback(cl_device_id deviceID,
cl_context context,
cl_command_queue queue,
int num_elements);
extern int test_simple_embedded_header_link(cl_device_id deviceID,
cl_context context,
cl_command_queue queue,
int num_elements);
extern int test_execute_after_simple_compile_and_link(cl_device_id deviceID,
cl_context context,
cl_command_queue queue,
int num_elements);
extern int test_execute_after_simple_compile_and_link_no_device_info(
cl_device_id deviceID, cl_context context, cl_command_queue queue,
int num_elements);
extern int test_execute_after_simple_compile_and_link_with_defines(
cl_device_id deviceID, cl_context context, cl_command_queue queue,
int num_elements);
extern int test_execute_after_simple_compile_and_link_with_callbacks(
cl_device_id deviceID, cl_context context, cl_command_queue queue,
int num_elements);
extern int test_execute_after_simple_library_with_link(cl_device_id deviceID,
cl_context context,
cl_command_queue queue,
int num_elements);
extern int test_execute_after_two_file_link(cl_device_id deviceID,
cl_context context,
cl_command_queue queue,
int num_elements);
extern int test_execute_after_embedded_header_link(cl_device_id deviceID,
cl_context context,
cl_command_queue queue,
int num_elements);
extern int test_execute_after_included_header_link(cl_device_id deviceID,
cl_context context,
cl_command_queue queue,
int num_elements);
extern int test_execute_after_serialize_reload_object(cl_device_id deviceID,
cl_context context,
cl_command_queue queue,
int num_elements);
extern int test_execute_after_serialize_reload_library(cl_device_id deviceID,
cl_context context,
cl_command_queue queue,
int num_elements);
extern int test_simple_library_only(cl_device_id deviceID, cl_context context,
cl_command_queue queue, int num_elements);
extern int test_simple_library_with_callback(cl_device_id deviceID,
cl_context context,
cl_command_queue queue,
int num_elements);
extern int test_simple_library_with_link(cl_device_id deviceID,
cl_context context,
cl_command_queue queue,
int num_elements);
extern int test_two_file_link(cl_device_id deviceID, cl_context context,
cl_command_queue queue, int num_elements);
extern int test_multi_file_libraries(cl_device_id deviceID, cl_context context,
cl_command_queue queue, int num_elements);
extern int test_multiple_libraries(cl_device_id deviceID, cl_context context,
cl_command_queue queue, int num_elements);
extern int test_multiple_files(cl_device_id deviceID, cl_context context,
cl_command_queue queue, int num_elements);
extern int test_multiple_files_multiple_libraries(cl_device_id deviceID,
cl_context context,
cl_command_queue queue,
int num_elements);
extern int test_multiple_embedded_headers(cl_device_id deviceID,
cl_context context,
cl_command_queue queue,
int num_elements);
extern int test_program_binary_type(cl_device_id deviceID, cl_context context,
cl_command_queue queue, int num_elements);
extern int test_compile_and_link_status_options_log(cl_device_id deviceID,
cl_context context,
cl_command_queue queue,
int num_elements);
extern int test_pragma_unroll(cl_device_id deviceID, cl_context context,
cl_command_queue queue, int num_elements);
extern int test_features_macro(cl_device_id deviceID, cl_context context,
cl_command_queue queue, int num_elements);
extern int test_features_macro_coupling(cl_device_id deviceID,
cl_context context,
cl_command_queue queue,
int num_elements);
extern int test_unload_valid(cl_device_id deviceID, cl_context context,
cl_command_queue queue, int num_elements);
extern int test_unload_invalid(cl_device_id deviceID, cl_context context,
cl_command_queue queue, int num_elements);
extern int test_unload_repeated(cl_device_id deviceID, cl_context context,
cl_command_queue queue, int num_elements);
extern int test_unload_compile_unload_link(cl_device_id deviceID,
cl_context context,
cl_command_queue queue,
int num_elements);
extern int test_unload_build_unload_create_kernel(cl_device_id deviceID,
cl_context context,
cl_command_queue queue,
int num_elements);
extern int test_unload_link_different(cl_device_id deviceID, cl_context context,
cl_command_queue queue, int num_elements);
extern int test_unload_build_threaded(cl_device_id deviceID, cl_context context,
cl_command_queue queue, int num_elements);
extern int test_unload_build_info(cl_device_id deviceID, cl_context context,
cl_command_queue queue, int num_elements);
extern int test_unload_program_binaries(cl_device_id deviceID,
cl_context context,
cl_command_queue queue,
int num_elements);

32
test_conformance/compiler/testBase.h
View File

@@ -16,14 +16,34 @@
#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 <vector>
#include <CL/cl.h>
#include "harness/conversions.h"
#include "harness/testHarness.h"
#include "harness/typeWrappers.h"
// This is a macro rather than a function to be able to use and act like the
// existing test_error macro.
//
// Not all compiler tests need to use this macro, only those that don't use the
// test harness compiler helpers.
#define check_compiler_available(DEVICE) \
{ \
cl_bool compilerAvailable = CL_FALSE; \
cl_int error = clGetDeviceInfo((DEVICE), CL_DEVICE_COMPILER_AVAILABLE, \
sizeof(compilerAvailable), \
&compilerAvailable, NULL); \
test_error(error, "Unable to query CL_DEVICE_COMPILER_AVAILABLE"); \
if (compilerAvailable == CL_FALSE) \
{ \
log_info("Skipping test - no compiler is available.\n"); \
return TEST_SKIPPED_ITSELF; \
} \
}
#include "procs.h"
// scope guard helper to ensure proper releasing of sub devices
struct SubDevicesScopeGuarded

7
test_conformance/compiler/test_async_build.cpp
View File

@@ -94,8 +94,7 @@ void CL_CALLBACK test_notify_build_complete(cl_program program, void *userData)
}
}
int test_async_build(cl_device_id deviceID, cl_context context,
cl_command_queue queue, int num_elements)
REGISTER_TEST(async_build)
{
cl_int error;
@@ -119,9 +118,9 @@ int test_async_build(cl_device_id deviceID, cl_context context,
test_error(error, "Unable to create program from source");
// Start an asynchronous build, registering the completion callback
TestData testData = { deviceID, testDef.expectedStatus };
TestData testData = { device, testDef.expectedStatus };
callbackResult = 0;
error = clBuildProgram(program, 1, &deviceID, NULL,
error = clBuildProgram(program, 1, &device, NULL,
test_notify_build_complete, (void *)&testData);
// Allow implementations to return synchronous build failures.
// They still need to call the callback.

91
test_conformance/compiler/test_build_helpers.cpp
View File

@@ -88,7 +88,7 @@ __kernel void sample_test_C(__global float *src, __global int *dst)
)";
int test_load_program_source(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
REGISTER_TEST(load_program_source)
{
int error;
clProgramWrapper program;
@@ -132,7 +132,7 @@ int test_load_program_source(cl_device_id deviceID, cl_context context, cl_comma
return 0;
}
int test_load_multistring_source(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
REGISTER_TEST(load_multistring_source)
{
int error;
clProgramWrapper program;
@@ -159,7 +159,7 @@ int test_load_multistring_source(cl_device_id deviceID, cl_context context, cl_c
}
/* Try compiling */
error = clBuildProgram( program, 1, &deviceID, NULL, NULL, NULL );
error = clBuildProgram(program, 1, &device, NULL, NULL, NULL);
test_error( error, "Unable to build multi-line program source" );
/* Should probably check binary here to verify the same results... */
@@ -169,7 +169,7 @@ int test_load_multistring_source(cl_device_id deviceID, cl_context context, cl_c
return 0;
}
int test_load_two_kernel_source(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
REGISTER_TEST(load_two_kernel_source)
{
int error;
cl_program program;
@@ -195,7 +195,7 @@ int test_load_two_kernel_source(cl_device_id deviceID, cl_context context, cl_co
}
/* Try compiling */
error = clBuildProgram( program, 1, &deviceID, NULL, NULL, NULL );
error = clBuildProgram(program, 1, &device, NULL, NULL, NULL);
test_error( error, "Unable to build two-kernel program source" );
/* Should probably check binary here to verify the same results... */
@@ -207,7 +207,7 @@ int test_load_two_kernel_source(cl_device_id deviceID, cl_context context, cl_co
return 0;
}
int test_load_null_terminated_source(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
REGISTER_TEST(load_null_terminated_source)
{
int error;
cl_program program;
@@ -222,7 +222,7 @@ int test_load_null_terminated_source(cl_device_id deviceID, cl_context context,
}
/* Try compiling */
error = clBuildProgram( program, 1, &deviceID, NULL, NULL, NULL );
error = clBuildProgram(program, 1, &device, NULL, NULL, NULL);
test_error( error, "Unable to build null-terminated program source" );
/* Should probably check binary here to verify the same results... */
@@ -234,7 +234,7 @@ int test_load_null_terminated_source(cl_device_id deviceID, cl_context context,
return 0;
}
int test_load_null_terminated_multi_line_source(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
REGISTER_TEST(load_null_terminated_multi_line_source)
{
int error;
cl_program program;
@@ -252,7 +252,7 @@ int test_load_null_terminated_multi_line_source(cl_device_id deviceID, cl_contex
}
/* Try compiling */
error = clBuildProgram( program, 1, &deviceID, NULL, NULL, NULL );
error = clBuildProgram(program, 1, &device, NULL, NULL, NULL);
test_error( error, "Unable to build null-terminated program source" );
/* Should probably check binary here to verify the same results... */
@@ -264,7 +264,7 @@ int test_load_null_terminated_multi_line_source(cl_device_id deviceID, cl_contex
return 0;
}
int test_load_discreet_length_source(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
REGISTER_TEST(load_discreet_length_source)
{
int error;
cl_program program;
@@ -295,7 +295,7 @@ int test_load_discreet_length_source(cl_device_id deviceID, cl_context context,
}
/* Try compiling */
error = clBuildProgram( program, 1, &deviceID, NULL, NULL, NULL );
error = clBuildProgram(program, 1, &device, NULL, NULL, NULL);
test_error( error, "Unable to build null-terminated program source" );
/* Should probably check binary here to verify the same results... */
@@ -307,7 +307,7 @@ int test_load_discreet_length_source(cl_device_id deviceID, cl_context context,
return 0;
}
int test_load_null_terminated_partial_multi_line_source(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
REGISTER_TEST(load_null_terminated_partial_multi_line_source)
{
int error;
cl_program program;
@@ -337,7 +337,7 @@ int test_load_null_terminated_partial_multi_line_source(cl_device_id deviceID, c
}
/* Try compiling */
error = clBuildProgram( program, 1, &deviceID, NULL, NULL, NULL );
error = clBuildProgram(program, 1, &device, NULL, NULL, NULL);
test_error( error, "Unable to build null-terminated program source" );
/* Should probably check binary here to verify the same results... */
@@ -349,7 +349,7 @@ int test_load_null_terminated_partial_multi_line_source(cl_device_id deviceID, c
return 0;
}
int test_get_program_info(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
REGISTER_TEST(get_program_info)
{
int error;
cl_program program;
@@ -373,7 +373,7 @@ int test_get_program_info(cl_device_id deviceID, cl_context context, cl_command_
test_error( error, "Unable to get device of program" );
/* Object comparability test. */
test_assert_error(device1 == deviceID,
test_assert_error(device1 == device,
"Unexpected result returned by CL_PROGRAM_DEVICES query");
cl_uint devCount;
@@ -439,9 +439,7 @@ int test_get_program_info(cl_device_id deviceID, cl_context context, cl_command_
return 0;
}
int test_get_program_info_kernel_names(cl_device_id deviceID,
cl_context context,
cl_command_queue queue, int num_elements)
REGISTER_TEST(get_program_info_kernel_names)
{
int error = CL_SUCCESS;
size_t total_kernels = 0;
@@ -474,8 +472,7 @@ int test_get_program_info_kernel_names(cl_device_id deviceID,
// returned. Query CL_PROGRAM_KERNEL_NAMES and check that the right
// kernel names are returned.
{
error =
clBuildProgram(program, 1, &deviceID, nullptr, nullptr, nullptr);
error = clBuildProgram(program, 1, &device, nullptr, nullptr, nullptr);
test_error(error, "clBuildProgram failed");
error = clGetProgramInfo(program, CL_PROGRAM_NUM_KERNELS,
@@ -517,7 +514,7 @@ int test_get_program_info_kernel_names(cl_device_id deviceID,
// kernel names are returned.
{
const char *build_options = "-DUSE_SAMPLE_TEST_B";
error = clBuildProgram(program, 1, &deviceID, build_options, nullptr,
error = clBuildProgram(program, 1, &device, build_options, nullptr,
nullptr);
test_error(error, "clBuildProgram failed");
@@ -553,14 +550,12 @@ int test_get_program_info_kernel_names(cl_device_id deviceID,
return CL_SUCCESS;
}
int test_get_program_info_mult_devices(cl_device_id deviceID,
cl_context context,
cl_command_queue queue, int num_elements)
REGISTER_TEST(get_program_info_mult_devices)
{
size_t size = 0;
// query multi-device context and perform objects comparability test
cl_int err = clGetDeviceInfo(deviceID, CL_DEVICE_PARTITION_PROPERTIES, 0,
cl_int err = clGetDeviceInfo(device, CL_DEVICE_PARTITION_PROPERTIES, 0,
nullptr, &size);
test_error_fail(err, "clGetDeviceInfo failed");
@@ -572,7 +567,7 @@ int test_get_program_info_mult_devices(cl_device_id deviceID,
std::vector<cl_device_partition_property> supported_props(
size / sizeof(cl_device_partition_property), 0);
err = clGetDeviceInfo(deviceID, CL_DEVICE_PARTITION_PROPERTIES,
err = clGetDeviceInfo(device, CL_DEVICE_PARTITION_PROPERTIES,
supported_props.size()
* sizeof(cl_device_partition_property),
supported_props.data(), &size);
@@ -585,7 +580,7 @@ int test_get_program_info_mult_devices(cl_device_id deviceID,
}
cl_uint maxComputeUnits = 0;
err = clGetDeviceInfo(deviceID, CL_DEVICE_MAX_COMPUTE_UNITS,
err = clGetDeviceInfo(device, CL_DEVICE_MAX_COMPUTE_UNITS,
sizeof(maxComputeUnits), &maxComputeUnits, nullptr);
test_error_ret(err, "Unable to get maximal number of compute units",
TEST_FAIL);
@@ -607,14 +602,14 @@ int test_get_program_info_mult_devices(cl_device_id deviceID,
if (sup_prop == prop[0])
{
// how many sub-devices can we create?
err = clCreateSubDevices(deviceID, prop.data(), 0, nullptr,
err = clCreateSubDevices(device, prop.data(), 0, nullptr,
&num_devices);
test_error_fail(err, "clCreateSubDevices failed");
if (num_devices < 2) continue;
// get the list of subDevices
scope_guard.reset(new SubDevicesScopeGuarded(num_devices));
err = clCreateSubDevices(deviceID, prop.data(), num_devices,
err = clCreateSubDevices(device, prop.data(), num_devices,
scope_guard->sub_devices.data(),
&num_devices);
test_error_fail(err, "clCreateSubDevices failed");
@@ -681,7 +676,7 @@ int test_get_program_info_mult_devices(cl_device_id deviceID,
return TEST_PASS;
}
int test_get_program_source(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
REGISTER_TEST(get_program_source)
{
cl_program program;
int error;
@@ -736,7 +731,7 @@ int test_get_program_source(cl_device_id deviceID, cl_context context, cl_comman
return 0;
}
int test_get_program_build_info(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
REGISTER_TEST(get_program_build_info)
{
cl_program program;
int error;
@@ -753,7 +748,8 @@ int test_get_program_build_info(cl_device_id deviceID, cl_context context, cl_co
}
/* Make sure getting the length works */
error = clGetProgramBuildInfo( program, deviceID, CL_PROGRAM_BUILD_STATUS, 0, NULL, &length );
error = clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_STATUS, 0,
NULL, &length);
test_error( error, "Unable to get program build status length" );
if( length != sizeof( status ) )
{
@@ -762,10 +758,11 @@ int test_get_program_build_info(cl_device_id deviceID, cl_context context, cl_co
}
/* Now actually build it and verify the status */
error = clBuildProgram( program, 1, &deviceID, NULL, NULL, NULL );
error = clBuildProgram(program, 1, &device, NULL, NULL, NULL);
test_error( error, "Unable to build program source" );
error = clGetProgramBuildInfo( program, deviceID, CL_PROGRAM_BUILD_STATUS, sizeof( status ), &status, NULL );
error = clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_STATUS,
sizeof(status), &status, NULL);
test_error( error, "Unable to get program build status" );
if( status != CL_BUILD_SUCCESS )
{
@@ -776,7 +773,8 @@ int test_get_program_build_info(cl_device_id deviceID, cl_context context, cl_co
/***** Build log *****/
/* Try getting the length */
error = clGetProgramBuildInfo( program, deviceID, CL_PROGRAM_BUILD_LOG, 0, NULL, &length );
error = clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, 0,
NULL, &length);
test_error( error, "Unable to get program build log length" );
log_info("Build log is %zu long.\n", length);
@@ -784,7 +782,8 @@ int test_get_program_build_info(cl_device_id deviceID, cl_context context, cl_co
buffer = (char*)malloc(length);
/* Try normal source */
error = clGetProgramBuildInfo( program, deviceID, CL_PROGRAM_BUILD_LOG, length, buffer, NULL );
error = clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, length,
buffer, NULL);
test_error( error, "Unable to get program build log" );
if( buffer[length-1] != '\0' )
@@ -794,23 +793,27 @@ int test_get_program_build_info(cl_device_id deviceID, cl_context context, cl_co
}
/* Try both at once */
error = clGetProgramBuildInfo( program, deviceID, CL_PROGRAM_BUILD_LOG, length, buffer, &newLength );
error = clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, length,
buffer, &newLength);
test_error( error, "Unable to get program build log" );
free(buffer);
/***** Build options *****/
error = clGetProgramBuildInfo( program, deviceID, CL_PROGRAM_BUILD_OPTIONS, 0, NULL, &length );
error = clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_OPTIONS, 0,
NULL, &length);
test_error( error, "Unable to get program build options length" );
buffer = (char*)malloc(length);
/* Try normal source */
error = clGetProgramBuildInfo( program, deviceID, CL_PROGRAM_BUILD_OPTIONS, length, buffer, NULL );
error = clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_OPTIONS,
length, buffer, NULL);
test_error( error, "Unable to get program build options" );
/* Try both at once */
error = clGetProgramBuildInfo( program, deviceID, CL_PROGRAM_BUILD_OPTIONS, length, buffer, &newLength );
error = clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_OPTIONS,
length, buffer, &newLength);
test_error( error, "Unable to get program build options" );
free(buffer);
@@ -826,19 +829,21 @@ int test_get_program_build_info(cl_device_id deviceID, cl_context context, cl_co
return -1;
}
error = clBuildProgram( program, 1, &deviceID, "-cl-opt-disable", NULL, NULL );
error = clBuildProgram(program, 1, &device, "-cl-opt-disable", NULL, NULL);
if( error != CL_SUCCESS )
{
print_error( error, "Building with valid options failed!" );
return -1;
}
error = clGetProgramBuildInfo( program, deviceID, CL_PROGRAM_BUILD_OPTIONS, 0, NULL, &length );
error = clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_OPTIONS, 0,
NULL, &length);
test_error( error, "Unable to get program build options" );
buffer = (char*)malloc(length);
error = clGetProgramBuildInfo( program, deviceID, CL_PROGRAM_BUILD_OPTIONS, length, buffer, NULL );
error = clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_OPTIONS,
length, buffer, NULL);
test_error( error, "Unable to get program build options" );
if( strcmp( (char *)buffer, "-cl-opt-disable" ) != 0 )
{

48
test_conformance/compiler/test_build_options.cpp
View File

@@ -88,7 +88,7 @@ cl_int get_result_from_program( cl_context context, cl_command_queue queue, cl_p
return CL_SUCCESS;
}
int test_options_build_optimizations(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
REGISTER_TEST(options_build_optimizations)
{
int error;
cl_build_status status;
@@ -105,10 +105,12 @@ int test_options_build_optimizations(cl_device_id deviceID, cl_context context,
/* Build with the macro defined */
log_info("Testing optimization option '%s'\n", optimization_options[i]);
error = clBuildProgram( program, 1, &deviceID, optimization_options[i], NULL, NULL );
error = clBuildProgram(program, 1, &device, optimization_options[i],
NULL, NULL);
test_error( error, "Test program did not properly build" );
error = clGetProgramBuildInfo( program, deviceID, CL_PROGRAM_BUILD_STATUS, sizeof( status ), &status, NULL );
error = clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_STATUS,
sizeof(status), &status, NULL);
test_error( error, "Unable to get program build status" );
if( (int)status != CL_BUILD_SUCCESS )
@@ -121,7 +123,7 @@ int test_options_build_optimizations(cl_device_id deviceID, cl_context context,
return 0;
}
int test_options_build_macro(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
REGISTER_TEST(options_build_macro)
{
int error;
clProgramWrapper program;
@@ -136,10 +138,11 @@ int test_options_build_macro(cl_device_id deviceID, cl_context context, cl_comma
}
/* Build with the macro defined */
error = clBuildProgram( program, 1, &deviceID, "-DTEST_MACRO=1 ", NULL, NULL );
error = clBuildProgram(program, 1, &device, "-DTEST_MACRO=1 ", NULL, NULL);
test_error( error, "Test program did not properly build" );
error = clGetProgramBuildInfo( program, deviceID, CL_PROGRAM_BUILD_STATUS, sizeof( status ), &status, NULL );
error = clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_STATUS,
sizeof(status), &status, NULL);
test_error( error, "Unable to get program build status" );
if( (int)status != CL_BUILD_SUCCESS )
@@ -162,7 +165,7 @@ int test_options_build_macro(cl_device_id deviceID, cl_context context, cl_comma
}
// Rebuild with a different value for the define macro, to make sure caching behaves properly
error = clBuildProgram( program, 1, &deviceID, "-DTEST_MACRO=5 ", NULL, NULL );
error = clBuildProgram(program, 1, &device, "-DTEST_MACRO=5 ", NULL, NULL);
test_error( error, "Test program did not properly rebuild" );
error = get_result_from_program( context, queue, program, &secondResult );
@@ -180,7 +183,7 @@ int test_options_build_macro(cl_device_id deviceID, cl_context context, cl_comma
return 0;
}
int test_options_build_macro_existence(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
REGISTER_TEST(options_build_macro_existence)
{
int error;
clProgramWrapper program;
@@ -195,7 +198,7 @@ int test_options_build_macro_existence(cl_device_id deviceID, cl_context context
}
/* Build without the macro defined */
error = clBuildProgram( program, 1, &deviceID, NULL, NULL, NULL );
error = clBuildProgram(program, 1, &device, NULL, NULL, NULL);
test_error( error, "Test program did not properly build" );
// Go ahead and run the program to verify results
@@ -211,7 +214,7 @@ int test_options_build_macro_existence(cl_device_id deviceID, cl_context context
}
// Now compile again with the macro defined and verify a change in results
error = clBuildProgram( program, 1, &deviceID, "-DTEST_MACRO", NULL, NULL );
error = clBuildProgram(program, 1, &device, "-DTEST_MACRO", NULL, NULL);
test_error( error, "Test program did not properly build" );
error = get_result_from_program( context, queue, program, &secondResult );
@@ -229,7 +232,7 @@ int test_options_build_macro_existence(cl_device_id deviceID, cl_context context
return 0;
}
int test_options_include_directory(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
REGISTER_TEST(options_include_directory)
{
int error;
@@ -252,10 +255,12 @@ int test_options_include_directory(cl_device_id deviceID, cl_context context, cl
include_dir = "-I " + path + sep + "includeTestDirectory";
// log_info("%s\n", include_dir);
error = clBuildProgram( program, 1, &deviceID, include_dir.c_str(), NULL, NULL );
error =
clBuildProgram(program, 1, &device, include_dir.c_str(), NULL, NULL);
test_error( error, "Test program did not properly build" );
error = clGetProgramBuildInfo( program, deviceID, CL_PROGRAM_BUILD_STATUS, sizeof( status ), &status, NULL );
error = clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_STATUS,
sizeof(status), &status, NULL);
test_error( error, "Unable to get program build status" );
if( (int)status != CL_BUILD_SUCCESS )
@@ -278,7 +283,8 @@ int test_options_include_directory(cl_device_id deviceID, cl_context context, cl
// Rebuild with a different include directory
include_dir = "-I " + path + sep + "secondIncludeTestDirectory";
error = clBuildProgram( program, 1, &deviceID, include_dir.c_str(), NULL, NULL );
error =
clBuildProgram(program, 1, &device, include_dir.c_str(), NULL, NULL);
test_error( error, "Test program did not properly rebuild" );
error = get_result_from_program( context, queue, program, &secondResult );
@@ -334,7 +340,7 @@ cl_int get_float_result_from_program( cl_context context, cl_command_queue queue
return CL_SUCCESS;
}
int test_options_denorm_cache(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
REGISTER_TEST(options_denorm_cache)
{
int error;
@@ -344,7 +350,8 @@ int test_options_denorm_cache(cl_device_id deviceID, cl_context context, cl_comm
// If denorms aren't even supported, testing this flag is pointless
cl_device_fp_config floatCaps = 0;
error = clGetDeviceInfo( deviceID, CL_DEVICE_SINGLE_FP_CONFIG, sizeof(floatCaps), &floatCaps, NULL);
error = clGetDeviceInfo(device, CL_DEVICE_SINGLE_FP_CONFIG,
sizeof(floatCaps), &floatCaps, NULL);
test_error( error, "Unable to get device FP config" );
if( ( floatCaps & CL_FP_DENORM ) == 0 )
{
@@ -356,10 +363,12 @@ int test_options_denorm_cache(cl_device_id deviceID, cl_context context, cl_comm
test_error( error, "Unable to create test program" );
// Build first WITH the denorm flush flag
error = clBuildProgram( program, 1, &deviceID, "-cl-denorms-are-zero", NULL, NULL );
error =
clBuildProgram(program, 1, &device, "-cl-denorms-are-zero", NULL, NULL);
test_error( error, "Test program did not properly build" );
error = clGetProgramBuildInfo( program, deviceID, CL_PROGRAM_BUILD_STATUS, sizeof( status ), &status, NULL );
error = clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_STATUS,
sizeof(status), &status, NULL);
test_error( error, "Unable to get program build status" );
if( (int)status != CL_BUILD_SUCCESS )
@@ -382,7 +391,7 @@ int test_options_denorm_cache(cl_device_id deviceID, cl_context context, cl_comm
// valid, there isn't anything we can to do validate results for now
// Rebuild without flushing flag set
error = clBuildProgram( program, 1, &deviceID, NULL, NULL, NULL );
error = clBuildProgram(program, 1, &device, NULL, NULL, NULL);
test_error( error, "Test program did not properly rebuild" );
error = get_float_result_from_program( context, queue, program, *input, *input, &secondResult );
@@ -406,4 +415,3 @@ int test_options_denorm_cache(cl_device_id deviceID, cl_context context, cl_comm
return 0;
}

405
test_conformance/compiler/test_compile.cpp
View File

File diff suppressed because it is too large Load Diff

2
test_conformance/compiler/test_compiler_defines_for_extensions.cpp
View File

@@ -131,7 +131,7 @@ bool string_has_prefix(const char *str, const char *prefix)
return strncmp(str, prefix, strlen(prefix)) == 0;
}
int test_compiler_defines_for_extensions(cl_device_id device, cl_context context, cl_command_queue queue, int n_elems )
REGISTER_TEST(compiler_defines_for_extensions)
{
int error;

115
test_conformance/compiler/test_feature_macro.cpp
View File

@@ -205,10 +205,10 @@ int feature_macro_verify_results(std::string test_macro_name,
return error;
}
int test_feature_macro_atomic_order_acq_rel(cl_device_id deviceID,
cl_context context,
std::string test_macro_name,
cl_bool& supported)
static int test_feature_macro_atomic_order_acq_rel(cl_device_id deviceID,
cl_context context,
std::string test_macro_name,
cl_bool& supported)
{
cl_int error = TEST_FAIL;
cl_bool api_status;
@@ -233,10 +233,10 @@ int test_feature_macro_atomic_order_acq_rel(cl_device_id deviceID,
compiler_status, supported);
}
int test_feature_macro_atomic_order_seq_cst(cl_device_id deviceID,
cl_context context,
std::string test_macro_name,
cl_bool& supported)
static int test_feature_macro_atomic_order_seq_cst(cl_device_id deviceID,
cl_context context,
std::string test_macro_name,
cl_bool& supported)
{
cl_int error = TEST_FAIL;
cl_bool api_status;
@@ -262,10 +262,10 @@ int test_feature_macro_atomic_order_seq_cst(cl_device_id deviceID,
compiler_status, supported);
}
int test_feature_macro_atomic_scope_device(cl_device_id deviceID,
cl_context context,
std::string test_macro_name,
cl_bool& supported)
static int test_feature_macro_atomic_scope_device(cl_device_id deviceID,
cl_context context,
std::string test_macro_name,
cl_bool& supported)
{
cl_int error = TEST_FAIL;
cl_bool api_status;
@@ -289,10 +289,9 @@ int test_feature_macro_atomic_scope_device(cl_device_id deviceID,
compiler_status, supported);
}
int test_feature_macro_atomic_scope_all_devices(cl_device_id deviceID,
cl_context context,
std::string test_macro_name,
cl_bool& supported)
static int test_feature_macro_atomic_scope_all_devices(
cl_device_id deviceID, cl_context context, std::string test_macro_name,
cl_bool& supported)
{
cl_int error = TEST_FAIL;
cl_bool api_status;
@@ -316,10 +315,10 @@ int test_feature_macro_atomic_scope_all_devices(cl_device_id deviceID,
compiler_status, supported);
}
int test_feature_macro_3d_image_writes(cl_device_id deviceID,
cl_context context,
std::string test_macro_name,
cl_bool& supported)
static int test_feature_macro_3d_image_writes(cl_device_id deviceID,
cl_context context,
std::string test_macro_name,
cl_bool& supported)
{
cl_int error = TEST_FAIL;
cl_bool api_status;
@@ -343,9 +342,10 @@ int test_feature_macro_3d_image_writes(cl_device_id deviceID,
compiler_status, supported);
}
int test_feature_macro_device_enqueue(cl_device_id deviceID, cl_context context,
std::string test_macro_name,
cl_bool& supported)
static int test_feature_macro_device_enqueue(cl_device_id deviceID,
cl_context context,
std::string test_macro_name,
cl_bool& supported)
{
cl_int error = TEST_FAIL;
cl_bool api_status;
@@ -371,10 +371,10 @@ int test_feature_macro_device_enqueue(cl_device_id deviceID, cl_context context,
compiler_status, supported);
}
int test_feature_macro_generic_address_space(cl_device_id deviceID,
cl_context context,
std::string test_macro_name,
cl_bool& supported)
static int test_feature_macro_generic_address_space(cl_device_id deviceID,
cl_context context,
std::string test_macro_name,
cl_bool& supported)
{
cl_int error = TEST_FAIL;
cl_bool api_status;
@@ -398,8 +398,9 @@ int test_feature_macro_generic_address_space(cl_device_id deviceID,
compiler_status, supported);
}
int test_feature_macro_pipes(cl_device_id deviceID, cl_context context,
std::string test_macro_name, cl_bool& supported)
static int test_feature_macro_pipes(cl_device_id deviceID, cl_context context,
std::string test_macro_name,
cl_bool& supported)
{
cl_int error = TEST_FAIL;
cl_bool api_status;
@@ -423,7 +424,7 @@ int test_feature_macro_pipes(cl_device_id deviceID, cl_context context,
compiler_status, supported);
}
int test_feature_macro_program_scope_global_variables(
static int test_feature_macro_program_scope_global_variables(
cl_device_id deviceID, cl_context context, std::string test_macro_name,
cl_bool& supported)
{
@@ -449,10 +450,10 @@ int test_feature_macro_program_scope_global_variables(
compiler_status, supported);
}
int test_feature_macro_read_write_images(cl_device_id deviceID,
cl_context context,
std::string test_macro_name,
cl_bool& supported)
static int test_feature_macro_read_write_images(cl_device_id deviceID,
cl_context context,
std::string test_macro_name,
cl_bool& supported)
{
cl_int error = TEST_FAIL;
cl_bool api_status;
@@ -476,9 +477,10 @@ int test_feature_macro_read_write_images(cl_device_id deviceID,
compiler_status, supported);
}
int test_feature_macro_subgroups(cl_device_id deviceID, cl_context context,
std::string test_macro_name,
cl_bool& supported)
static int test_feature_macro_subgroups(cl_device_id deviceID,
cl_context context,
std::string test_macro_name,
cl_bool& supported)
{
cl_int error = TEST_FAIL;
cl_bool api_status;
@@ -502,7 +504,7 @@ int test_feature_macro_subgroups(cl_device_id deviceID, cl_context context,
compiler_status, supported);
}
int test_feature_macro_work_group_collective_functions(
static int test_feature_macro_work_group_collective_functions(
cl_device_id deviceID, cl_context context, std::string test_macro_name,
cl_bool& supported)
{
@@ -529,8 +531,9 @@ int test_feature_macro_work_group_collective_functions(
compiler_status, supported);
}
int test_feature_macro_images(cl_device_id deviceID, cl_context context,
std::string test_macro_name, cl_bool& supported)
static int test_feature_macro_images(cl_device_id deviceID, cl_context context,
std::string test_macro_name,
cl_bool& supported)
{
cl_int error = TEST_FAIL;
cl_bool api_status;
@@ -554,8 +557,9 @@ int test_feature_macro_images(cl_device_id deviceID, cl_context context,
compiler_status, supported);
}
int test_feature_macro_fp64(cl_device_id deviceID, cl_context context,
std::string test_macro_name, cl_bool& supported)
static int test_feature_macro_fp64(cl_device_id deviceID, cl_context context,
std::string test_macro_name,
cl_bool& supported)
{
cl_int error = TEST_FAIL;
cl_bool api_status;
@@ -580,7 +584,7 @@ int test_feature_macro_fp64(cl_device_id deviceID, cl_context context,
compiler_status, supported);
}
int test_feature_macro_integer_dot_product_input_4x8bit_packed(
static int test_feature_macro_integer_dot_product_input_4x8bit_packed(
cl_device_id deviceID, cl_context context, std::string test_macro_name,
cl_bool& supported)
{
@@ -616,7 +620,7 @@ int test_feature_macro_integer_dot_product_input_4x8bit_packed(
compiler_status, supported);
}
int test_feature_macro_integer_dot_product_input_4x8bit(
static int test_feature_macro_integer_dot_product_input_4x8bit(
cl_device_id deviceID, cl_context context, std::string test_macro_name,
cl_bool& supported)
{
@@ -652,8 +656,9 @@ int test_feature_macro_integer_dot_product_input_4x8bit(
compiler_status, supported);
}
int test_feature_macro_int64(cl_device_id deviceID, cl_context context,
std::string test_macro_name, cl_bool& supported)
static int test_feature_macro_int64(cl_device_id deviceID, cl_context context,
std::string test_macro_name,
cl_bool& supported)
{
cl_int error = TEST_FAIL;
cl_bool api_status;
@@ -722,8 +727,8 @@ int test_feature_macro_int64(cl_device_id deviceID, cl_context context,
compiler_status, supported);
}
int test_consistency_c_features_list(cl_device_id deviceID,
std::vector<std::string> vec_to_cmp)
static int test_consistency_c_features_list(cl_device_id deviceID,
std::vector<std::string> vec_to_cmp)
{
log_info("\nComparison list of features: CL_DEVICE_OPENCL_C_FEATURES vs "
"API/compiler queries.\n");
@@ -792,13 +797,12 @@ int test_consistency_c_features_list(cl_device_id deviceID,
#define NEW_FEATURE_MACRO_TEST(feat) \
test_macro_name = "__opencl_c_" #feat; \
error |= test_feature_macro_##feat(deviceID, context, test_macro_name, \
error |= test_feature_macro_##feat(device, context, test_macro_name, \
supported); \
if (supported) supported_features_vec.push_back(test_macro_name);
int test_features_macro(cl_device_id deviceID, cl_context context,
cl_command_queue queue, int num_elements)
REGISTER_TEST_VERSION(features_macro, Version(3, 0))
{
// Note: Not checking that the feature array is empty for the compiler not
@@ -806,7 +810,7 @@ int test_features_macro(cl_device_id deviceID, cl_context context,
// support compilation from OpenCL C source, this query may return an empty
// array." It "may" return an empty array implies that an implementation
// also "may not".
check_compiler_available(deviceID);
check_compiler_available(device);
int error = TEST_PASS;
cl_bool supported = CL_FALSE;
@@ -830,17 +834,16 @@ int test_features_macro(cl_device_id deviceID, cl_context context,
NEW_FEATURE_MACRO_TEST(integer_dot_product_input_4x8bit);
NEW_FEATURE_MACRO_TEST(integer_dot_product_input_4x8bit_packed);
error |= test_consistency_c_features_list(deviceID, supported_features_vec);
error |= test_consistency_c_features_list(device, supported_features_vec);
return error;
}
// This test checks that a supported feature comes with other required features
int test_features_macro_coupling(cl_device_id deviceID, cl_context context,
cl_command_queue queue, int num_elements)
REGISTER_TEST(features_macro_coupling)
{
OpenCLCFeatures features;
int error = get_device_cl_c_features(deviceID, features);
int error = get_device_cl_c_features(device, features);
if (error)
{
log_error("Couldn't query OpenCL C features for the device!\n");

12
test_conformance/compiler/test_image_macro.cpp
View File

@@ -36,21 +36,22 @@ const char * image_not_supported_source = "kernel void not_enabled(global int *
"\r\n } \r\n";
int test_image_macro(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
REGISTER_TEST(image_macro)
{
cl_bool image_support;
char buf[256];
int status;
cl_program program;
status = clGetDeviceInfo( deviceID, CL_DEVICE_NAME, sizeof( buf ), buf, NULL );
status = clGetDeviceInfo(device, CL_DEVICE_NAME, sizeof(buf), buf, NULL);
if( status )
{
log_error( "getting device info (name): %d\n", status );
exit(-1);
}
status = clGetDeviceInfo( deviceID, CL_DEVICE_IMAGE_SUPPORT, sizeof( image_support ), &image_support, NULL );
status = clGetDeviceInfo(device, CL_DEVICE_IMAGE_SUPPORT,
sizeof(image_support), &image_support, NULL);
if( status )
{
log_error( "getting device info (image support): %d\n", status );
@@ -67,7 +68,7 @@ int test_image_macro(cl_device_id deviceID, cl_context context, cl_command_queue
return status;
}
status = clBuildProgram( program, 1, &deviceID, NULL, NULL, NULL );
status = clBuildProgram(program, 1, &device, NULL, NULL, NULL);
if( status )
log_error("CL_DEVICE_IMAGE_SUPPORT is set, __IMAGE_SUPPORT__ macro not set \n");
else
@@ -82,7 +83,7 @@ int test_image_macro(cl_device_id deviceID, cl_context context, cl_command_queue
return status;
}
status = clBuildProgram( program, 1, &deviceID, NULL, NULL, NULL );
status = clBuildProgram(program, 1, &device, NULL, NULL, NULL);
if( status )
log_error("CL_DEVICE_IMAGE_SUPPORT not set, __IMAGE_SUPPORT__ macro is set \n");
else
@@ -98,4 +99,3 @@ int test_image_macro(cl_device_id deviceID, cl_context context, cl_command_queue
return status;
}

3
test_conformance/compiler/test_opencl_c_versions.cpp
View File

@@ -286,8 +286,7 @@ static int test_CL_DEVICE_OPENCL_C_VERSION_versions(cl_device_id device,
return TEST_PASS;
}
int test_opencl_c_versions(cl_device_id device, cl_context context,
cl_command_queue queue, int num_elements)
REGISTER_TEST(opencl_c_versions)
{
check_compiler_available(device);

91
test_conformance/compiler/test_pragma_unroll.cpp
View File

@@ -248,49 +248,60 @@ const char *pragma_unroll_kernels[] = {
"}\n",
};
int test_pragma_unroll(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
const size_t ELEMENT_NUM = 100;
const size_t KERNEL_NUM = 24;
REGISTER_TEST_VERSION(pragma_unroll, Version(2, 0))
{
const size_t ELEMENT_NUM = 100;
const size_t KERNEL_NUM = 24;
cl_int error;
cl_int error;
//execute all kernels and check if the results are as expected
for (size_t kernelIdx = 0; kernelIdx < KERNEL_NUM; ++kernelIdx) {
clProgramWrapper program;
clKernelWrapper kernel;
if (create_single_kernel_helper(
context, &program, &kernel, 1,
(const char **)&pragma_unroll_kernels[kernelIdx], "pragma_unroll"))
// execute all kernels and check if the results are as expected
for (size_t kernelIdx = 0; kernelIdx < KERNEL_NUM; ++kernelIdx)
{
log_error("The program we attempted to compile was: \n%s\n",
pragma_unroll_kernels[kernelIdx]);
return -1;
clProgramWrapper program;
clKernelWrapper kernel;
if (create_single_kernel_helper(
context, &program, &kernel, 1,
(const char **)&pragma_unroll_kernels[kernelIdx],
"pragma_unroll"))
{
log_error("The program we attempted to compile was: \n%s\n",
pragma_unroll_kernels[kernelIdx]);
return -1;
}
clMemWrapper buffer =
clCreateBuffer(context, CL_MEM_READ_WRITE,
ELEMENT_NUM * sizeof(cl_uint), NULL, &error);
test_error(error, "clCreateBuffer failed");
error = clSetKernelArg(kernel, 0, sizeof(buffer), &buffer);
test_error(error, "clSetKernelArg failed");
// only one thread should be enough to verify if kernel is fully
// functional
size_t workSize = 1;
error = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &workSize, NULL,
0, NULL, NULL);
test_error(error, "clEnqueueNDRangeKernel failed");
std::vector<cl_uint> results(ELEMENT_NUM, 0);
error = clEnqueueReadBuffer(queue, buffer, CL_TRUE, 0,
ELEMENT_NUM * sizeof(cl_uint), &results[0],
0, NULL, NULL);
test_error(error, "clEnqueueReadBuffer failed");
for (size_t i = 0; i < ELEMENT_NUM; ++i)
{
if (results[i] != i)
{
log_error("Kernel %zu returned invalid result. Test: %d, "
"expected: %zu\n",
kernelIdx + 1, results[i], i);
return -1;
}
}
}
clMemWrapper buffer = clCreateBuffer(context, CL_MEM_READ_WRITE, ELEMENT_NUM * sizeof(cl_uint), NULL, &error);
test_error(error, "clCreateBuffer failed");
error = clSetKernelArg(kernel, 0, sizeof(buffer), &buffer);
test_error(error, "clSetKernelArg failed");
//only one thread should be enough to verify if kernel is fully functional
size_t workSize = 1;
error = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &workSize, NULL, 0, NULL, NULL);
test_error(error, "clEnqueueNDRangeKernel failed");
std::vector<cl_uint> results(ELEMENT_NUM, 0);
error = clEnqueueReadBuffer(queue, buffer, CL_TRUE, 0, ELEMENT_NUM * sizeof(cl_uint), &results[0], 0, NULL, NULL);
test_error(error, "clEnqueueReadBuffer failed");
for (size_t i = 0; i < ELEMENT_NUM; ++i) {
if (results[i] != i) {
log_error(
"Kernel %zu returned invalid result. Test: %d, expected: %zu\n",
kernelIdx + 1, results[i], i);
return -1;
}
}
}
return 0;
return 0;
}

451
test_conformance/compiler/test_preprocessor.cpp
View File

@@ -37,85 +37,100 @@ const char *define_kernel_code[] = {
"}\n"};
REGISTER_TEST(preprocessor_define_udef)
{
cl_int error;
clKernelWrapper kernel;
clProgramWrapper program;
clMemWrapper buffer[3];
cl_int *srcData, *resultData;
int i;
MTdata d;
error = create_single_kernel_helper(context, &program, &kernel, 1,
define_kernel_code, "define_test");
if (error) return -1;
buffer[0] = clCreateBuffer(context, CL_MEM_READ_WRITE,
num_elements * sizeof(cl_int), NULL, &error);
test_error(error, "clCreateBuffer failed");
buffer[1] = clCreateBuffer(context, CL_MEM_READ_WRITE,
num_elements * sizeof(cl_int), NULL, &error);
test_error(error, "clCreateBuffer failed");
buffer[2] = clCreateBuffer(context, CL_MEM_READ_WRITE,
num_elements * sizeof(cl_int), NULL, &error);
test_error(error, "clCreateBuffer failed");
srcData = (cl_int *)malloc(sizeof(cl_int) * num_elements);
if (srcData == NULL)
{
log_error("Failed to allocate storage for source data (%d cl_ints).\n",
num_elements);
return -1;
}
d = init_genrand(gRandomSeed);
for (i = 0; i < num_elements; i++)
srcData[i] = (int)get_random_float(-1024, 1024, d);
free_mtdata(d);
d = NULL;
resultData = (cl_int *)malloc(sizeof(cl_int) * num_elements);
if (resultData == NULL)
{
free(srcData);
log_error("Failed to allocate storage for result data (%d cl_ints).\n",
num_elements);
return -1;
}
error = clSetKernelArg(kernel, 0, sizeof(buffer[0]), &buffer[0]);
test_error(error, "clSetKernelArg failed");
error = clSetKernelArg(kernel, 1, sizeof(buffer[1]), &buffer[1]);
test_error(error, "clSetKernelArg failed");
error = clSetKernelArg(kernel, 2, sizeof(buffer[2]), &buffer[2]);
test_error(error, "clSetKernelArg failed");
int test_preprocessor_define_udef(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
error = clEnqueueWriteBuffer(queue, buffer[0], CL_TRUE, 0,
num_elements * sizeof(cl_int), srcData, 0,
NULL, NULL);
test_error(error, "clEnqueueWriteBuffer failed");
cl_int error;
clKernelWrapper kernel;
clProgramWrapper program;
clMemWrapper buffer[3];
cl_int *srcData, *resultData;
int i;
MTdata d;
size_t threads[3] = { (size_t)num_elements, 0, 0 };
error = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, threads, NULL, 0,
NULL, NULL);
test_error(error, "clEnqueueNDRangeKernel failed");
error = create_single_kernel_helper(context, &program, &kernel, 1, define_kernel_code, "define_test");
if (error)
return -1;
error = clEnqueueReadBuffer(queue, buffer[1], CL_TRUE, 0,
num_elements * sizeof(cl_int), resultData, 0,
NULL, NULL);
test_error(error, "clEnqueueReadBuffer failed");
buffer[0] = clCreateBuffer(context, CL_MEM_READ_WRITE, num_elements*sizeof(cl_int), NULL, &error);
test_error( error, "clCreateBuffer failed");
buffer[1] = clCreateBuffer(context, CL_MEM_READ_WRITE, num_elements*sizeof(cl_int), NULL, &error);
test_error( error, "clCreateBuffer failed");
buffer[2] = clCreateBuffer(context, CL_MEM_READ_WRITE, num_elements*sizeof(cl_int), NULL, &error);
test_error( error, "clCreateBuffer failed");
for (i = 0; i < num_elements; i++)
if (resultData[i] != srcData[i] * 2)
{
free(srcData);
free(resultData);
return -1;
}
srcData = (cl_int*)malloc(sizeof(cl_int)*num_elements);
if (srcData == NULL) {
log_error("Failed to allocate storage for source data (%d cl_ints).\n", num_elements);
return -1;
}
error = clEnqueueReadBuffer(queue, buffer[2], CL_TRUE, 0,
num_elements * sizeof(cl_int), resultData, 0,
NULL, NULL);
test_error(error, "clEnqueueReadBuffer failed");
d = init_genrand( gRandomSeed );
for (i=0; i<num_elements; i++)
srcData[i] = (int)get_random_float(-1024, 1024,d);
free_mtdata(d); d = NULL;
for (i = 0; i < num_elements; i++)
if (resultData[i] != srcData[i] * 4)
{
free(srcData);
free(resultData);
return -1;
}
resultData = (cl_int*)malloc(sizeof(cl_int)*num_elements);
if (resultData == NULL) {
free(srcData);
log_error("Failed to allocate storage for result data (%d cl_ints).\n", num_elements);
return -1;
}
error = clSetKernelArg(kernel, 0, sizeof(buffer[0]), &buffer[0]);
test_error(error, "clSetKernelArg failed");
error = clSetKernelArg(kernel, 1, sizeof(buffer[1]), &buffer[1]);
test_error(error, "clSetKernelArg failed");
error = clSetKernelArg(kernel, 2, sizeof(buffer[2]), &buffer[2]);
test_error(error, "clSetKernelArg failed");
error = clEnqueueWriteBuffer(queue, buffer[0], CL_TRUE, 0, num_elements*sizeof(cl_int), srcData, 0, NULL, NULL);
test_error(error, "clEnqueueWriteBuffer failed");
size_t threads[3] = { (size_t)num_elements, 0, 0 };
error = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, threads, NULL, 0, NULL, NULL);
test_error(error, "clEnqueueNDRangeKernel failed");
error = clEnqueueReadBuffer(queue, buffer[1], CL_TRUE, 0, num_elements*sizeof(cl_int), resultData, 0, NULL, NULL);
test_error(error, "clEnqueueReadBuffer failed");
for (i=0; i<num_elements; i++)
if (resultData[i] != srcData[i]*2) {
free(srcData);
free(resultData);
return -1;
}
error = clEnqueueReadBuffer(queue, buffer[2], CL_TRUE, 0, num_elements*sizeof(cl_int), resultData, 0, NULL, NULL);
test_error(error, "clEnqueueReadBuffer failed");
for (i=0; i<num_elements; i++)
if (resultData[i] != srcData[i]*4) {
free(srcData);
free(resultData);
return -1;
}
free(srcData);
free(resultData);
return 0;
free(resultData);
return 0;
}
@@ -133,68 +148,75 @@ const char *include_kernel_code =
"}\n";
int test_preprocessor_include(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
REGISTER_TEST(preprocessor_include)
{
cl_int error;
clKernelWrapper kernel;
clProgramWrapper program;
clMemWrapper buffer[2];
cl_int *resultData;
int i;
cl_int error;
clKernelWrapper kernel;
clProgramWrapper program;
clMemWrapper buffer[2];
cl_int *resultData;
int i;
char include_dir[4096] = { 0 };
char include_kernel[4096] = { 0 };
char include_dir[4096] = {0};
char include_kernel[4096] = {0};
char const *sep = get_dir_sep();
char const *path = get_exe_dir();
char const * sep = get_dir_sep();
char const * path = get_exe_dir();
/* Build with the include directory defined */
sprintf(include_dir, "%s%sincludeTestDirectory%stestIncludeFile.h", path,
sep, sep);
sprintf(include_kernel, include_kernel_code, include_dir);
free((void *)sep);
free((void *)path);
/* Build with the include directory defined */
sprintf(include_dir,"%s%sincludeTestDirectory%stestIncludeFile.h", path, sep, sep);
sprintf(include_kernel, include_kernel_code, include_dir);
free( (void *) sep );
free( (void *) path );
const char *test_kernel[] = { include_kernel, 0 };
error = create_single_kernel_helper(context, &program, &kernel, 1,
test_kernel, "include_test");
if (error) return -1;
const char* test_kernel[] = { include_kernel, 0 };
error = create_single_kernel_helper(context, &program, &kernel, 1, test_kernel, "include_test");
if (error)
return -1;
buffer[0] = clCreateBuffer(context, CL_MEM_READ_WRITE,
num_elements * sizeof(cl_int), NULL, &error);
test_error(error, "clCreateBuffer failed");
buffer[1] = clCreateBuffer(context, CL_MEM_READ_WRITE,
num_elements * sizeof(cl_int), NULL, &error);
test_error(error, "clCreateBuffer failed");
buffer[0] = clCreateBuffer(context, CL_MEM_READ_WRITE, num_elements*sizeof(cl_int), NULL, &error);
test_error( error, "clCreateBuffer failed");
buffer[1] = clCreateBuffer(context, CL_MEM_READ_WRITE, num_elements*sizeof(cl_int), NULL, &error);
test_error( error, "clCreateBuffer failed");
resultData = (cl_int*)malloc(sizeof(cl_int)*num_elements);
if (resultData == NULL) {
log_error("Failed to allocate storage for result data (%d cl_ints).\n", num_elements);
return -1;
}
error = clSetKernelArg(kernel, 0, sizeof(buffer[0]), &buffer[0]);
test_error(error, "clSetKernelArg failed");
error = clSetKernelArg(kernel, 1, sizeof(buffer[1]), &buffer[1]);
test_error(error, "clSetKernelArg failed");
size_t threads[3] = { (size_t)num_elements, 0, 0 };
error = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, threads, NULL, 0, NULL, NULL);
test_error(error, "clEnqueueNDRangeKernel failed");
error = clEnqueueReadBuffer(queue, buffer[1], CL_TRUE, 0, num_elements*sizeof(cl_int), resultData, 0, NULL, NULL);
test_error(error, "clEnqueueReadBuffer failed");
for (i=0; i<num_elements; i++)
if (resultData[i] != 12) {
free(resultData);
return -1;
resultData = (cl_int *)malloc(sizeof(cl_int) * num_elements);
if (resultData == NULL)
{
log_error("Failed to allocate storage for result data (%d cl_ints).\n",
num_elements);
return -1;
}
free(resultData);
return 0;
error = clSetKernelArg(kernel, 0, sizeof(buffer[0]), &buffer[0]);
test_error(error, "clSetKernelArg failed");
error = clSetKernelArg(kernel, 1, sizeof(buffer[1]), &buffer[1]);
test_error(error, "clSetKernelArg failed");
size_t threads[3] = { (size_t)num_elements, 0, 0 };
error = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, threads, NULL, 0,
NULL, NULL);
test_error(error, "clEnqueueNDRangeKernel failed");
error = clEnqueueReadBuffer(queue, buffer[1], CL_TRUE, 0,
num_elements * sizeof(cl_int), resultData, 0,
NULL, NULL);
test_error(error, "clEnqueueReadBuffer failed");
for (i = 0; i < num_elements; i++)
if (resultData[i] != 12)
{
free(resultData);
return -1;
}
free(resultData);
return 0;
}
const char *line_error_kernel_code[] = {
"__kernel void line_error_test(__global int *dstA)\n"
"{\n"
@@ -206,59 +228,79 @@ const char *line_error_kernel_code[] = {
"}\n"};
int test_preprocessor_line_error(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
REGISTER_TEST(preprocessor_line_error)
{
cl_int error, error2;
clKernelWrapper kernel;
clProgramWrapper program;
clMemWrapper buffer[2];
cl_int error, error2;
clKernelWrapper kernel;
clProgramWrapper program;
clMemWrapper buffer[2];
char buildLog[1024 * 128];
char buildLog[ 1024 * 128 ];
log_info("test_preprocessor_line_error may report spurious ERRORS in the "
"conformance log.\n");
log_info("test_preprocessor_line_error may report spurious ERRORS in the conformance log.\n");
/* Create the program object from source */
program = clCreateProgramWithSource(context, 1, line_error_kernel_code,
NULL, &error);
test_error(error, "clCreateProgramWithSource failed");
/* Create the program object from source */
program = clCreateProgramWithSource( context, 1, line_error_kernel_code, NULL, &error );
test_error(error, "clCreateProgramWithSource failed");
/* Compile the program */
error2 = clBuildProgram( program, 0, NULL, NULL, NULL, NULL );
if (error2) {
log_info("Build error detected at clBuildProgram.");
} else {
log_info("Error not reported by clBuildProgram.\n");
}
cl_build_status status;
error = clGetProgramBuildInfo(program, deviceID, CL_PROGRAM_BUILD_STATUS, sizeof(status), &status, NULL);
test_error(error, "clGetProgramBuildInfo failed for CL_PROGRAM_BUILD_STATUS");
if (status != CL_BUILD_ERROR) {
log_error("Build status did not return CL_BUILD_ERROR for a program with #error defined.\n");
return -1;
} else if (status == CL_BUILD_ERROR || error2) {
error2 = clGetProgramBuildInfo( program, deviceID, CL_PROGRAM_BUILD_LOG, sizeof( buildLog ), buildLog, NULL );
test_error( error2, "Unable to get program build log" );
log_info("Build failed as expected with #error in source:\n");
log_info( "Build log is: ------------\n" );
log_info( "%s\n", buildLog );
log_info( "Original source is: ------------\n" );
log_info( "%s", line_error_kernel_code[0] );
log_info( "\n----------\n" );
if (strstr(buildLog, "fictitious/file/name.c")) {
log_info("Found file name from #line param in log output.\n");
} else {
log_info("WARNING: Did not find file name from #line param in log output.\n");
/* Compile the program */
error2 = clBuildProgram(program, 0, NULL, NULL, NULL, NULL);
if (error2)
{
log_info("Build error detected at clBuildProgram.");
}
else
{
log_info("Error not reported by clBuildProgram.\n");
}
if (strstr(buildLog, "124")) {
log_info("Found line number from #line param in log output.\n");
} else {
log_info("WARNING: Did not find line number from #line param in log output.\n");
cl_build_status status;
error = clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_STATUS,
sizeof(status), &status, NULL);
test_error(error,
"clGetProgramBuildInfo failed for CL_PROGRAM_BUILD_STATUS");
if (status != CL_BUILD_ERROR)
{
log_error("Build status did not return CL_BUILD_ERROR for a program "
"with #error defined.\n");
return -1;
}
else if (status == CL_BUILD_ERROR || error2)
{
error2 = clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG,
sizeof(buildLog), buildLog, NULL);
test_error(error2, "Unable to get program build log");
log_info("test_preprocessor_line_error PASSED.\n");
log_info("Build failed as expected with #error in source:\n");
log_info("Build log is: ------------\n");
log_info("%s\n", buildLog);
log_info("Original source is: ------------\n");
log_info("%s", line_error_kernel_code[0]);
log_info("\n----------\n");
if (strstr(buildLog, "fictitious/file/name.c"))
{
log_info("Found file name from #line param in log output.\n");
}
else
{
log_info("WARNING: Did not find file name from #line param in log "
"output.\n");
}
if (strstr(buildLog, "124"))
{
log_info("Found line number from #line param in log output.\n");
}
else
{
log_info("WARNING: Did not find line number from #line param in "
"log output.\n");
}
log_info("test_preprocessor_line_error PASSED.\n");
return 0;
}
@@ -281,7 +323,6 @@ int test_preprocessor_line_error(cl_device_id deviceID, cl_context context, cl_c
}
const char *pragma_kernel_code[] = {
"__kernel void pragma_test(__global int *dstA)\n"
"{\n"
@@ -294,47 +335,51 @@ const char *pragma_kernel_code[] = {
"}\n"};
int test_preprocessor_pragma(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
REGISTER_TEST(preprocessor_pragma)
{
cl_int error;
clKernelWrapper kernel;
clProgramWrapper program;
clMemWrapper buffer[2];
cl_int *resultData;
int i;
cl_int error;
clKernelWrapper kernel;
clProgramWrapper program;
clMemWrapper buffer[2];
cl_int *resultData;
int i;
error = create_single_kernel_helper(context, &program, &kernel, 1,
pragma_kernel_code, "pragma_test");
if (error) return -1;
error = create_single_kernel_helper(context, &program, &kernel, 1, pragma_kernel_code, "pragma_test");
if (error)
return -1;
buffer[0] = clCreateBuffer(context, CL_MEM_READ_WRITE,
num_elements * sizeof(cl_int), NULL, &error);
test_error(error, "clCreateBuffer failed");
buffer[0] = clCreateBuffer(context, CL_MEM_READ_WRITE, num_elements*sizeof(cl_int), NULL, &error);
test_error( error, "clCreateBuffer failed");
error = clSetKernelArg(kernel, 0, sizeof(buffer[0]), &buffer[0]);
test_error(error, "clSetKernelArg failed");
error = clSetKernelArg(kernel, 0, sizeof(buffer[0]), &buffer[0]);
test_error(error, "clSetKernelArg failed");
size_t threads[3] = { (size_t)num_elements, 0, 0 };
error = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, threads, NULL, 0,
NULL, NULL);
test_error(error, "clEnqueueNDRangeKernel failed");
size_t threads[3] = { (size_t)num_elements, 0, 0 };
error = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, threads, NULL, 0, NULL, NULL);
test_error(error, "clEnqueueNDRangeKernel failed");
resultData = (cl_int*)malloc(sizeof(cl_int)*num_elements);
if (resultData == NULL) {
log_error("Failed to allocate storage for result data (%d cl_ints).\n", num_elements);
return -1;
}
error = clEnqueueReadBuffer(queue, buffer[0], CL_TRUE, 0, num_elements*sizeof(cl_int), resultData, 0, NULL, NULL);
test_error(error, "clEnqueueReadBuffer failed");
for (i=0; i<num_elements; i++)
if (resultData[i] != i) {
free(resultData);
return -1;
resultData = (cl_int *)malloc(sizeof(cl_int) * num_elements);
if (resultData == NULL)
{
log_error("Failed to allocate storage for result data (%d cl_ints).\n",
num_elements);
return -1;
}
free(resultData);
return 0;
error = clEnqueueReadBuffer(queue, buffer[0], CL_TRUE, 0,
num_elements * sizeof(cl_int), resultData, 0,
NULL, NULL);
test_error(error, "clEnqueueReadBuffer failed");
for (i = 0; i < num_elements; i++)
if (resultData[i] != i)
{
free(resultData);
return -1;
}
free(resultData);
return 0;
}

40
test_conformance/compiler/test_unload_platform_compiler.cpp
View File

@@ -378,7 +378,7 @@ static void unload_platform_compiler(const cl_platform_id platform)
}
/* Test calling the function with a valid platform */
int test_unload_valid(cl_device_id device, cl_context, cl_command_queue, int)
REGISTER_TEST(unload_valid)
{
const cl_platform_id platform = device_platform(device);
const long int err = clUnloadPlatformCompiler(platform);
@@ -393,7 +393,8 @@ int test_unload_valid(cl_device_id device, cl_context, cl_command_queue, int)
}
/* Test calling the function with invalid platform */
int test_unload_invalid(cl_device_id, cl_context, cl_command_queue, int)
/* Disabling temporarily, see GitHub #977
REGISTER_TEST(unload_invalid)
{
const long int err = clUnloadPlatformCompiler(nullptr);
@@ -405,10 +406,10 @@ int test_unload_invalid(cl_device_id, cl_context, cl_command_queue, int)
return 0;
}
*/
/* Test calling the function multiple times in a row */
int test_unload_repeated(cl_device_id device, cl_context context,
cl_command_queue, int)
REGISTER_TEST(unload_repeated)
{
check_compiler_available(device);
@@ -438,8 +439,7 @@ int test_unload_repeated(cl_device_id device, cl_context context,
}
/* Test calling the function between compilation and linking of programs */
int test_unload_compile_unload_link(cl_device_id device, cl_context context,
cl_command_queue, int)
REGISTER_TEST(unload_compile_unload_link)
{
check_compiler_available(device);
@@ -469,9 +469,7 @@ int test_unload_compile_unload_link(cl_device_id device, cl_context context,
}
/* Test calling the function between program build and kernel creation */
int test_unload_build_unload_create_kernel(cl_device_id device,
cl_context context, cl_command_queue,
int)
REGISTER_TEST(unload_build_unload_create_kernel)
{
check_compiler_available(device);
@@ -501,8 +499,7 @@ int test_unload_build_unload_create_kernel(cl_device_id device,
/* Test linking together two programs that were built with a call to the unload
* function in between */
int test_unload_link_different(cl_device_id device, cl_context context,
cl_command_queue, int)
REGISTER_TEST(unload_link_different)
{
check_compiler_available(device);
@@ -587,7 +584,7 @@ int test_unload_link_different(cl_device_id device, cl_context context,
return 1;
}
const clCommandQueueWrapper queue =
const clCommandQueueWrapper test_queue =
clCreateCommandQueue(context, device, 0, &err);
if (CL_SUCCESS != err)
{
@@ -616,8 +613,8 @@ int test_unload_link_different(cl_device_id device, cl_context context,
}
static const size_t work_size = 1;
err = clEnqueueNDRangeKernel(queue, kernel, 1, nullptr, &work_size, nullptr,
0, nullptr, nullptr);
err = clEnqueueNDRangeKernel(test_queue, kernel, 1, nullptr, &work_size,
nullptr, 0, nullptr, nullptr);
if (CL_SUCCESS != err)
{
log_error("Test failure: clEnqueueNDRangeKernel() == %ld\n",
@@ -625,8 +622,8 @@ int test_unload_link_different(cl_device_id device, cl_context context,
return 1;
}
err = clEnqueueReadBuffer(queue, buffer, CL_BLOCKING, 0, sizeof(cl_uint),
&value, 0, nullptr, nullptr);
err = clEnqueueReadBuffer(test_queue, buffer, CL_BLOCKING, 0,
sizeof(cl_uint), &value, 0, nullptr, nullptr);
if (CL_SUCCESS != err)
{
log_error("Test failure: clEnqueueReadBuffer() == %ld\n",
@@ -634,7 +631,7 @@ int test_unload_link_different(cl_device_id device, cl_context context,
return 1;
}
err = clFinish(queue);
err = clFinish(test_queue);
if (CL_SUCCESS != err) throw unload_test_failure("clFinish()", err);
if (42 != value)
@@ -649,8 +646,7 @@ int test_unload_link_different(cl_device_id device, cl_context context,
/* Test calling the function in a thread while others threads are building
* programs */
int test_unload_build_threaded(cl_device_id device, cl_context context,
cl_command_queue, int)
REGISTER_TEST(unload_build_threaded)
{
using clock = std::chrono::steady_clock;
@@ -740,8 +736,7 @@ int test_unload_build_threaded(cl_device_id device, cl_context context,
}
/* Test grabbing program build information after calling the unload function */
int test_unload_build_info(cl_device_id device, cl_context context,
cl_command_queue, int)
REGISTER_TEST(unload_build_info)
{
check_compiler_available(device);
@@ -903,8 +898,7 @@ int test_unload_build_info(cl_device_id device, cl_context context,
/* Test calling the unload function between program building and fetching the
* program binaries */
int test_unload_program_binaries(cl_device_id device, cl_context context,
cl_command_queue, int)
REGISTER_TEST(unload_program_binaries)
{
check_compiler_available(device);

15
test_conformance/conversions/test_conversions.cpp
View File

@@ -92,8 +92,7 @@ char appName[64] = "ctest";
int gMultithread = 1;
int test_conversions(cl_device_id device, cl_context context,
cl_command_queue queue, int num_elements)
REGISTER_TEST(conversions)
{
if (argCount)
{
@@ -108,13 +107,6 @@ int test_conversions(cl_device_id device, cl_context context,
}
test_definition test_list[] = {
ADD_TEST(conversions),
};
const int test_num = ARRAY_SIZE(test_list);
int main(int argc, const char **argv)
{
int error;
@@ -148,8 +140,9 @@ int main(int argc, const char **argv)
gMTdata = init_genrand(gRandomSeed);
const char *arg[] = { argv[0] };
int ret =
runTestHarnessWithCheck(1, arg, test_num, test_list, true, 0, InitCL);
int ret = runTestHarnessWithCheck(
1, arg, test_registry::getInstance().num_tests(),
test_registry::getInstance().definitions(), true, 0, InitCL);
free_mtdata(gMTdata);
if (gQueue)

67
test_conformance/generic_address_space/advanced_tests.cpp
View File

@@ -293,7 +293,8 @@ private:
const ExtraKernelArgMemType _extraKernelArgMemType;
};
int test_library_function(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
REGISTER_TEST(library_function)
{
const std::string LIBRARY_FUNCTION = common::CONFORMANCE_VERIFY_FENCE +
NL
NL "bool helperFunction(float *floatp, float val) {"
@@ -340,10 +341,11 @@ __kernel void testKernel(__global uint *results) {
CAdvancedTest test(LIBRARY_FUNCTION, KERNEL_FUNCTION);
return test.Execute(deviceID, context, queue, num_elements);
return test.Execute(device, context, queue, num_elements);
}
int test_generic_variable_volatile(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
REGISTER_TEST(generic_variable_volatile)
{
std::vector<std::string> KERNEL_FUNCTIONS;
KERNEL_FUNCTIONS.push_back(common::CONFORMANCE_VERIFY_FENCE +
@@ -420,10 +422,11 @@ int test_generic_variable_volatile(cl_device_id deviceID, cl_context context, cl
CAdvancedTest test(KERNEL_FUNCTIONS);
return test.Execute(deviceID, context, queue, num_elements);
return test.Execute(device, context, queue, num_elements);
}
int test_generic_variable_const(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
REGISTER_TEST(generic_variable_const)
{
std::vector<std::string> KERNEL_FUNCTIONS;
KERNEL_FUNCTIONS.push_back(common::CONFORMANCE_VERIFY_FENCE +
@@ -474,10 +477,11 @@ int test_generic_variable_const(cl_device_id deviceID, cl_context context, cl_co
CAdvancedTest test(KERNEL_FUNCTIONS);
return test.Execute(deviceID, context, queue, num_elements);
return test.Execute(device, context, queue, num_elements);
}
int test_generic_variable_gentype(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
REGISTER_TEST(generic_variable_gentype)
{
const std::string KERNEL_FUNCTION_TEMPLATE = common::CONFORMANCE_VERIFY_FENCE +
NL
NL "%s"
@@ -531,7 +535,8 @@ int test_generic_variable_gentype(cl_device_id deviceID, cl_context context, cl_
const std::string cl_khr_fp64_pragma = "#pragma OPENCL EXTENSION cl_khr_fp64 : enable";
// Add double floating types if they are supported
if (is_extension_available(deviceID, "cl_khr_fp64")) {
if (is_extension_available(device, "cl_khr_fp64"))
{
for (size_t j = 0; j < sizeof(vector_sizes) / sizeof(vector_sizes[0]); j++) {
for (size_t k = 0; k < sizeof(address_spaces) / sizeof(address_spaces[0]); k++) {
char temp_kernel[1024];
@@ -551,7 +556,8 @@ int test_generic_variable_gentype(cl_device_id deviceID, cl_context context, cl_
const std::string cl_khr_fp16_pragma = "#pragma OPENCL EXTENSION cl_khr_fp16 : enable";
// Add half floating types if they are supported
if (is_extension_available(deviceID, "cl_khr_fp16")) {
if (is_extension_available(device, "cl_khr_fp16"))
{
for (size_t j = 0; j < sizeof(vector_sizes) / sizeof(vector_sizes[0]); j++) {
for (size_t k = 0; k < sizeof(address_spaces) / sizeof(address_spaces[0]); k++) {
char temp_kernel[1024];
@@ -586,7 +592,7 @@ int test_generic_variable_gentype(cl_device_id deviceID, cl_context context, cl_
CAdvancedTest test(KERNEL_FUNCTIONS);
return test.Execute(deviceID, context, queue, num_elements);
return test.Execute(device, context, queue, num_elements);
}
void create_math_kernels(std::vector<std::string>& KERNEL_FUNCTIONS) {
@@ -919,19 +925,21 @@ void create_vstore_kernels(std::vector<std::string>& KERNEL_FUNCTIONS, cl_device
}
}
int test_builtin_functions(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
REGISTER_TEST(builtin_functions)
{
std::vector<std::string> KERNEL_FUNCTIONS;
create_math_kernels(KERNEL_FUNCTIONS);
create_vload_kernels(KERNEL_FUNCTIONS, deviceID);
create_vstore_kernels(KERNEL_FUNCTIONS, deviceID);
create_vload_kernels(KERNEL_FUNCTIONS, device);
create_vstore_kernels(KERNEL_FUNCTIONS, device);
CAdvancedTest test(KERNEL_FUNCTIONS);
return test.Execute(deviceID, context, queue, num_elements);
return test.Execute(device, context, queue, num_elements);
}
int test_generic_advanced_casting(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
REGISTER_TEST(generic_advanced_casting)
{
std::vector<std::string> KERNEL_FUNCTIONS;
KERNEL_FUNCTIONS.push_back(
@@ -980,18 +988,20 @@ int test_generic_advanced_casting(cl_device_id deviceID, cl_context context, cl_
CAdvancedTest test(KERNEL_FUNCTIONS);
return test.Execute(deviceID, context, queue, num_elements);
return test.Execute(device, context, queue, num_elements);
}
int test_generic_ptr_to_host_mem_svm(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
REGISTER_TEST(generic_ptr_to_host_mem_svm)
{
cl_int result = CL_SUCCESS;
/* Test SVM capabilities and select matching tests */
cl_device_svm_capabilities caps;
auto version = get_device_cl_version(deviceID);
auto version = get_device_cl_version(device);
auto expected_min_version = Version(2, 0);
cl_int error = clGetDeviceInfo(deviceID, CL_DEVICE_SVM_CAPABILITIES, sizeof(caps), &caps, NULL);
cl_int error = clGetDeviceInfo(device, CL_DEVICE_SVM_CAPABILITIES,
sizeof(caps), &caps, NULL);
test_error(error, "clGetDeviceInfo(CL_DEVICE_SVM_CAPABILITIES) failed");
if ((version < expected_min_version)
@@ -1000,35 +1010,40 @@ int test_generic_ptr_to_host_mem_svm(cl_device_id deviceID, cl_context context,
if (caps & CL_DEVICE_SVM_COARSE_GRAIN_BUFFER) {
CAdvancedTest test_global_svm_ptr(common::GLOBAL_KERNEL_FUNCTION, ARG_TYPE_COARSE_GRAINED_SVM);
result |= test_global_svm_ptr.Execute(deviceID, context, queue, num_elements);
result |=
test_global_svm_ptr.Execute(device, context, queue, num_elements);
}
if (caps & CL_DEVICE_SVM_FINE_GRAIN_BUFFER) {
CAdvancedTest test_global_svm_ptr(common::GLOBAL_KERNEL_FUNCTION, ARG_TYPE_FINE_GRAINED_BUFFER_SVM);
result |= test_global_svm_ptr.Execute(deviceID, context, queue, num_elements);
result |=
test_global_svm_ptr.Execute(device, context, queue, num_elements);
}
if (caps & CL_DEVICE_SVM_FINE_GRAIN_SYSTEM) {
CAdvancedTest test_global_svm_ptr(common::GLOBAL_KERNEL_FUNCTION, ARG_TYPE_FINE_GRAINED_SYSTEM_SVM);
result |= test_global_svm_ptr.Execute(deviceID, context, queue, num_elements);
result |=
test_global_svm_ptr.Execute(device, context, queue, num_elements);
}
if (caps & CL_DEVICE_SVM_ATOMICS) {
CAdvancedTest test_global_svm_ptr(common::GLOBAL_KERNEL_FUNCTION, ARG_TYPE_ATOMICS_SVM);
result |= test_global_svm_ptr.Execute(deviceID, context, queue, num_elements);
result |=
test_global_svm_ptr.Execute(device, context, queue, num_elements);
}
return result;
}
int test_generic_ptr_to_host_mem(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
REGISTER_TEST(generic_ptr_to_host_mem)
{
cl_int result = CL_SUCCESS;
CAdvancedTest test_global_ptr(common::GLOBAL_KERNEL_FUNCTION, ARG_TYPE_HOST_PTR);
result |= test_global_ptr.Execute(deviceID, context, queue, num_elements);
result |= test_global_ptr.Execute(device, context, queue, num_elements);
CAdvancedTest test_local_ptr(common::LOCAL_KERNEL_FUNCTION, ARG_TYPE_HOST_LOCAL);
result |= test_local_ptr.Execute(deviceID, context, queue, num_elements / 64);
result |= test_local_ptr.Execute(device, context, queue, num_elements / 64);
return result;
}

14
test_conformance/generic_address_space/atomic_tests.cpp
View File

@@ -91,10 +91,9 @@ kernel void testKernel(global atomic_int* globalPtr, local atomic_int* localPtr)
)OpenCLC";
}
int test_generic_atomics_invariant(cl_device_id deviceID, cl_context context,
cl_command_queue queue, int)
REGISTER_TEST(generic_atomics_invariant)
{
const auto version = get_device_cl_version(deviceID);
const auto version = get_device_cl_version(device);
if (version < Version(2, 0)) return TEST_SKIPPED_ITSELF;
@@ -108,7 +107,7 @@ int test_generic_atomics_invariant(cl_device_id deviceID, cl_context context,
size_t wgSize, retSize;
// Attempt to find the simd unit size for the device.
err = clGetKernelWorkGroupInfo(kernel, deviceID,
err = clGetKernelWorkGroupInfo(kernel, device,
CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE,
sizeof(wgSize), &wgSize, &retSize);
test_error(err, "clGetKernelWorkGroupInfo failed");
@@ -154,10 +153,9 @@ int test_generic_atomics_invariant(cl_device_id deviceID, cl_context context,
return CL_SUCCESS;
}
int test_generic_atomics_variant(cl_device_id deviceID, cl_context context,
cl_command_queue queue, int)
REGISTER_TEST(generic_atomics_variant)
{
const auto version = get_device_cl_version(deviceID);
const auto version = get_device_cl_version(device);
if (version < Version(2, 0)) return TEST_SKIPPED_ITSELF;
@@ -171,7 +169,7 @@ int test_generic_atomics_variant(cl_device_id deviceID, cl_context context,
size_t wgSize, retSize;
// Attempt to find the simd unit size for the device.
err = clGetKernelWorkGroupInfo(kernel, deviceID,
err = clGetKernelWorkGroupInfo(kernel, device,
CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE,
sizeof(wgSize), &wgSize, &retSize);
test_error(err, "clGetKernelWorkGroupInfo failed");

60
test_conformance/generic_address_space/basic_tests.cpp
View File

@@ -97,7 +97,8 @@ private:
const std::vector<std::string> _kernels;
};
int test_function_get_fence(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
REGISTER_TEST(function_get_fence)
{
const std::string KERNEL_FUNCTION = common::CONFORMANCE_VERIFY_FENCE +
NL
NL "__global int gint = 1;"
@@ -142,10 +143,11 @@ int test_function_get_fence(cl_device_id deviceID, cl_context context, cl_comman
CBasicTest test(KERNEL_FUNCTION);
return test.Execute(deviceID, context, queue, num_elements);
return test.Execute(device, context, queue, num_elements);
}
int test_function_to_address_space(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
REGISTER_TEST(function_to_address_space)
{
const std::string KERNEL_FUNCTION =
NL
NL "__global int gint = 1;"
@@ -190,10 +192,11 @@ int test_function_to_address_space(cl_device_id deviceID, cl_context context, cl
CBasicTest test(KERNEL_FUNCTION);
return test.Execute(deviceID, context, queue, num_elements);
return test.Execute(device, context, queue, num_elements);
}
int test_variable_get_fence(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
REGISTER_TEST(variable_get_fence)
{
const std::string KERNEL_FUNCTION = common::CONFORMANCE_VERIFY_FENCE +
NL
NL "__global int gint = 1;"
@@ -223,10 +226,11 @@ int test_variable_get_fence(cl_device_id deviceID, cl_context context, cl_comman
CBasicTest test(KERNEL_FUNCTION);
return test.Execute(deviceID, context, queue, num_elements);
return test.Execute(device, context, queue, num_elements);
}
int test_variable_to_address_space(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
REGISTER_TEST(variable_to_address_space)
{
const std::string KERNEL_FUNCTION =
NL
NL "__global int gint = 1;"
@@ -256,10 +260,11 @@ int test_variable_to_address_space(cl_device_id deviceID, cl_context context, cl
CBasicTest test(KERNEL_FUNCTION);
return test.Execute(deviceID, context, queue, num_elements);
return test.Execute(device, context, queue, num_elements);
}
int test_casting(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
REGISTER_TEST(casting)
{
std::vector<std::string> KERNEL_FUNCTIONS;
// pointers to global, local or private are implicitly convertible to generic
@@ -345,10 +350,11 @@ int test_casting(cl_device_id deviceID, cl_context context, cl_command_queue que
CBasicTest test(KERNEL_FUNCTIONS);
return test.Execute(deviceID, context, queue, num_elements);
return test.Execute(device, context, queue, num_elements);
}
int test_conditional_casting(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
REGISTER_TEST(conditional_casting)
{
const std::string KERNEL_FUNCTION = common::CONFORMANCE_VERIFY_FENCE +
NL
NL "__global int gint = 1;"
@@ -376,10 +382,11 @@ int test_conditional_casting(cl_device_id deviceID, cl_context context, cl_comma
CBasicTest test(KERNEL_FUNCTION);
return test.Execute(deviceID, context, queue, num_elements);
return test.Execute(device, context, queue, num_elements);
}
int test_chain_casting(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
REGISTER_TEST(chain_casting)
{
const std::string KERNEL_FUNCTION = common::CONFORMANCE_VERIFY_FENCE +
NL
NL "__global int gint = 1;"
@@ -407,10 +414,11 @@ int test_chain_casting(cl_device_id deviceID, cl_context context, cl_command_que
NL;
CBasicTest test(KERNEL_FUNCTION);
return test.Execute(deviceID, context, queue, num_elements);
return test.Execute(device, context, queue, num_elements);
}
int test_ternary_operator_casting(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
REGISTER_TEST(ternary_operator_casting)
{
const std::string KERNEL_FUNCTION = common::CONFORMANCE_VERIFY_FENCE +
NL
NL "__global int gint = 1;"
@@ -435,10 +443,11 @@ int test_ternary_operator_casting(cl_device_id deviceID, cl_context context, cl_
CBasicTest test(KERNEL_FUNCTION);
return test.Execute(deviceID, context, queue, num_elements);
return test.Execute(device, context, queue, num_elements);
}
int test_language_struct(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
REGISTER_TEST(language_struct)
{
std::vector<std::string> KERNEL_FUNCTIONS;
// implicit private struct
@@ -588,10 +597,11 @@ int test_language_struct(cl_device_id deviceID, cl_context context, cl_command_q
CBasicTest test(KERNEL_FUNCTIONS);
return test.Execute(deviceID, context, queue, num_elements);
return test.Execute(device, context, queue, num_elements);
}
int test_language_union(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
REGISTER_TEST(language_union)
{
std::vector<std::string> KERNEL_FUNCTIONS;
KERNEL_FUNCTIONS.push_back(common::CONFORMANCE_VERIFY_FENCE +
@@ -726,10 +736,11 @@ int test_language_union(cl_device_id deviceID, cl_context context, cl_command_qu
CBasicTest test(KERNEL_FUNCTIONS);
return test.Execute(deviceID, context, queue, num_elements);
return test.Execute(device, context, queue, num_elements);
}
int test_multiple_calls_same_function(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
REGISTER_TEST(multiple_calls_same_function)
{
const std::string KERNEL_FUNCTION =
NL
NL "int shift2(const int *ptr, int arg) {"
@@ -759,10 +770,11 @@ int test_multiple_calls_same_function(cl_device_id deviceID, cl_context context,
CBasicTest test(KERNEL_FUNCTION);
return test.Execute(deviceID, context, queue, num_elements);
return test.Execute(device, context, queue, num_elements);
}
int test_compare_pointers(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
REGISTER_TEST(compare_pointers)
{
std::vector<std::string> KERNEL_FUNCTIONS;
KERNEL_FUNCTIONS.push_back(
@@ -884,5 +896,5 @@ int test_compare_pointers(cl_device_id deviceID, cl_context context, cl_command_
CBasicTest test(KERNEL_FUNCTIONS);
return test.Execute(deviceID, context, queue, num_elements);
return test.Execute(device, context, queue, num_elements);
}

66
test_conformance/generic_address_space/main.cpp
View File

@@ -18,66 +18,8 @@
#include <iostream>
// basic tests
extern int test_function_get_fence(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_function_to_address_space(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_variable_get_fence(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_variable_to_address_space(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_casting(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_conditional_casting(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_chain_casting(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_ternary_operator_casting(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_language_struct(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_language_union(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_multiple_calls_same_function(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_compare_pointers(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
// advanced tests
extern int test_library_function(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_generic_variable_volatile(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_generic_variable_const(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_generic_variable_gentype(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_builtin_functions(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_generic_advanced_casting(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_generic_ptr_to_host_mem(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_generic_ptr_to_host_mem_svm(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_max_number_of_params(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
// atomic tests
int test_generic_atomics_invariant(cl_device_id deviceID, cl_context context,
cl_command_queue queue, int num_elements);
int test_generic_atomics_variant(cl_device_id deviceID, cl_context context,
cl_command_queue queue, int num_elements);
test_definition test_list[] = {
// basic tests
ADD_TEST(function_get_fence),
ADD_TEST(function_to_address_space),
ADD_TEST(variable_get_fence),
ADD_TEST(variable_to_address_space),
ADD_TEST(casting),
ADD_TEST(conditional_casting),
ADD_TEST(chain_casting),
ADD_TEST(ternary_operator_casting),
ADD_TEST(language_struct),
ADD_TEST(language_union),
ADD_TEST(multiple_calls_same_function),
ADD_TEST(compare_pointers),
// advanced tests
ADD_TEST(library_function),
ADD_TEST(generic_variable_volatile),
ADD_TEST(generic_variable_const),
ADD_TEST(generic_variable_gentype),
ADD_TEST(builtin_functions),
ADD_TEST(generic_advanced_casting),
ADD_TEST(generic_ptr_to_host_mem),
ADD_TEST(generic_ptr_to_host_mem_svm),
ADD_TEST(max_number_of_params),
// atomic tests
ADD_TEST(generic_atomics_invariant),
ADD_TEST(generic_atomics_variant),
};
const int test_num = ARRAY_SIZE( test_list );
test_status InitCL(cl_device_id device) {
test_status InitCL(cl_device_id device)
{
auto version = get_device_cl_version(device);
auto expected_min_version = Version(2, 0);
@@ -134,5 +76,7 @@ test_status InitCL(cl_device_id device) {
int main(int argc, const char *argv[])
{
return runTestHarnessWithCheck(argc, argv, test_num, test_list, false, false, InitCL);
return runTestHarnessWithCheck(
argc, argv, test_registry::getInstance().num_tests(),
test_registry::getInstance().definitions(), false, false, InitCL);
}

13
test_conformance/generic_address_space/stress_tests.cpp
View File

@@ -99,15 +99,20 @@ private:
const std::vector<std::string> _kernels;
};
int test_max_number_of_params(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
REGISTER_TEST(max_number_of_params)
{
cl_int error;
size_t deviceMaxParameterSize;
error = clGetDeviceInfo(deviceID, CL_DEVICE_MAX_PARAMETER_SIZE, sizeof(deviceMaxParameterSize), &deviceMaxParameterSize, NULL);
error = clGetDeviceInfo(device, CL_DEVICE_MAX_PARAMETER_SIZE,
sizeof(deviceMaxParameterSize),
&deviceMaxParameterSize, NULL);
test_error(error, "clGetDeviceInfo failed");
size_t deviceAddressBits;
error = clGetDeviceInfo(deviceID, CL_DEVICE_ADDRESS_BITS, sizeof(deviceAddressBits), &deviceAddressBits, NULL);
error =
clGetDeviceInfo(device, CL_DEVICE_ADDRESS_BITS,
sizeof(deviceAddressBits), &deviceAddressBits, NULL);
test_error(error, "clGetDeviceInfo failed");
size_t maxParams = deviceMaxParameterSize / (deviceAddressBits / 8);
@@ -174,5 +179,5 @@ int test_max_number_of_params(cl_device_id deviceID, cl_context context, cl_comm
CStressTest test(KERNEL_FUNCTION);
return test.Execute(deviceID, context, queue, num_elements);
return test.Execute(device, context, queue, num_elements);
}

17
test_conformance/geometrics/main.cpp
View File

@@ -17,27 +17,14 @@
#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( geom_cross ),
ADD_TEST( geom_dot ),
ADD_TEST( geom_distance ),
ADD_TEST( geom_fast_distance ),
ADD_TEST( geom_length ),
ADD_TEST( geom_fast_length ),
ADD_TEST( geom_normalize ),
ADD_TEST( geom_fast_normalize ),
};
const int test_num = ARRAY_SIZE( test_list );
int main(int argc, const char *argv[])
{
return runTestHarness(argc, argv, test_num, test_list, false, 0);
return runTestHarness(argc, argv, test_registry::getInstance().num_tests(),
test_registry::getInstance().definitions(), false, 0);
}

35
test_conformance/geometrics/procs.h
View File

@@ -1,35 +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/typeWrappers.h"
extern int create_program_and_kernel(const char *source, const char *kernel_name, cl_program *program_ret, cl_kernel *kernel_ret);
extern int test_geom_cross(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_geom_dot(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_geom_distance(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_geom_fast_distance(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_geom_length(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_geom_fast_length(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_geom_normalize(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_geom_fast_normalize(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_geom_cross_double(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements, MTdata d);
extern int test_geom_dot_double(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements, MTdata d);
extern int test_geom_distance_double(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements, MTdata d);
extern int test_geom_length_double(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements, MTdata d);
extern int test_geom_normalize_double(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements, MTdata d);

18
test_conformance/geometrics/testBase.h
View File

@@ -16,14 +16,30 @@
#ifndef _testBase_h
#define _testBase_h
#include <CL/cl.h>
#include "harness/compat.h"
#include "harness/mt19937.h"
#include <stdio.h>
#include <string.h>
#include <sys/types.h>
#include <sys/stat.h>
#include "procs.h"
extern int test_geom_cross_double(cl_device_id deviceID, cl_context context,
cl_command_queue queue, int num_elements,
MTdata d);
extern int test_geom_dot_double(cl_device_id deviceID, cl_context context,
cl_command_queue queue, int num_elements,
MTdata d);
extern int test_geom_distance_double(cl_device_id deviceID, cl_context context,
cl_command_queue queue, int num_elements,
MTdata d);
extern int test_geom_length_double(cl_device_id deviceID, cl_context context,
cl_command_queue queue, int num_elements,
MTdata d);
extern int test_geom_normalize_double(cl_device_id deviceID, cl_context context,
cl_command_queue queue, int num_elements,
MTdata d);
#endif // _testBase_h

49
test_conformance/geometrics/test_geometrics.cpp
View File

@@ -148,15 +148,13 @@ void cross_product( const float *vecA, const float *vecB, float *outVector, floa
}
int test_geom_cross(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements )
REGISTER_TEST(geom_cross)
{
int vecsize;
RandomSeed seed(gRandomSeed);
/* Get the default rounding mode */
cl_device_fp_config defaultRoundingMode = get_default_rounding_mode(deviceID);
cl_device_fp_config defaultRoundingMode = get_default_rounding_mode(device);
if( 0 == defaultRoundingMode )
return -1;
@@ -269,12 +267,16 @@ int test_geom_cross(cl_device_id deviceID, cl_context context, cl_command_queue
}
} // for(vecsize=...
if(!is_extension_available(deviceID, "cl_khr_fp64")) {
if (!is_extension_available(device, "cl_khr_fp64"))
{
log_info("Extension cl_khr_fp64 not supported; skipping double tests.\n");
return 0;
} else {
}
else
{
log_info("Testing doubles...\n");
return test_geom_cross_double( deviceID, context, queue, num_elements, seed);
return test_geom_cross_double(device, context, queue, num_elements,
seed);
}
}
@@ -497,7 +499,7 @@ double verifyDot( float *srcA, float *srcB, size_t vecSize )
return total;
}
int test_geom_dot(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
REGISTER_TEST(geom_dot)
{
size_t sizes[] = { 1, 2, 3, 4, 0 };
unsigned int size;
@@ -516,14 +518,14 @@ int test_geom_dot(cl_device_id deviceID, cl_context context, cl_command_queue qu
if (retVal)
return retVal;
if(!is_extension_available(deviceID, "cl_khr_fp64"))
if (!is_extension_available(device, "cl_khr_fp64"))
{
log_info("Extension cl_khr_fp64 not supported; skipping double tests.\n");
return 0;
}
log_info("Testing doubles...\n");
return test_geom_dot_double( deviceID, context, queue, num_elements, seed);
return test_geom_dot_double(device, context, queue, num_elements, seed);
}
double verifyFastDistance( float *srcA, float *srcB, size_t vecSize )
@@ -542,7 +544,7 @@ double verifyFastDistance( float *srcA, float *srcB, size_t vecSize )
return sqrt( total );
}
int test_geom_fast_distance(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
REGISTER_TEST(geom_fast_distance)
{
size_t sizes[] = { 1, 2, 3, 4, 0 };
unsigned int size;
@@ -589,7 +591,7 @@ double verifyDistance( float *srcA, float *srcB, size_t vecSize )
return sqrt( total );
}
int test_geom_distance(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
REGISTER_TEST(geom_distance)
{
size_t sizes[] = { 1, 2, 3, 4, 0 };
unsigned int size;
@@ -616,13 +618,14 @@ int test_geom_distance(cl_device_id deviceID, cl_context context, cl_command_que
if (retVal)
return retVal;
if(!is_extension_available(deviceID, "cl_khr_fp64"))
if (!is_extension_available(device, "cl_khr_fp64"))
{
log_info("Extension cl_khr_fp64 not supported; skipping double tests.\n");
return 0;
} else {
log_info("Testing doubles...\n");
return test_geom_distance_double( deviceID, context, queue, num_elements, seed);
return test_geom_distance_double(device, context, queue, num_elements,
seed);
}
}
@@ -754,7 +757,7 @@ double verifyLength( float *srcA, size_t vecSize )
return sqrt( total );
}
int test_geom_length(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
REGISTER_TEST(geom_length)
{
size_t sizes[] = { 1, 2, 3, 4, 0 };
unsigned int size;
@@ -781,7 +784,7 @@ int test_geom_length(cl_device_id deviceID, cl_context context, cl_command_queue
if (retVal)
return retVal;
if(!is_extension_available(deviceID, "cl_khr_fp64"))
if (!is_extension_available(device, "cl_khr_fp64"))
{
log_info("Extension cl_khr_fp64 not supported; skipping double tests.\n");
return 0;
@@ -789,7 +792,8 @@ int test_geom_length(cl_device_id deviceID, cl_context context, cl_command_queue
else
{
log_info("Testing doubles...\n");
return test_geom_length_double( deviceID, context, queue, num_elements, seed);
return test_geom_length_double(device, context, queue, num_elements,
seed);
}
}
@@ -809,7 +813,7 @@ double verifyFastLength( float *srcA, size_t vecSize )
return sqrt( total );
}
int test_geom_fast_length(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
REGISTER_TEST(geom_fast_length)
{
size_t sizes[] = { 1, 2, 3, 4, 0 };
unsigned int size;
@@ -1059,7 +1063,7 @@ void verifyNormalize( float *srcA, float *dst, size_t vecSize )
dst[i] = (float)( (double)srcA[i] / value );
}
int test_geom_normalize(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
REGISTER_TEST(geom_normalize)
{
size_t sizes[] = { 1, 2, 3, 4, 0 };
unsigned int size;
@@ -1084,18 +1088,19 @@ int test_geom_normalize(cl_device_id deviceID, cl_context context, cl_command_qu
if (retVal)
return retVal;
if(!is_extension_available(deviceID, "cl_khr_fp64"))
if (!is_extension_available(device, "cl_khr_fp64"))
{
log_info("Extension cl_khr_fp64 not supported; skipping double tests.\n");
return 0;
} else {
log_info("Testing doubles...\n");
return test_geom_normalize_double( deviceID, context, queue, num_elements, seed);
return test_geom_normalize_double(device, context, queue, num_elements,
seed);
}
}
int test_geom_fast_normalize(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
REGISTER_TEST(geom_fast_normalize)
{
size_t sizes[] = { 1, 2, 3, 4, 0 };
unsigned int size;

5
test_conformance/math_brute_force/function_list.cpp
View File

@@ -291,13 +291,12 @@ const Func functionList[] = {
ENTRY(erfc, 16.0f, 16.0f, 4.0f, 4.0f, FTZ_OFF, unaryF),
ENTRY(erf, 16.0f, 16.0f, 4.0f, 4.0f, FTZ_OFF, unaryF),
// relaxed error is overwritten in unary.c as it is 3+floor(fabs(2*x))
// floor(fabs(2*x)) is added to the relaxed error in unary.c
ENTRY_EXT(exp, 3.0f, 4.0f, 2.0f, 3.0f, 3.0f, FTZ_OFF, unaryF, 4.0f),
// relaxed error is overwritten in unary.c as it is 3+floor(fabs(2*x))
// floor(fabs(2*x)) is added to the relaxed error in unary.c
ENTRY_EXT(exp2, 3.0f, 4.0f, 2.0f, 3.0f, 3.0f, FTZ_OFF, unaryF, 4.0f),
// relaxed error is overwritten in unary.c as it is 3+floor(fabs(2*x)) in derived mode;
// in non-derived mode it uses the ulp error for half_exp10.
ENTRY_EXT(exp10, 3.0f, 4.0f, 2.0f, 3.0f, 8192.0f, FTZ_OFF, unaryF, 8192.0f),

165
test_conformance/math_brute_force/main.cpp
View File

@@ -263,61 +263,112 @@ static int doTest(const char *name)
return error;
}
#define DO_TEST(name) \
REGISTER_TEST_VERSION(name, Version(1, 0)) { return doTest(#name); }
#define TEST_LAMBDA(name) \
[](cl_device_id, cl_context, cl_command_queue, int) { \
return doTest(#name); \
}
// Redefine ADD_TEST to use TEST_LAMBDA.
#undef ADD_TEST
#define ADD_TEST(name) \
{ \
TEST_LAMBDA(name), #name, Version(1, 0) \
}
static test_definition test_list[] = {
ADD_TEST(acos), ADD_TEST(acosh), ADD_TEST(acospi),
ADD_TEST(asin), ADD_TEST(asinh), ADD_TEST(asinpi),
ADD_TEST(atan), ADD_TEST(atanh), ADD_TEST(atanpi),
ADD_TEST(atan2), ADD_TEST(atan2pi), ADD_TEST(cbrt),
ADD_TEST(ceil), ADD_TEST(copysign), ADD_TEST(cos),
ADD_TEST(cosh), ADD_TEST(cospi), ADD_TEST(exp),
ADD_TEST(exp2), ADD_TEST(exp10), ADD_TEST(expm1),
ADD_TEST(fabs), ADD_TEST(fdim), ADD_TEST(floor),
ADD_TEST(fma), ADD_TEST(fmax), ADD_TEST(fmin),
ADD_TEST(fmod), ADD_TEST(fract), ADD_TEST(frexp),
ADD_TEST(hypot), ADD_TEST(ilogb), ADD_TEST(isequal),
ADD_TEST(isfinite), ADD_TEST(isgreater), ADD_TEST(isgreaterequal),
ADD_TEST(isinf), ADD_TEST(isless), ADD_TEST(islessequal),
ADD_TEST(islessgreater), ADD_TEST(isnan), ADD_TEST(isnormal),
ADD_TEST(isnotequal), ADD_TEST(isordered), ADD_TEST(isunordered),
ADD_TEST(ldexp), ADD_TEST(lgamma), ADD_TEST(lgamma_r),
ADD_TEST(log), ADD_TEST(log2), ADD_TEST(log10),
ADD_TEST(log1p), ADD_TEST(logb), ADD_TEST(mad),
ADD_TEST(maxmag), ADD_TEST(minmag), ADD_TEST(modf),
ADD_TEST(nan), ADD_TEST(nextafter), ADD_TEST(pow),
ADD_TEST(pown), ADD_TEST(powr), ADD_TEST(remainder),
ADD_TEST(remquo), ADD_TEST(rint), ADD_TEST(rootn),
ADD_TEST(round), ADD_TEST(rsqrt), ADD_TEST(signbit),
ADD_TEST(sin), ADD_TEST(sincos), ADD_TEST(sinh),
ADD_TEST(sinpi), ADD_TEST(sqrt), ADD_TEST(sqrt_cr),
ADD_TEST(tan), ADD_TEST(tanh), ADD_TEST(tanpi),
ADD_TEST(trunc), ADD_TEST(half_cos), ADD_TEST(half_divide),
ADD_TEST(half_exp), ADD_TEST(half_exp2), ADD_TEST(half_exp10),
ADD_TEST(half_log), ADD_TEST(half_log2), ADD_TEST(half_log10),
ADD_TEST(half_powr), ADD_TEST(half_recip), ADD_TEST(half_rsqrt),
ADD_TEST(half_sin), ADD_TEST(half_sqrt), ADD_TEST(half_tan),
ADD_TEST(add), ADD_TEST(subtract), ADD_TEST(reciprocal),
ADD_TEST(divide), ADD_TEST(divide_cr), ADD_TEST(multiply),
ADD_TEST(assignment), ADD_TEST(not ), ADD_TEST(erf),
ADD_TEST(erfc),
};
#undef ADD_TEST
#undef TEST_LAMBDA
static const int test_num = ARRAY_SIZE(test_list);
DO_TEST(acos)
DO_TEST(acosh)
DO_TEST(acospi)
DO_TEST(asin)
DO_TEST(asinh)
DO_TEST(asinpi)
DO_TEST(atan)
DO_TEST(atanh)
DO_TEST(atanpi)
DO_TEST(atan2)
DO_TEST(atan2pi)
DO_TEST(cbrt)
DO_TEST(ceil)
DO_TEST(copysign)
DO_TEST(cos)
DO_TEST(cosh)
DO_TEST(cospi)
DO_TEST(exp)
DO_TEST(exp2)
DO_TEST(exp10)
DO_TEST(expm1)
DO_TEST(fabs)
DO_TEST(fdim)
DO_TEST(floor)
DO_TEST(fma)
DO_TEST(fmax)
DO_TEST(fmin)
DO_TEST(fmod)
DO_TEST(fract)
DO_TEST(frexp)
DO_TEST(hypot)
DO_TEST(ilogb)
DO_TEST(isequal)
DO_TEST(isfinite)
DO_TEST(isgreater)
DO_TEST(isgreaterequal)
DO_TEST(isinf)
DO_TEST(isless)
DO_TEST(islessequal)
DO_TEST(islessgreater)
DO_TEST(isnan)
DO_TEST(isnormal)
DO_TEST(isnotequal)
DO_TEST(isordered)
DO_TEST(isunordered)
DO_TEST(ldexp)
DO_TEST(lgamma)
DO_TEST(lgamma_r)
DO_TEST(log)
DO_TEST(log2)
DO_TEST(log10)
DO_TEST(log1p)
DO_TEST(logb)
DO_TEST(mad)
DO_TEST(maxmag)
DO_TEST(minmag)
DO_TEST(modf)
DO_TEST(nan)
DO_TEST(nextafter)
DO_TEST(pow)
DO_TEST(pown)
DO_TEST(powr)
DO_TEST(remainder)
DO_TEST(remquo)
DO_TEST(rint)
DO_TEST(rootn)
DO_TEST(round)
DO_TEST(rsqrt)
DO_TEST(signbit)
DO_TEST(sin)
DO_TEST(sincos)
DO_TEST(sinh)
DO_TEST(sinpi)
DO_TEST(sqrt)
DO_TEST(sqrt_cr)
DO_TEST(tan)
DO_TEST(tanh)
DO_TEST(tanpi)
DO_TEST(trunc)
DO_TEST(half_cos)
DO_TEST(half_divide)
DO_TEST(half_exp)
DO_TEST(half_exp2)
DO_TEST(half_exp10)
DO_TEST(half_log)
DO_TEST(half_log2)
DO_TEST(half_log10)
DO_TEST(half_powr)
DO_TEST(half_recip)
DO_TEST(half_rsqrt)
DO_TEST(half_sin)
DO_TEST(half_sqrt)
DO_TEST(half_tan)
DO_TEST(add)
DO_TEST(subtract)
DO_TEST(reciprocal)
DO_TEST(divide)
DO_TEST(divide_cr)
DO_TEST(multiply)
DO_TEST(assignment)
DO_TEST(not )
DO_TEST(erf)
DO_TEST(erfc)
#pragma mark -
@@ -355,8 +406,10 @@ int main(int argc, const char *argv[])
FPU_mode_type oldMode;
DisableFTZ(&oldMode);
int ret = runTestHarnessWithCheck(gTestNames.size(), gTestNames.data(),
test_num, test_list, true, 0, InitCL);
int ret = runTestHarnessWithCheck(
gTestNames.size(), gTestNames.data(),
test_registry::getInstance().num_tests(),
test_registry::getInstance().definitions(), true, 0, InitCL);
RestoreFPState(&oldMode);

8
test_conformance/math_brute_force/unary_float.cpp
View File

@@ -300,13 +300,7 @@ cl_int Test(cl_uint job_id, cl_uint thread_id, void *data)
if (strcmp(fname, "exp") == 0 || strcmp(fname, "exp2") == 0)
{
// For full profile, ULP depends on input value.
// For embedded profile, ULP comes from functionList.
if (!gIsEmbedded)
{
ulps = 3.0f + floor(fabs(2 * s[j]));
}
ulps += floor(fabs(2 * s[j]));
fail = !(fabsf(err) <= ulps);
}
if (strcmp(fname, "tan") == 0)

30
test_conformance/multiple_device_context/main.cpp
View File

@@ -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
@@ -13,34 +13,10 @@
// 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"
#include "harness/mt19937.h"
#if !defined(_WIN32)
#include <unistd.h>
#endif
test_definition test_list[] = {
ADD_TEST( context_multiple_contexts_same_device ),
ADD_TEST( context_two_contexts_same_device ),
ADD_TEST( context_three_contexts_same_device ),
ADD_TEST( context_four_contexts_same_device ),
ADD_TEST( two_devices ),
ADD_TEST( max_devices ),
ADD_TEST( hundred_queues ),
};
const int test_num = ARRAY_SIZE( test_list );
int main(int argc, const char *argv[])
{
return runTestHarness(argc, argv, test_num, test_list, true, 0);
return runTestHarness(argc, argv, test_registry::getInstance().num_tests(),
test_registry::getInstance().definitions(), true, 0);
}

31
test_conformance/multiple_device_context/procs.h
View File

@@ -1,31 +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/typeWrappers.h"
#include "harness/mt19937.h"
extern int test_context_multiple_contexts_same_device(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_context_two_contexts_same_device(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_context_three_contexts_same_device(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_context_four_contexts_same_device(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_two_devices(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_max_devices(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
extern int test_hundred_queues(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements);

16
test_conformance/multiple_device_context/testBase.h
View File

@@ -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
@@ -16,16 +16,8 @@
#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"
#include "harness/conversions.h"
#include "harness/testHarness.h"
#include "harness/kernelHelpers.h"
#endif // _testBase_h

22
test_conformance/multiple_device_context/test_multiple_contexts.cpp
View File

@@ -414,8 +414,9 @@ cl_int UseTestItem( const TestItem *item, cl_int *err )
}
int test_context_multiple_contexts_same_device(cl_device_id deviceID, size_t maxCount, size_t minCount )
static int test_context_multiple_contexts_same_device(cl_device_id deviceID,
size_t maxCount,
size_t minCount)
{
size_t i, j;
cl_int err = CL_SUCCESS;
@@ -525,23 +526,22 @@ exit:
// sane limit, currently 200), attempting to use each along the way. We keep track of how many we could make before
// a failure occurred. We then free everything and attempt to go do it again a few times. If you are able to make
// that many contexts 5 times over, then you pass.
int test_context_multiple_contexts_same_device(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
REGISTER_TEST(context_multiple_contexts_same_device)
{
return test_context_multiple_contexts_same_device(deviceID, 200, 1);
return test_context_multiple_contexts_same_device(device, 200, 1);
}
int test_context_two_contexts_same_device(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
REGISTER_TEST(context_two_contexts_same_device)
{
return test_context_multiple_contexts_same_device( deviceID, 2, 2 );
return test_context_multiple_contexts_same_device(device, 2, 2);
}
int test_context_three_contexts_same_device(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
REGISTER_TEST(context_three_contexts_same_device)
{
return test_context_multiple_contexts_same_device( deviceID, 3, 3 );
return test_context_multiple_contexts_same_device(device, 3, 3);
}
int test_context_four_contexts_same_device(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
REGISTER_TEST(context_four_contexts_same_device)
{
return test_context_multiple_contexts_same_device( deviceID, 4, 4 );
return test_context_multiple_contexts_same_device(device, 4, 4);
}

12
test_conformance/multiple_device_context/test_multiple_devices.cpp
View File

@@ -35,7 +35,8 @@ const char *test_kernels[] = { "__kernel void kernelA(__global uint *dst)\n"
#define MAX_DEVICES 32
#define MAX_QUEUES 1000
int test_device_set(size_t deviceCount, size_t queueCount, cl_device_id *devices, int num_elements)
static int test_device_set(size_t deviceCount, size_t queueCount,
cl_device_id *devices, int num_elements)
{
int error;
clContextWrapper context;
@@ -187,7 +188,7 @@ int test_device_set(size_t deviceCount, size_t queueCount, cl_device_id *devices
return 0;
}
int test_two_devices(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
REGISTER_TEST(two_devices)
{
cl_platform_id platform;
cl_device_id devices[2];
@@ -215,7 +216,7 @@ int test_two_devices(cl_device_id deviceID, cl_context context, cl_command_queue
return test_device_set( 2, 2, devices, num_elements );
}
int test_max_devices(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
REGISTER_TEST(max_devices)
{
cl_platform_id platform;
cl_device_id devices[MAX_DEVICES];
@@ -235,8 +236,7 @@ int test_max_devices(cl_device_id deviceID, cl_context context, cl_command_queue
return test_device_set( deviceCount, deviceCount, devices, num_elements );
}
int test_hundred_queues(cl_device_id device, cl_context contextIgnore, cl_command_queue queueIgnore, int num_elements)
REGISTER_TEST(hundred_queues)
{
return test_device_set( 1, 100, &device, num_elements );
return test_device_set(1, 100, &device, num_elements);
}