diff --git a/test_conformance/integer_ops/CMakeLists.txt b/test_conformance/integer_ops/CMakeLists.txt index 5344eabc..7bc991f8 100644 --- a/test_conformance/integer_ops/CMakeLists.txt +++ b/test_conformance/integer_ops/CMakeLists.txt @@ -12,6 +12,9 @@ set(${MODULE_NAME}_SOURCES verification_and_generation_functions.cpp test_popcount.cpp test_integer_dot_product.cpp + test_extended_bit_ops_extract.cpp + test_extended_bit_ops_insert.cpp + test_extended_bit_ops_reverse.cpp ) include(../CMakeCommon.txt) diff --git a/test_conformance/integer_ops/main.cpp b/test_conformance/integer_ops/main.cpp index e57cffd9..59840de7 100644 --- a/test_conformance/integer_ops/main.cpp +++ b/test_conformance/integer_ops/main.cpp @@ -1,6 +1,6 @@ // -// Copyright (c) 2017 The Khronos Group Inc. -// +// Copyright (c) 2017-2022 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 @@ -59,6 +59,10 @@ test_definition test_list[] = { ADD_TEST(integer_mul24), ADD_TEST(integer_mad24), + ADD_TEST(extended_bit_ops_extract), + ADD_TEST(extended_bit_ops_insert), + ADD_TEST(extended_bit_ops_reverse), + ADD_TEST(long_math), ADD_TEST(long_logic), ADD_TEST(long_shift), diff --git a/test_conformance/integer_ops/procs.h b/test_conformance/integer_ops/procs.h index 82311fb9..31f6ae5f 100644 --- a/test_conformance/integer_ops/procs.h +++ b/test_conformance/integer_ops/procs.h @@ -1,6 +1,6 @@ // -// Copyright (c) 2017 The Khronos Group Inc. -// +// Copyright (c) 2017-2022 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 @@ -61,6 +61,18 @@ extern int test_integer_sub_sat(cl_device_id deviceID, cl_context context, cl_co extern int test_integer_mul24(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); extern int test_integer_mad24(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); +extern int test_extended_bit_ops_extract(cl_device_id device_id, + cl_context context, + cl_command_queue commands, + int num_elements); +extern int test_extended_bit_ops_insert(cl_device_id device_id, + cl_context context, + cl_command_queue commands, + int num_elements); +extern int test_extended_bit_ops_reverse(cl_device_id device_id, + cl_context context, + cl_command_queue commands, + int num_elements); extern int test_long_math(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); extern int test_long_logic(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); diff --git a/test_conformance/integer_ops/test_extended_bit_ops_extract.cpp b/test_conformance/integer_ops/test_extended_bit_ops_extract.cpp new file mode 100644 index 00000000..9b4e0950 --- /dev/null +++ b/test_conformance/integer_ops/test_extended_bit_ops_extract.cpp @@ -0,0 +1,287 @@ +// +// Copyright (c) 2022 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 +#include +#include +#include +#include + +#include "procs.h" +#include "harness/integer_ops_test_info.h" +#include "harness/testHarness.h" + +template +static typename std::make_unsigned::type +arithmetic_shift_right(T tx, cl_uint count) +{ + typedef typename std::make_unsigned::type unsigned_t; + unsigned_t x = static_cast(tx); + + // To implement an arithmetic shift right: + // - If the sign bit is not set, shift as usual. + // - Otherwise, flip all of the bits, shift, then flip back. + unsigned_t s = -(x >> (sizeof(x) * 8 - 1)); + unsigned_t result = (s ^ x) >> count ^ s; + + return result; +} + +template +static typename std::make_unsigned::type +cpu_bit_extract_signed(T tbase, cl_uint offset, cl_uint count) +{ + typedef typename std::make_signed::type unsigned_t; + + assert(offset <= sizeof(T) * 8); + assert(count <= sizeof(T) * 8); + assert(offset + count <= sizeof(T) * 8); + + unsigned_t base = static_cast(tbase); + unsigned_t result; + + if (count == 0) + { + result = 0; + } + else + { + result = base << (sizeof(T) * 8 - offset - count); + result = arithmetic_shift_right(result, sizeof(T) * 8 - count); + } + + return result; +} + +template +static typename std::make_unsigned::type +cpu_bit_extract_unsigned(T tbase, cl_uint offset, cl_uint count) +{ + typedef typename std::make_unsigned::type unsigned_t; + + assert(offset <= sizeof(T) * 8); + assert(count <= sizeof(T) * 8); + assert(offset + count <= sizeof(T) * 8); + + unsigned_t base = static_cast(tbase); + unsigned_t result; + + if (count == 0) + { + result = 0; + } + else + { + result = base << (sizeof(T) * 8 - offset - count); + result = result >> (sizeof(T) * 8 - count); + } + + return result; +} + +template +static void +calculate_reference(std::vector::type>& sref, + std::vector::type>& uref, + const std::vector& base) +{ + sref.resize(base.size()); + uref.resize(base.size()); + for (size_t i = 0; i < base.size(); i++) + { + cl_uint offset = (i / N) / (sizeof(T) * 8 + 1); + cl_uint count = (i / N) % (sizeof(T) * 8 + 1); + if (offset + count > sizeof(T) * 8) + { + count = (sizeof(T) * 8) - offset; + } + sref[i] = cpu_bit_extract_signed(base[i], offset, count); + uref[i] = cpu_bit_extract_unsigned(base[i], offset, count); + } +} + +static constexpr const char* kernel_source = R"CLC( +__kernel void test_bitfield_extract(__global SIGNED_TYPE* sdst, __global UNSIGNED_TYPE* udst, __global TYPE* base) +{ + int index = get_global_id(0); + uint offset = index / (sizeof(BASETYPE) * 8 + 1); + uint count = index % (sizeof(BASETYPE) * 8 + 1); + if (offset + count > sizeof(BASETYPE) * 8) { + count = (sizeof(BASETYPE) * 8) - offset; + } + sdst[index] = bitfield_extract_signed(base[index], offset, count); + udst[index] = bitfield_extract_unsigned(base[index], offset, count); +} +)CLC"; + +static constexpr const char* kernel_source_vec3 = R"CLC( +__kernel void test_bitfield_extract(__global SIGNED_BASETYPE* sdst, __global UNSIGNED_BASETYPE* udst, __global BASETYPE* base) +{ + int index = get_global_id(0); + uint offset = index / (sizeof(BASETYPE) * 8 + 1); + uint count = index % (sizeof(BASETYPE) * 8 + 1); + if (offset + count > sizeof(BASETYPE) * 8) { + count = (sizeof(BASETYPE) * 8) - offset; + } + TYPE b = vload3(index, base); + SIGNED_TYPE s = bitfield_extract_signed(b, offset, count); + UNSIGNED_TYPE u = bitfield_extract_unsigned(b, offset, count); + vstore3(s, index, sdst); + vstore3(u, index, udst); +} +)CLC"; + +template +static int test_vectype(cl_device_id device, cl_context context, + cl_command_queue queue) +{ + // Because converting from an unsigned type to a signed type is + // implementation-defined if the most significant bit is set until C++ 20, + // compute all reference results using unsigned types. + typedef typename std::make_unsigned::type unsigned_t; + + cl_int error = CL_SUCCESS; + + clProgramWrapper program; + clKernelWrapper kernel; + + std::string buildOptions; + buildOptions += " -DTYPE="; + buildOptions += + TestInfo::deviceTypeName + ((N > 1) ? std::to_string(N) : ""); + buildOptions += " -DSIGNED_TYPE="; + buildOptions += + TestInfo::deviceTypeNameSigned + ((N > 1) ? std::to_string(N) : ""); + buildOptions += " -DUNSIGNED_TYPE="; + buildOptions += TestInfo::deviceTypeNameUnsigned + + ((N > 1) ? std::to_string(N) : ""); + buildOptions += " -DBASETYPE="; + buildOptions += TestInfo::deviceTypeName; + buildOptions += " -DSIGNED_BASETYPE="; + buildOptions += TestInfo::deviceTypeNameSigned; + buildOptions += " -DUNSIGNED_BASETYPE="; + buildOptions += TestInfo::deviceTypeNameUnsigned; + + const size_t ELEMENTS_TO_TEST = (sizeof(T) * 8 + 1) * (sizeof(T) * 8 + 1); + + std::vector base(ELEMENTS_TO_TEST * N); + fill_vector_with_random_data(base); + + std::vector sreference; + std::vector ureference; + calculate_reference(sreference, ureference, base); + + const char* source = (N == 3) ? kernel_source_vec3 : kernel_source; + error = create_single_kernel_helper(context, &program, &kernel, 1, &source, + "test_bitfield_extract", + buildOptions.c_str()); + test_error(error, "Unable to create test_bitfield_insert kernel"); + + clMemWrapper sdst = + clCreateBuffer(context, 0, sreference.size() * sizeof(T), NULL, &error); + test_error(error, "Unable to create signed output buffer"); + + clMemWrapper udst = + clCreateBuffer(context, 0, ureference.size() * sizeof(T), NULL, &error); + test_error(error, "Unable to create unsigned output buffer"); + + clMemWrapper src_base = + clCreateBuffer(context, CL_MEM_COPY_HOST_PTR, base.size() * sizeof(T), + base.data(), &error); + test_error(error, "Unable to create base buffer"); + + error = clSetKernelArg(kernel, 0, sizeof(sdst), &sdst); + test_error(error, "Unable to set signed output buffer kernel arg"); + + error = clSetKernelArg(kernel, 1, sizeof(udst), &udst); + test_error(error, "Unable to set unsigned output buffer kernel arg"); + + error = clSetKernelArg(kernel, 2, sizeof(src_base), &src_base); + test_error(error, "Unable to set base buffer kernel arg"); + + size_t global_work_size[] = { sreference.size() / N }; + error = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, global_work_size, + NULL, 0, NULL, NULL); + test_error(error, "Unable to enqueue test kernel"); + + error = clFinish(queue); + test_error(error, "clFinish failed after test kernel"); + + std::vector sresults(sreference.size(), 99); + error = clEnqueueReadBuffer(queue, sdst, CL_TRUE, 0, + sresults.size() * sizeof(T), sresults.data(), 0, + NULL, NULL); + test_error(error, "Unable to read signed data after test kernel"); + + if (sresults != sreference) + { + log_error("Signed result buffer did not match reference buffer!\n"); + return TEST_FAIL; + } + + std::vector uresults(ureference.size(), 99); + error = clEnqueueReadBuffer(queue, udst, CL_TRUE, 0, + uresults.size() * sizeof(T), uresults.data(), 0, + NULL, NULL); + test_error(error, "Unable to read unsigned data after test kernel"); + + if (uresults != ureference) + { + log_error("Unsigned result buffer did not match reference buffer!\n"); + return TEST_FAIL; + } + + return TEST_PASS; +} + +template +static int test_type(cl_device_id device, cl_context context, + cl_command_queue queue) +{ + log_info(" testing type %s\n", TestInfo::deviceTypeName); + + return test_vectype(device, context, queue) + | test_vectype(device, context, queue) + | test_vectype(device, context, queue) + | test_vectype(device, context, queue) + | test_vectype(device, context, queue) + | test_vectype(device, context, queue); +} + +int test_extended_bit_ops_extract(cl_device_id device, cl_context context, + cl_command_queue queue, int num_elements) +{ + if (is_extension_available(device, "cl_khr_extended_bit_ops")) + { + int result = TEST_PASS; + + result |= test_type(device, context, queue); + result |= test_type(device, context, queue); + result |= test_type(device, context, queue); + result |= test_type(device, context, queue); + result |= test_type(device, context, queue); + result |= test_type(device, context, queue); + if (gHasLong) + { + result |= test_type(device, context, queue); + result |= test_type(device, context, queue); + } + return result; + } + + log_info("cl_khr_extended_bit_ops is not supported\n"); + return TEST_SKIPPED_ITSELF; +} diff --git a/test_conformance/integer_ops/test_extended_bit_ops_insert.cpp b/test_conformance/integer_ops/test_extended_bit_ops_insert.cpp new file mode 100644 index 00000000..e6d8522c --- /dev/null +++ b/test_conformance/integer_ops/test_extended_bit_ops_insert.cpp @@ -0,0 +1,214 @@ +// +// Copyright (c) 2022 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 +#include +#include +#include +#include + +#include "procs.h" +#include "harness/integer_ops_test_info.h" +#include "harness/testHarness.h" + +template +static typename std::make_unsigned::type +cpu_bit_insert(T tbase, T tinsert, cl_uint offset, cl_uint count) +{ + assert(offset <= sizeof(T) * 8); + assert(count <= sizeof(T) * 8); + assert(offset + count <= sizeof(T) * 8); + + cl_ulong base = static_cast(tbase); + cl_ulong insert = static_cast(tinsert); + + cl_ulong mask = (count < 64) ? ((1ULL << count) - 1) << offset : ~0ULL; + cl_ulong result = ((insert << offset) & mask) | (base & ~mask); + + return static_cast::type>(result); +} + +template +static void +calculate_reference(std::vector::type>& ref, + const std::vector& base, const std::vector& insert) +{ + ref.resize(base.size()); + for (size_t i = 0; i < base.size(); i++) + { + cl_uint offset = (i / N) / (sizeof(T) * 8 + 1); + cl_uint count = (i / N) % (sizeof(T) * 8 + 1); + if (offset + count > sizeof(T) * 8) + { + count = (sizeof(T) * 8) - offset; + } + ref[i] = cpu_bit_insert(base[i], insert[i], offset, count); + } +} + +static constexpr const char* kernel_source = R"CLC( +__kernel void test_bitfield_insert(__global TYPE* dst, __global TYPE* base, __global TYPE* insert) +{ + int index = get_global_id(0); + uint offset = index / (sizeof(BASETYPE) * 8 + 1); + uint count = index % (sizeof(BASETYPE) * 8 + 1); + if (offset + count > sizeof(BASETYPE) * 8) { + count = (sizeof(BASETYPE) * 8) - offset; + } + dst[index] = bitfield_insert(base[index], insert[index], offset, count); +} +)CLC"; + +static constexpr const char* kernel_source_vec3 = R"CLC( +__kernel void test_bitfield_insert(__global BASETYPE* dst, __global BASETYPE* base, __global BASETYPE* insert) +{ + int index = get_global_id(0); + uint offset = index / (sizeof(BASETYPE) * 8 + 1); + uint count = index % (sizeof(BASETYPE) * 8 + 1); + if (offset + count > sizeof(BASETYPE) * 8) { + count = (sizeof(BASETYPE) * 8) - offset; + } + TYPE b = vload3(index, base); + TYPE i = vload3(index, insert); + TYPE d = bitfield_insert(b, i, offset, count); + vstore3(d, index, dst); +} +)CLC"; + +template +static int test_vectype(cl_device_id device, cl_context context, + cl_command_queue queue) +{ + // Because converting from an unsigned type to a signed type is + // implementation-defined if the most significant bit is set until C++ 20, + // compute all reference results using unsigned types. + typedef typename std::make_unsigned::type unsigned_t; + + cl_int error = CL_SUCCESS; + + clProgramWrapper program; + clKernelWrapper kernel; + + std::string buildOptions{ "-DTYPE=" }; + buildOptions += TestInfo::deviceTypeName; + if (N > 1) + { + buildOptions += std::to_string(N); + } + buildOptions += " -DBASETYPE="; + buildOptions += TestInfo::deviceTypeName; + + const size_t ELEMENTS_TO_TEST = (sizeof(T) * 8 + 1) * (sizeof(T) * 8 + 1); + + std::vector base(ELEMENTS_TO_TEST * N); + std::fill(base.begin(), base.end(), static_cast(0xA5A5A5A5A5A5A5A5ULL)); + + std::vector insert(ELEMENTS_TO_TEST * N); + fill_vector_with_random_data(insert); + + std::vector reference; + calculate_reference(reference, base, insert); + + const char* source = (N == 3) ? kernel_source_vec3 : kernel_source; + error = create_single_kernel_helper(context, &program, &kernel, 1, &source, + "test_bitfield_insert", + buildOptions.c_str()); + test_error(error, "Unable to create test_bitfield_insert kernel"); + + clMemWrapper dst = + clCreateBuffer(context, 0, reference.size() * sizeof(T), NULL, &error); + test_error(error, "Unable to create output buffer"); + + clMemWrapper src_base = + clCreateBuffer(context, CL_MEM_COPY_HOST_PTR, base.size() * sizeof(T), + base.data(), &error); + test_error(error, "Unable to create base buffer"); + + clMemWrapper src_insert = + clCreateBuffer(context, CL_MEM_COPY_HOST_PTR, insert.size() * sizeof(T), + insert.data(), &error); + test_error(error, "Unable to create insert buffer"); + + error = clSetKernelArg(kernel, 0, sizeof(dst), &dst); + test_error(error, "Unable to set output buffer kernel arg"); + + error = clSetKernelArg(kernel, 1, sizeof(src_base), &src_base); + test_error(error, "Unable to set base buffer kernel arg"); + + error = clSetKernelArg(kernel, 2, sizeof(src_insert), &src_insert); + test_error(error, "Unable to set insert buffer kernel arg"); + + size_t global_work_size[] = { reference.size() / N }; + error = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, global_work_size, + NULL, 0, NULL, NULL); + test_error(error, "Unable to enqueue test kernel"); + + error = clFinish(queue); + test_error(error, "clFinish failed after test kernel"); + + std::vector results(reference.size(), 99); + error = + clEnqueueReadBuffer(queue, dst, CL_TRUE, 0, results.size() * sizeof(T), + results.data(), 0, NULL, NULL); + test_error(error, "Unable to read data after test kernel"); + + if (results != reference) + { + log_error("Result buffer did not match reference buffer!\n"); + return TEST_FAIL; + } + + return TEST_PASS; +} + +template +static int test_type(cl_device_id device, cl_context context, + cl_command_queue queue) +{ + log_info(" testing type %s\n", TestInfo::deviceTypeName); + + return test_vectype(device, context, queue) + | test_vectype(device, context, queue) + | test_vectype(device, context, queue) + | test_vectype(device, context, queue) + | test_vectype(device, context, queue) + | test_vectype(device, context, queue); +} + +int test_extended_bit_ops_insert(cl_device_id device, cl_context context, + cl_command_queue queue, int num_elements) +{ + if (is_extension_available(device, "cl_khr_extended_bit_ops")) + { + int result = TEST_PASS; + + result |= test_type(device, context, queue); + result |= test_type(device, context, queue); + result |= test_type(device, context, queue); + result |= test_type(device, context, queue); + result |= test_type(device, context, queue); + result |= test_type(device, context, queue); + if (gHasLong) + { + result |= test_type(device, context, queue); + result |= test_type(device, context, queue); + } + return result; + } + + log_info("cl_khr_extended_bit_ops is not supported\n"); + return TEST_SKIPPED_ITSELF; +} diff --git a/test_conformance/integer_ops/test_extended_bit_ops_reverse.cpp b/test_conformance/integer_ops/test_extended_bit_ops_reverse.cpp new file mode 100644 index 00000000..136f9d1d --- /dev/null +++ b/test_conformance/integer_ops/test_extended_bit_ops_reverse.cpp @@ -0,0 +1,177 @@ +// +// Copyright (c) 2022 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 +#include +#include +#include + +#include "procs.h" +#include "harness/integer_ops_test_info.h" +#include "harness/testHarness.h" + +template static T cpu_bit_reverse(T base) +{ + T result = 0; + + const size_t count = sizeof(T) * 8; + for (size_t i = 0; i < count; i++) + { + if (base & ((T)1 << i)) + { + result |= ((T)1 << (count - i - 1)); + } + } + return result; +} + +template +static void calculate_reference(std::vector& ref, const std::vector& base) +{ + ref.resize(base.size()); + for (size_t i = 0; i < base.size(); i++) + { + ref[i] = cpu_bit_reverse(base[i]); + } +} + +static constexpr const char* kernel_source = R"CLC( +__kernel void test_bit_reverse(__global TYPE* dst, __global TYPE* base) +{ + int index = get_global_id(0); + dst[index] = bit_reverse(base[index]); +} +)CLC"; + +static constexpr const char* kernel_source_vec3 = R"CLC( +__kernel void test_bit_reverse(__global BASETYPE* dst, __global BASETYPE* base) +{ + int index = get_global_id(0); + TYPE s = vload3(index, base); + TYPE d = bit_reverse(s); + vstore3(d, index, dst); +} +)CLC"; + +template +static int test_vectype(cl_device_id device, cl_context context, + cl_command_queue queue) +{ + cl_int error = CL_SUCCESS; + + clProgramWrapper program; + clKernelWrapper kernel; + + std::string buildOptions{ "-DTYPE=" }; + buildOptions += TestInfo::deviceTypeName; + if (N > 1) + { + buildOptions += std::to_string(N); + } + buildOptions += " -DBASETYPE="; + buildOptions += TestInfo::deviceTypeName; + + const size_t ELEMENTS_TO_TEST = 65536; + std::vector base(ELEMENTS_TO_TEST * N); + fill_vector_with_random_data(base); + + std::vector reference; + calculate_reference(reference, base); + + const char* source = (N == 3) ? kernel_source_vec3 : kernel_source; + error = + create_single_kernel_helper(context, &program, &kernel, 1, &source, + "test_bit_reverse", buildOptions.c_str()); + test_error(error, "Unable to create test_bit_reverse kernel"); + + clMemWrapper src; + clMemWrapper dst; + + dst = + clCreateBuffer(context, 0, reference.size() * sizeof(T), NULL, &error); + test_error(error, "Unable to create output buffer"); + + src = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR, base.size() * sizeof(T), + base.data(), &error); + test_error(error, "Unable to create base buffer"); + + error = clSetKernelArg(kernel, 0, sizeof(dst), &dst); + test_error(error, "Unable to set output buffer kernel arg"); + + error = clSetKernelArg(kernel, 1, sizeof(src), &src); + test_error(error, "Unable to set base buffer kernel arg"); + + size_t global_work_size[] = { reference.size() / N }; + error = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, global_work_size, + NULL, 0, NULL, NULL); + test_error(error, "Unable to enqueue test kernel"); + + error = clFinish(queue); + test_error(error, "clFinish failed after test kernel"); + + std::vector results(reference.size(), 99); + error = + clEnqueueReadBuffer(queue, dst, CL_TRUE, 0, results.size() * sizeof(T), + results.data(), 0, NULL, NULL); + test_error(error, "Unable to read data after test kernel"); + + if (results != reference) + { + log_error("Result buffer did not match reference buffer!\n"); + return TEST_FAIL; + } + + return TEST_PASS; +} + +template +static int test_type(cl_device_id device, cl_context context, + cl_command_queue queue) +{ + log_info(" testing type %s\n", TestInfo::deviceTypeName); + + return test_vectype(device, context, queue) + | test_vectype(device, context, queue) + | test_vectype(device, context, queue) + | test_vectype(device, context, queue) + | test_vectype(device, context, queue) + | test_vectype(device, context, queue); +} + +int test_extended_bit_ops_reverse(cl_device_id device, cl_context context, + cl_command_queue queue, int num_elements) +{ + if (is_extension_available(device, "cl_khr_extended_bit_ops")) + { + int result = TEST_PASS; + + result |= test_type(device, context, queue); + result |= test_type(device, context, queue); + result |= test_type(device, context, queue); + result |= test_type(device, context, queue); + result |= test_type(device, context, queue); + result |= test_type(device, context, queue); + if (gHasLong) + { + result |= test_type(device, context, queue); + result |= test_type(device, context, queue); + } + return result; + } + + log_info("cl_khr_extended_bit_ops is not supported\n"); + return TEST_SKIPPED_ITSELF; +}