mirror of
https://github.com/KhronosGroup/OpenCL-CTS.git
synced 2026-03-19 14:09:03 +00:00
2fcdde96d2
* updated reduce test * switched all reduce tests to new framework * switch over scans to new framework * remove old files * minor fixes * add type type name to the kernel name * fix Windows build and warnings * address review comments
457 lines
15 KiB
C++
457 lines
15 KiB
C++
//
|
|
// 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
|
|
//
|
|
// http://www.apache.org/licenses/LICENSE-2.0
|
|
//
|
|
// Unless required by applicable law or agreed to in writing, software
|
|
// distributed under the License is distributed on an "AS IS" BASIS,
|
|
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
|
|
// See the License for the specific language governing permissions and
|
|
// limitations under the License.
|
|
//
|
|
#include "harness/compat.h"
|
|
|
|
#include <algorithm>
|
|
#include <limits>
|
|
#include <vector>
|
|
|
|
#include "procs.h"
|
|
|
|
static std::string make_kernel_string(const std::string &type,
|
|
const std::string &kernelName,
|
|
const std::string &func)
|
|
{
|
|
// Build a kernel string of the form:
|
|
// __kernel void KERNEL_NAME(global TYPE *input, global TYPE *output) {
|
|
// int tid = get_global_id(0);
|
|
// output[tid] = FUNC(input[tid]);
|
|
// }
|
|
|
|
std::ostringstream os;
|
|
os << "__kernel void " << kernelName << "(global " << type
|
|
<< " *input, global " << type << " *output) {\n";
|
|
os << " int tid = get_global_id(0);\n";
|
|
os << " output[tid] = " << func << "(input[tid]);\n";
|
|
os << "}\n";
|
|
return os.str();
|
|
}
|
|
|
|
template <typename T> struct TestTypeInfo
|
|
{
|
|
};
|
|
|
|
template <> struct TestTypeInfo<cl_int>
|
|
{
|
|
static constexpr const char *deviceName = "int";
|
|
};
|
|
|
|
template <> struct TestTypeInfo<cl_uint>
|
|
{
|
|
static constexpr const char *deviceName = "uint";
|
|
};
|
|
|
|
template <> struct TestTypeInfo<cl_long>
|
|
{
|
|
static constexpr const char *deviceName = "long";
|
|
};
|
|
|
|
template <> struct TestTypeInfo<cl_ulong>
|
|
{
|
|
static constexpr const char *deviceName = "ulong";
|
|
};
|
|
|
|
template <typename T> struct Add
|
|
{
|
|
using Type = T;
|
|
static constexpr const char *opName = "add";
|
|
static constexpr T identityValue = 0;
|
|
static T combine(T a, T b) { return a + b; }
|
|
};
|
|
|
|
template <typename T> struct Max
|
|
{
|
|
using Type = T;
|
|
static constexpr const char *opName = "max";
|
|
static constexpr T identityValue = std::numeric_limits<T>::min();
|
|
static T combine(T a, T b) { return std::max(a, b); }
|
|
};
|
|
|
|
template <typename T> struct Min
|
|
{
|
|
using Type = T;
|
|
static constexpr const char *opName = "min";
|
|
static constexpr T identityValue = std::numeric_limits<T>::max();
|
|
static T combine(T a, T b) { return std::min(a, b); }
|
|
};
|
|
|
|
template <typename C> struct Reduce
|
|
{
|
|
using Type = typename C::Type;
|
|
|
|
static constexpr const char *testName = "work_group_reduce";
|
|
static constexpr const char *testOpName = C::opName;
|
|
static constexpr const char *deviceTypeName =
|
|
TestTypeInfo<Type>::deviceName;
|
|
static constexpr const char *kernelName = "test_wg_reduce";
|
|
static int verify(Type *inptr, Type *outptr, size_t n_elems,
|
|
size_t max_wg_size)
|
|
{
|
|
for (size_t i = 0; i < n_elems; i += max_wg_size)
|
|
{
|
|
size_t wg_size = std::min(max_wg_size, n_elems - i);
|
|
|
|
Type result = C::identityValue;
|
|
for (size_t j = 0; j < wg_size; j++)
|
|
{
|
|
result = C::combine(result, inptr[i + j]);
|
|
}
|
|
|
|
for (size_t j = 0; j < wg_size; j++)
|
|
{
|
|
if (result != outptr[i + j])
|
|
{
|
|
log_info("%s_%s: Error at %zu\n", testName, testOpName,
|
|
i + j);
|
|
return -1;
|
|
}
|
|
}
|
|
}
|
|
return 0;
|
|
}
|
|
};
|
|
|
|
template <typename C> struct ScanInclusive
|
|
{
|
|
using Type = typename C::Type;
|
|
|
|
static constexpr const char *testName = "work_group_scan_inclusive";
|
|
static constexpr const char *testOpName = C::opName;
|
|
static constexpr const char *deviceTypeName =
|
|
TestTypeInfo<Type>::deviceName;
|
|
static constexpr const char *kernelName = "test_wg_scan_inclusive";
|
|
static int verify(Type *inptr, Type *outptr, size_t n_elems,
|
|
size_t max_wg_size)
|
|
{
|
|
for (size_t i = 0; i < n_elems; i += max_wg_size)
|
|
{
|
|
size_t wg_size = std::min(max_wg_size, n_elems - i);
|
|
|
|
Type result = C::identityValue;
|
|
for (size_t j = 0; j < wg_size; ++j)
|
|
{
|
|
result = C::combine(result, inptr[i + j]);
|
|
if (result != outptr[i + j])
|
|
{
|
|
log_info("%s_%s: Error at %zu\n", testName, testOpName,
|
|
i + j);
|
|
return -1;
|
|
}
|
|
}
|
|
}
|
|
return 0;
|
|
}
|
|
};
|
|
|
|
template <typename C> struct ScanExclusive
|
|
{
|
|
using Type = typename C::Type;
|
|
|
|
static constexpr const char *testName = "work_group_scan_exclusive";
|
|
static constexpr const char *testOpName = C::opName;
|
|
static constexpr const char *deviceTypeName =
|
|
TestTypeInfo<Type>::deviceName;
|
|
static constexpr const char *kernelName = "test_wg_scan_exclusive";
|
|
static int verify(Type *inptr, Type *outptr, size_t n_elems,
|
|
size_t max_wg_size)
|
|
{
|
|
for (size_t i = 0; i < n_elems; i += max_wg_size)
|
|
{
|
|
size_t wg_size = std::min(max_wg_size, n_elems - i);
|
|
|
|
Type result = C::identityValue;
|
|
for (size_t j = 0; j < wg_size; ++j)
|
|
{
|
|
if (result != outptr[i + j])
|
|
{
|
|
log_info("%s_%s: Error at %zu\n", testName, testOpName,
|
|
i + j);
|
|
return -1;
|
|
}
|
|
result = C::combine(result, inptr[i + j]);
|
|
}
|
|
}
|
|
return 0;
|
|
}
|
|
};
|
|
|
|
template <typename TestInfo>
|
|
static int run_test(cl_device_id device, cl_context context,
|
|
cl_command_queue queue, int n_elems)
|
|
{
|
|
using T = typename TestInfo::Type;
|
|
|
|
cl_int err = CL_SUCCESS;
|
|
|
|
clProgramWrapper program;
|
|
clKernelWrapper kernel;
|
|
|
|
std::string funcName = TestInfo::testName;
|
|
funcName += "_";
|
|
funcName += TestInfo::testOpName;
|
|
|
|
std::string kernelName = TestInfo::kernelName;
|
|
kernelName += "_";
|
|
kernelName += TestInfo::testOpName;
|
|
kernelName += "_";
|
|
kernelName += TestInfo::deviceTypeName;
|
|
|
|
std::string kernelString =
|
|
make_kernel_string(TestInfo::deviceTypeName, kernelName, funcName);
|
|
|
|
const char *kernel_source = kernelString.c_str();
|
|
err = create_single_kernel_helper(context, &program, &kernel, 1,
|
|
&kernel_source, kernelName.c_str());
|
|
test_error(err, "Unable to create test kernel");
|
|
|
|
size_t wg_size[1];
|
|
err = get_max_allowed_1d_work_group_size_on_device(device, kernel, wg_size);
|
|
test_error(err, "get_max_allowed_1d_work_group_size_on_device failed");
|
|
|
|
clMemWrapper src = clCreateBuffer(context, CL_MEM_READ_WRITE,
|
|
sizeof(T) * n_elems, NULL, &err);
|
|
test_error(err, "Unable to create source buffer");
|
|
|
|
clMemWrapper dst = clCreateBuffer(context, CL_MEM_READ_WRITE,
|
|
sizeof(T) * n_elems, NULL, &err);
|
|
test_error(err, "Unable to create destination buffer");
|
|
|
|
std::vector<T> input_ptr(n_elems);
|
|
|
|
MTdataHolder d(gRandomSeed);
|
|
for (int i = 0; i < n_elems; i++)
|
|
{
|
|
input_ptr[i] = (T)genrand_int64(d);
|
|
}
|
|
|
|
err = clEnqueueWriteBuffer(queue, src, CL_TRUE, 0, sizeof(T) * n_elems,
|
|
input_ptr.data(), 0, NULL, NULL);
|
|
test_error(err, "clWriteBuffer to initialize src buffer failed");
|
|
|
|
err = clSetKernelArg(kernel, 0, sizeof(src), &src);
|
|
test_error(err, "Unable to set src buffer kernel arg");
|
|
err |= clSetKernelArg(kernel, 1, sizeof(dst), &dst);
|
|
test_error(err, "Unable to set dst buffer kernel arg");
|
|
|
|
size_t global_work_size[] = { (size_t)n_elems };
|
|
err = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, global_work_size,
|
|
wg_size, 0, NULL, NULL);
|
|
test_error(err, "Unable to enqueue test kernel");
|
|
|
|
std::vector<T> output_ptr(n_elems);
|
|
|
|
cl_uint dead = 0xdeaddead;
|
|
memset_pattern4(output_ptr.data(), &dead, sizeof(T) * n_elems);
|
|
err = clEnqueueReadBuffer(queue, dst, CL_TRUE, 0, sizeof(T) * n_elems,
|
|
output_ptr.data(), 0, NULL, NULL);
|
|
test_error(err, "clEnqueueReadBuffer to read read dst buffer failed");
|
|
|
|
if (TestInfo::verify(input_ptr.data(), output_ptr.data(), n_elems,
|
|
wg_size[0]))
|
|
{
|
|
log_error("%s_%s %s failed\n", TestInfo::testName, TestInfo::testOpName,
|
|
TestInfo::deviceTypeName);
|
|
return TEST_FAIL;
|
|
}
|
|
|
|
log_info("%s_%s %s passed\n", TestInfo::testName, TestInfo::testOpName,
|
|
TestInfo::deviceTypeName);
|
|
return TEST_PASS;
|
|
}
|
|
|
|
int test_work_group_reduce_add(cl_device_id device, cl_context context,
|
|
cl_command_queue queue, int n_elems)
|
|
{
|
|
int result = TEST_PASS;
|
|
|
|
result |= run_test<Reduce<Add<cl_int>>>(device, context, queue, n_elems);
|
|
result |= run_test<Reduce<Add<cl_uint>>>(device, context, queue, n_elems);
|
|
|
|
if (gHasLong)
|
|
{
|
|
result |=
|
|
run_test<Reduce<Add<cl_long>>>(device, context, queue, n_elems);
|
|
result |=
|
|
run_test<Reduce<Add<cl_ulong>>>(device, context, queue, n_elems);
|
|
}
|
|
|
|
return result;
|
|
}
|
|
|
|
int test_work_group_reduce_max(cl_device_id device, cl_context context,
|
|
cl_command_queue queue, int n_elems)
|
|
{
|
|
int result = TEST_PASS;
|
|
|
|
result |= run_test<Reduce<Max<cl_int>>>(device, context, queue, n_elems);
|
|
result |= run_test<Reduce<Max<cl_uint>>>(device, context, queue, n_elems);
|
|
|
|
if (gHasLong)
|
|
{
|
|
result |=
|
|
run_test<Reduce<Max<cl_long>>>(device, context, queue, n_elems);
|
|
result |=
|
|
run_test<Reduce<Max<cl_ulong>>>(device, context, queue, n_elems);
|
|
}
|
|
|
|
return result;
|
|
}
|
|
|
|
int test_work_group_reduce_min(cl_device_id device, cl_context context,
|
|
cl_command_queue queue, int n_elems)
|
|
{
|
|
int result = TEST_PASS;
|
|
|
|
result |= run_test<Reduce<Min<cl_int>>>(device, context, queue, n_elems);
|
|
result |= run_test<Reduce<Min<cl_uint>>>(device, context, queue, n_elems);
|
|
|
|
if (gHasLong)
|
|
{
|
|
result |=
|
|
run_test<Reduce<Min<cl_long>>>(device, context, queue, n_elems);
|
|
result |=
|
|
run_test<Reduce<Min<cl_ulong>>>(device, context, queue, n_elems);
|
|
}
|
|
|
|
return result;
|
|
}
|
|
|
|
int test_work_group_scan_inclusive_add(cl_device_id device, cl_context context,
|
|
cl_command_queue queue, int n_elems)
|
|
{
|
|
int result = TEST_PASS;
|
|
|
|
result |=
|
|
run_test<ScanInclusive<Add<cl_int>>>(device, context, queue, n_elems);
|
|
result |=
|
|
run_test<ScanInclusive<Add<cl_uint>>>(device, context, queue, n_elems);
|
|
|
|
if (gHasLong)
|
|
{
|
|
result |= run_test<ScanInclusive<Add<cl_long>>>(device, context, queue,
|
|
n_elems);
|
|
result |= run_test<ScanInclusive<Add<cl_ulong>>>(device, context, queue,
|
|
n_elems);
|
|
}
|
|
|
|
return result;
|
|
}
|
|
|
|
int test_work_group_scan_inclusive_max(cl_device_id device, cl_context context,
|
|
cl_command_queue queue, int n_elems)
|
|
{
|
|
int result = TEST_PASS;
|
|
|
|
result |=
|
|
run_test<ScanInclusive<Max<cl_int>>>(device, context, queue, n_elems);
|
|
result |=
|
|
run_test<ScanInclusive<Max<cl_uint>>>(device, context, queue, n_elems);
|
|
|
|
if (gHasLong)
|
|
{
|
|
result |= run_test<ScanInclusive<Max<cl_long>>>(device, context, queue,
|
|
n_elems);
|
|
result |= run_test<ScanInclusive<Max<cl_ulong>>>(device, context, queue,
|
|
n_elems);
|
|
}
|
|
|
|
return result;
|
|
}
|
|
|
|
int test_work_group_scan_inclusive_min(cl_device_id device, cl_context context,
|
|
cl_command_queue queue, int n_elems)
|
|
{
|
|
int result = TEST_PASS;
|
|
|
|
result |=
|
|
run_test<ScanInclusive<Min<cl_int>>>(device, context, queue, n_elems);
|
|
result |=
|
|
run_test<ScanInclusive<Min<cl_uint>>>(device, context, queue, n_elems);
|
|
|
|
if (gHasLong)
|
|
{
|
|
result |= run_test<ScanInclusive<Min<cl_long>>>(device, context, queue,
|
|
n_elems);
|
|
result |= run_test<ScanInclusive<Min<cl_ulong>>>(device, context, queue,
|
|
n_elems);
|
|
}
|
|
|
|
return result;
|
|
}
|
|
|
|
int test_work_group_scan_exclusive_add(cl_device_id device, cl_context context,
|
|
cl_command_queue queue, int n_elems)
|
|
{
|
|
int result = TEST_PASS;
|
|
|
|
result |=
|
|
run_test<ScanExclusive<Add<cl_int>>>(device, context, queue, n_elems);
|
|
result |=
|
|
run_test<ScanExclusive<Add<cl_uint>>>(device, context, queue, n_elems);
|
|
|
|
if (gHasLong)
|
|
{
|
|
result |= run_test<ScanExclusive<Add<cl_long>>>(device, context, queue,
|
|
n_elems);
|
|
result |= run_test<ScanExclusive<Add<cl_ulong>>>(device, context, queue,
|
|
n_elems);
|
|
}
|
|
|
|
return result;
|
|
}
|
|
|
|
int test_work_group_scan_exclusive_max(cl_device_id device, cl_context context,
|
|
cl_command_queue queue, int n_elems)
|
|
{
|
|
int result = TEST_PASS;
|
|
|
|
result |=
|
|
run_test<ScanExclusive<Max<cl_int>>>(device, context, queue, n_elems);
|
|
result |=
|
|
run_test<ScanExclusive<Max<cl_uint>>>(device, context, queue, n_elems);
|
|
|
|
if (gHasLong)
|
|
{
|
|
result |= run_test<ScanExclusive<Max<cl_long>>>(device, context, queue,
|
|
n_elems);
|
|
result |= run_test<ScanExclusive<Max<cl_ulong>>>(device, context, queue,
|
|
n_elems);
|
|
}
|
|
|
|
return result;
|
|
}
|
|
|
|
int test_work_group_scan_exclusive_min(cl_device_id device, cl_context context,
|
|
cl_command_queue queue, int n_elems)
|
|
{
|
|
int result = TEST_PASS;
|
|
|
|
result |=
|
|
run_test<ScanExclusive<Min<cl_int>>>(device, context, queue, n_elems);
|
|
result |=
|
|
run_test<ScanExclusive<Min<cl_uint>>>(device, context, queue, n_elems);
|
|
|
|
if (gHasLong)
|
|
{
|
|
result |= run_test<ScanExclusive<Min<cl_long>>>(device, context, queue,
|
|
n_elems);
|
|
result |= run_test<ScanExclusive<Min<cl_ulong>>>(device, context, queue,
|
|
n_elems);
|
|
}
|
|
|
|
return result;
|
|
}
|