mirror of
https://github.com/KhronosGroup/OpenCL-CTS.git
synced 2026-03-19 06:09:01 +00:00
add tests for cl_khr_extended_bit_ops (#1232)
* very basic test infrastructure * move test to integer_ops * added bitfield_reverse test * added test for bitfield insert * add a separate file for bitfield_extract * cleaned up implementation and all tests are passing * rename helper file so it can be used by other tests * remove temporary hacks * fix formatting * address code review comments * remove duplicated code after rebase * fix formatting a few naming changes for consistency more changes for consistency * fix copyright dates * remove unused variable
This commit is contained in:
Ben Ashbaugh
@@ -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)
|
||||
|
||||
@@ -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),
|
||||
|
||||
@@ -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);
|
||||
|
||||
287
test_conformance/integer_ops/test_extended_bit_ops_extract.cpp
Normal file
287
test_conformance/integer_ops/test_extended_bit_ops_extract.cpp
Normal file
@@ -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 <algorithm>
|
||||
#include <numeric>
|
||||
#include <string>
|
||||
#include <type_traits>
|
||||
#include <vector>
|
||||
|
||||
#include "procs.h"
|
||||
#include "harness/integer_ops_test_info.h"
|
||||
#include "harness/testHarness.h"
|
||||
|
||||
template <typename T>
|
||||
static typename std::make_unsigned<T>::type
|
||||
arithmetic_shift_right(T tx, cl_uint count)
|
||||
{
|
||||
typedef typename std::make_unsigned<T>::type unsigned_t;
|
||||
unsigned_t x = static_cast<unsigned_t>(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 <typename T>
|
||||
static typename std::make_unsigned<T>::type
|
||||
cpu_bit_extract_signed(T tbase, cl_uint offset, cl_uint count)
|
||||
{
|
||||
typedef typename std::make_signed<T>::type unsigned_t;
|
||||
|
||||
assert(offset <= sizeof(T) * 8);
|
||||
assert(count <= sizeof(T) * 8);
|
||||
assert(offset + count <= sizeof(T) * 8);
|
||||
|
||||
unsigned_t base = static_cast<unsigned_t>(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 <typename T>
|
||||
static typename std::make_unsigned<T>::type
|
||||
cpu_bit_extract_unsigned(T tbase, cl_uint offset, cl_uint count)
|
||||
{
|
||||
typedef typename std::make_unsigned<T>::type unsigned_t;
|
||||
|
||||
assert(offset <= sizeof(T) * 8);
|
||||
assert(count <= sizeof(T) * 8);
|
||||
assert(offset + count <= sizeof(T) * 8);
|
||||
|
||||
unsigned_t base = static_cast<unsigned_t>(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 <typename T, size_t N>
|
||||
static void
|
||||
calculate_reference(std::vector<typename std::make_unsigned<T>::type>& sref,
|
||||
std::vector<typename std::make_unsigned<T>::type>& uref,
|
||||
const std::vector<T>& 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 <typename T, size_t N>
|
||||
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<T>::type unsigned_t;
|
||||
|
||||
cl_int error = CL_SUCCESS;
|
||||
|
||||
clProgramWrapper program;
|
||||
clKernelWrapper kernel;
|
||||
|
||||
std::string buildOptions;
|
||||
buildOptions += " -DTYPE=";
|
||||
buildOptions +=
|
||||
TestInfo<T>::deviceTypeName + ((N > 1) ? std::to_string(N) : "");
|
||||
buildOptions += " -DSIGNED_TYPE=";
|
||||
buildOptions +=
|
||||
TestInfo<T>::deviceTypeNameSigned + ((N > 1) ? std::to_string(N) : "");
|
||||
buildOptions += " -DUNSIGNED_TYPE=";
|
||||
buildOptions += TestInfo<T>::deviceTypeNameUnsigned
|
||||
+ ((N > 1) ? std::to_string(N) : "");
|
||||
buildOptions += " -DBASETYPE=";
|
||||
buildOptions += TestInfo<T>::deviceTypeName;
|
||||
buildOptions += " -DSIGNED_BASETYPE=";
|
||||
buildOptions += TestInfo<T>::deviceTypeNameSigned;
|
||||
buildOptions += " -DUNSIGNED_BASETYPE=";
|
||||
buildOptions += TestInfo<T>::deviceTypeNameUnsigned;
|
||||
|
||||
const size_t ELEMENTS_TO_TEST = (sizeof(T) * 8 + 1) * (sizeof(T) * 8 + 1);
|
||||
|
||||
std::vector<T> base(ELEMENTS_TO_TEST * N);
|
||||
fill_vector_with_random_data(base);
|
||||
|
||||
std::vector<unsigned_t> sreference;
|
||||
std::vector<unsigned_t> ureference;
|
||||
calculate_reference<T, N>(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<unsigned_t> 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<unsigned_t> 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 <typename T>
|
||||
static int test_type(cl_device_id device, cl_context context,
|
||||
cl_command_queue queue)
|
||||
{
|
||||
log_info(" testing type %s\n", TestInfo<T>::deviceTypeName);
|
||||
|
||||
return test_vectype<T, 1>(device, context, queue)
|
||||
| test_vectype<T, 2>(device, context, queue)
|
||||
| test_vectype<T, 3>(device, context, queue)
|
||||
| test_vectype<T, 4>(device, context, queue)
|
||||
| test_vectype<T, 8>(device, context, queue)
|
||||
| test_vectype<T, 16>(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<cl_char>(device, context, queue);
|
||||
result |= test_type<cl_uchar>(device, context, queue);
|
||||
result |= test_type<cl_short>(device, context, queue);
|
||||
result |= test_type<cl_ushort>(device, context, queue);
|
||||
result |= test_type<cl_int>(device, context, queue);
|
||||
result |= test_type<cl_uint>(device, context, queue);
|
||||
if (gHasLong)
|
||||
{
|
||||
result |= test_type<cl_long>(device, context, queue);
|
||||
result |= test_type<cl_ulong>(device, context, queue);
|
||||
}
|
||||
return result;
|
||||
}
|
||||
|
||||
log_info("cl_khr_extended_bit_ops is not supported\n");
|
||||
return TEST_SKIPPED_ITSELF;
|
||||
}
|
||||
214
test_conformance/integer_ops/test_extended_bit_ops_insert.cpp
Normal file
214
test_conformance/integer_ops/test_extended_bit_ops_insert.cpp
Normal file
@@ -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 <algorithm>
|
||||
#include <numeric>
|
||||
#include <string>
|
||||
#include <type_traits>
|
||||
#include <vector>
|
||||
|
||||
#include "procs.h"
|
||||
#include "harness/integer_ops_test_info.h"
|
||||
#include "harness/testHarness.h"
|
||||
|
||||
template <typename T>
|
||||
static typename std::make_unsigned<T>::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<cl_ulong>(tbase);
|
||||
cl_ulong insert = static_cast<cl_ulong>(tinsert);
|
||||
|
||||
cl_ulong mask = (count < 64) ? ((1ULL << count) - 1) << offset : ~0ULL;
|
||||
cl_ulong result = ((insert << offset) & mask) | (base & ~mask);
|
||||
|
||||
return static_cast<typename std::make_unsigned<T>::type>(result);
|
||||
}
|
||||
|
||||
template <typename T, size_t N>
|
||||
static void
|
||||
calculate_reference(std::vector<typename std::make_unsigned<T>::type>& ref,
|
||||
const std::vector<T>& base, const std::vector<T>& 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 <typename T, size_t N>
|
||||
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<T>::type unsigned_t;
|
||||
|
||||
cl_int error = CL_SUCCESS;
|
||||
|
||||
clProgramWrapper program;
|
||||
clKernelWrapper kernel;
|
||||
|
||||
std::string buildOptions{ "-DTYPE=" };
|
||||
buildOptions += TestInfo<T>::deviceTypeName;
|
||||
if (N > 1)
|
||||
{
|
||||
buildOptions += std::to_string(N);
|
||||
}
|
||||
buildOptions += " -DBASETYPE=";
|
||||
buildOptions += TestInfo<T>::deviceTypeName;
|
||||
|
||||
const size_t ELEMENTS_TO_TEST = (sizeof(T) * 8 + 1) * (sizeof(T) * 8 + 1);
|
||||
|
||||
std::vector<T> base(ELEMENTS_TO_TEST * N);
|
||||
std::fill(base.begin(), base.end(), static_cast<T>(0xA5A5A5A5A5A5A5A5ULL));
|
||||
|
||||
std::vector<T> insert(ELEMENTS_TO_TEST * N);
|
||||
fill_vector_with_random_data(insert);
|
||||
|
||||
std::vector<unsigned_t> reference;
|
||||
calculate_reference<T, N>(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<unsigned_t> 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 <typename T>
|
||||
static int test_type(cl_device_id device, cl_context context,
|
||||
cl_command_queue queue)
|
||||
{
|
||||
log_info(" testing type %s\n", TestInfo<T>::deviceTypeName);
|
||||
|
||||
return test_vectype<T, 1>(device, context, queue)
|
||||
| test_vectype<T, 2>(device, context, queue)
|
||||
| test_vectype<T, 3>(device, context, queue)
|
||||
| test_vectype<T, 4>(device, context, queue)
|
||||
| test_vectype<T, 8>(device, context, queue)
|
||||
| test_vectype<T, 16>(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<cl_char>(device, context, queue);
|
||||
result |= test_type<cl_uchar>(device, context, queue);
|
||||
result |= test_type<cl_short>(device, context, queue);
|
||||
result |= test_type<cl_ushort>(device, context, queue);
|
||||
result |= test_type<cl_int>(device, context, queue);
|
||||
result |= test_type<cl_uint>(device, context, queue);
|
||||
if (gHasLong)
|
||||
{
|
||||
result |= test_type<cl_long>(device, context, queue);
|
||||
result |= test_type<cl_ulong>(device, context, queue);
|
||||
}
|
||||
return result;
|
||||
}
|
||||
|
||||
log_info("cl_khr_extended_bit_ops is not supported\n");
|
||||
return TEST_SKIPPED_ITSELF;
|
||||
}
|
||||
177
test_conformance/integer_ops/test_extended_bit_ops_reverse.cpp
Normal file
177
test_conformance/integer_ops/test_extended_bit_ops_reverse.cpp
Normal file
@@ -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 <algorithm>
|
||||
#include <numeric>
|
||||
#include <string>
|
||||
#include <vector>
|
||||
|
||||
#include "procs.h"
|
||||
#include "harness/integer_ops_test_info.h"
|
||||
#include "harness/testHarness.h"
|
||||
|
||||
template <typename T> 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 <typename T>
|
||||
static void calculate_reference(std::vector<T>& ref, const std::vector<T>& 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 <typename T, size_t N>
|
||||
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<T>::deviceTypeName;
|
||||
if (N > 1)
|
||||
{
|
||||
buildOptions += std::to_string(N);
|
||||
}
|
||||
buildOptions += " -DBASETYPE=";
|
||||
buildOptions += TestInfo<T>::deviceTypeName;
|
||||
|
||||
const size_t ELEMENTS_TO_TEST = 65536;
|
||||
std::vector<T> base(ELEMENTS_TO_TEST * N);
|
||||
fill_vector_with_random_data(base);
|
||||
|
||||
std::vector<T> 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<T> 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 <typename T>
|
||||
static int test_type(cl_device_id device, cl_context context,
|
||||
cl_command_queue queue)
|
||||
{
|
||||
log_info(" testing type %s\n", TestInfo<T>::deviceTypeName);
|
||||
|
||||
return test_vectype<T, 1>(device, context, queue)
|
||||
| test_vectype<T, 2>(device, context, queue)
|
||||
| test_vectype<T, 3>(device, context, queue)
|
||||
| test_vectype<T, 4>(device, context, queue)
|
||||
| test_vectype<T, 8>(device, context, queue)
|
||||
| test_vectype<T, 16>(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<cl_char>(device, context, queue);
|
||||
result |= test_type<cl_uchar>(device, context, queue);
|
||||
result |= test_type<cl_short>(device, context, queue);
|
||||
result |= test_type<cl_ushort>(device, context, queue);
|
||||
result |= test_type<cl_int>(device, context, queue);
|
||||
result |= test_type<cl_uint>(device, context, queue);
|
||||
if (gHasLong)
|
||||
{
|
||||
result |= test_type<cl_long>(device, context, queue);
|
||||
result |= test_type<cl_ulong>(device, context, queue);
|
||||
}
|
||||
return result;
|
||||
}
|
||||
|
||||
log_info("cl_khr_extended_bit_ops is not supported\n");
|
||||
return TEST_SKIPPED_ITSELF;
|
||||
}
|
||||
Reference in New Issue
Block a user