mirror of
https://github.com/KhronosGroup/OpenCL-CTS.git
synced 2026-03-19 14:09:03 +00:00
e15c6eb760
Similar to #2219, we see "'fpclassify': ambiguous call" error in test_conformance\basic\test_fpmath.cpp due to missing constexpr at https://github.com/KhronosGroup/OpenCL-CTS/blob/9265cbb2c274/test_conformance/basic/test_fpmath.cpp#L104 This PR fixes the issue by moving utility function isnan_fp in testHarness.h and use it. Note this PR doesn't modify use of isnan in many tests where only float/double values are checked.
380 lines
11 KiB
C++
380 lines
11 KiB
C++
//
|
|
// 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/compat.h"
|
|
#include "harness/mathHelpers.h"
|
|
#include "harness/rounding_mode.h"
|
|
#include "harness/stringHelpers.h"
|
|
|
|
#include <CL/cl_half.h>
|
|
|
|
#include <stdio.h>
|
|
#include <stdlib.h>
|
|
#include <string.h>
|
|
#include <sys/types.h>
|
|
#include <sys/stat.h>
|
|
|
|
#include <algorithm>
|
|
#include <functional>
|
|
#include <map>
|
|
#include <string>
|
|
#include <vector>
|
|
|
|
#include "testBase.h"
|
|
|
|
extern cl_half_rounding_mode halfRoundingMode;
|
|
|
|
namespace {
|
|
|
|
const char *fp_kernel_code = R"(
|
|
%s
|
|
__kernel void test_fp(__global TYPE *srcA, __global TYPE *srcB, __global TYPE *dst)
|
|
{
|
|
int tid = get_global_id(0);
|
|
|
|
dst[tid] = srcA[tid] OP srcB[tid];
|
|
})";
|
|
|
|
#define HFF(num) cl_half_from_float(num, halfRoundingMode)
|
|
#define HTF(num) cl_half_to_float(num)
|
|
|
|
template <typename T> double toDouble(T val)
|
|
{
|
|
if (std::is_same<cl_half, T>::value)
|
|
return HTF(val);
|
|
else
|
|
return val;
|
|
}
|
|
|
|
cl_half half_plus(cl_half a, cl_half b)
|
|
{
|
|
return HFF(std::plus<float>()(HTF(a), HTF(b)));
|
|
}
|
|
|
|
cl_half half_minus(cl_half a, cl_half b)
|
|
{
|
|
return HFF(std::minus<float>()(HTF(a), HTF(b)));
|
|
}
|
|
|
|
cl_half half_mult(cl_half a, cl_half b)
|
|
{
|
|
return HFF(std::multiplies<float>()(HTF(a), HTF(b)));
|
|
}
|
|
|
|
template <typename T> struct TestDef
|
|
{
|
|
const char op;
|
|
std::function<T(T, T)> ref;
|
|
std::string type_str;
|
|
size_t vec_size;
|
|
};
|
|
|
|
template <typename T>
|
|
int verify_fp(std::vector<T> (&input)[2], std::vector<T> &output,
|
|
const TestDef<T> &test)
|
|
{
|
|
auto &inA = input[0];
|
|
auto &inB = input[1];
|
|
for (size_t i = 0; i < output.size(); i++)
|
|
{
|
|
T r = test.ref(inA[i], inB[i]);
|
|
bool both_nan = false;
|
|
|
|
both_nan = isnan_fp(r) && isnan_fp(output[i]);
|
|
|
|
// If not both nan, check if the result is the same
|
|
if (!both_nan && (r != output[i]))
|
|
{
|
|
log_error("FP math test for type: %s, vec size: %zu, failed at "
|
|
"index %zu, %a '%c' %a, expected %a, get %a\n",
|
|
test.type_str.c_str(), test.vec_size, i, toDouble(inA[i]),
|
|
test.op, toDouble(inB[i]), toDouble(r),
|
|
toDouble(output[i]));
|
|
return -1;
|
|
}
|
|
}
|
|
|
|
return 0;
|
|
}
|
|
|
|
template <typename T> void generate_random_inputs(std::vector<T> (&input)[2])
|
|
{
|
|
RandomSeed seed(gRandomSeed);
|
|
|
|
if (std::is_same<T, float>::value)
|
|
{
|
|
auto random_generator = [&seed]() {
|
|
return get_random_float(-MAKE_HEX_FLOAT(0x1.0p31f, 0x1, 31),
|
|
MAKE_HEX_FLOAT(0x1.0p31f, 0x1, 31), seed);
|
|
};
|
|
for (auto &v : input)
|
|
std::generate(v.begin(), v.end(), random_generator);
|
|
}
|
|
else if (std::is_same<T, double>::value)
|
|
{
|
|
auto random_generator = [&seed]() {
|
|
return get_random_double(-MAKE_HEX_DOUBLE(0x1.0p63, 0x1LL, 63),
|
|
MAKE_HEX_DOUBLE(0x1.0p63, 0x1LL, 63),
|
|
seed);
|
|
};
|
|
for (auto &v : input)
|
|
std::generate(v.begin(), v.end(), random_generator);
|
|
}
|
|
else
|
|
{
|
|
auto random_generator = [&seed]() {
|
|
return HFF(get_random_float(-MAKE_HEX_FLOAT(0x1.0p8f, 0x1, 8),
|
|
MAKE_HEX_FLOAT(0x1.0p8f, 0x1, 8),
|
|
seed));
|
|
};
|
|
for (auto &v : input)
|
|
std::generate(v.begin(), v.end(), random_generator);
|
|
}
|
|
}
|
|
|
|
struct TypesIterator
|
|
{
|
|
using TypeIter = std::tuple<cl_float, cl_half, cl_double>;
|
|
|
|
TypesIterator(cl_device_id deviceID, cl_context context,
|
|
cl_command_queue queue, int num_elems)
|
|
: context(context), queue(queue), fpConfigHalf(0), fpConfigFloat(0),
|
|
num_elements(num_elems)
|
|
{
|
|
// typeid().name one day
|
|
type2name[sizeof(cl_half)] = "half";
|
|
type2name[sizeof(cl_float)] = "float";
|
|
type2name[sizeof(cl_double)] = "double";
|
|
|
|
fp16Support = is_extension_available(deviceID, "cl_khr_fp16");
|
|
fp64Support = is_extension_available(deviceID, "cl_khr_fp64");
|
|
|
|
fpConfigFloat = get_default_rounding_mode(deviceID);
|
|
|
|
if (fp16Support)
|
|
fpConfigHalf =
|
|
get_default_rounding_mode(deviceID, CL_DEVICE_HALF_FP_CONFIG);
|
|
|
|
for_each_elem(it);
|
|
}
|
|
|
|
template <typename T> int test_fpmath(TestDef<T> &test)
|
|
{
|
|
constexpr size_t vecSizes[] = { 1, 2, 4, 8, 16 };
|
|
cl_int err = CL_SUCCESS;
|
|
|
|
std::ostringstream sstr;
|
|
if (std::is_same<T, double>::value)
|
|
sstr << "#pragma OPENCL EXTENSION cl_khr_fp64 : enable\n";
|
|
|
|
if (std::is_same<T, cl_half>::value)
|
|
sstr << "#pragma OPENCL EXTENSION cl_khr_fp16 : enable\n";
|
|
|
|
std::string program_source =
|
|
str_sprintf(std::string(fp_kernel_code), sstr.str().c_str());
|
|
|
|
for (unsigned i = 0; i < ARRAY_SIZE(vecSizes); i++)
|
|
{
|
|
test.vec_size = vecSizes[i];
|
|
|
|
std::ostringstream vecNameStr;
|
|
vecNameStr << test.type_str;
|
|
if (test.vec_size != 1) vecNameStr << test.vec_size;
|
|
|
|
clMemWrapper streams[3];
|
|
clProgramWrapper program;
|
|
clKernelWrapper kernel;
|
|
|
|
size_t length = sizeof(T) * num_elements * test.vec_size;
|
|
|
|
bool isRTZ = false;
|
|
RoundingMode oldMode = kDefaultRoundingMode;
|
|
|
|
|
|
// If we only support rtz mode
|
|
if (std::is_same<T, cl_half>::value)
|
|
{
|
|
if (CL_FP_ROUND_TO_ZERO == fpConfigHalf)
|
|
{
|
|
isRTZ = true;
|
|
oldMode = get_round();
|
|
}
|
|
}
|
|
else if (std::is_same<T, float>::value)
|
|
{
|
|
if (CL_FP_ROUND_TO_ZERO == fpConfigFloat)
|
|
{
|
|
isRTZ = true;
|
|
oldMode = get_round();
|
|
}
|
|
}
|
|
|
|
std::vector<T> inputs[]{
|
|
std::vector<T>(test.vec_size * num_elements),
|
|
std::vector<T>(test.vec_size * num_elements)
|
|
};
|
|
std::vector<T> output =
|
|
std::vector<T>(test.vec_size * num_elements);
|
|
|
|
generate_random_inputs<T>(inputs);
|
|
|
|
for (size_t i = 0; i < ARRAY_SIZE(streams); i++)
|
|
{
|
|
streams[i] = clCreateBuffer(context, CL_MEM_READ_WRITE, length,
|
|
NULL, &err);
|
|
test_error(err, "clCreateBuffer failed.");
|
|
}
|
|
for (size_t i = 0; i < ARRAY_SIZE(inputs); i++)
|
|
{
|
|
err =
|
|
clEnqueueWriteBuffer(queue, streams[i], CL_TRUE, 0, length,
|
|
inputs[i].data(), 0, NULL, NULL);
|
|
test_error(err, "clEnqueueWriteBuffer failed.");
|
|
}
|
|
|
|
std::string build_options = "-DTYPE=";
|
|
build_options.append(vecNameStr.str())
|
|
.append(" -DOP=")
|
|
.append(1, test.op);
|
|
|
|
const char *ptr = program_source.c_str();
|
|
err =
|
|
create_single_kernel_helper(context, &program, &kernel, 1, &ptr,
|
|
"test_fp", build_options.c_str());
|
|
|
|
test_error(err, "create_single_kernel_helper failed");
|
|
|
|
for (size_t i = 0; i < ARRAY_SIZE(streams); i++)
|
|
{
|
|
err =
|
|
clSetKernelArg(kernel, i, sizeof(streams[i]), &streams[i]);
|
|
test_error(err, "clSetKernelArgs failed.");
|
|
}
|
|
|
|
size_t threads[] = { static_cast<size_t>(num_elements) };
|
|
err = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, threads, NULL,
|
|
0, NULL, NULL);
|
|
test_error(err, "clEnqueueNDRangeKernel failed.");
|
|
|
|
err = clEnqueueReadBuffer(queue, streams[2], CL_TRUE, 0, length,
|
|
output.data(), 0, NULL, NULL);
|
|
test_error(err, "clEnqueueReadBuffer failed.");
|
|
|
|
if (isRTZ) set_round(kRoundTowardZero, kfloat);
|
|
|
|
err = verify_fp(inputs, output, test);
|
|
|
|
if (isRTZ) set_round(oldMode, kfloat);
|
|
|
|
test_error(err, "test verification failed");
|
|
log_info("FP '%c' '%s' test passed\n", test.op,
|
|
vecNameStr.str().c_str());
|
|
}
|
|
|
|
return err;
|
|
}
|
|
|
|
template <typename T> int test_fpmath_common()
|
|
{
|
|
int err = TEST_PASS;
|
|
if (std::is_same<cl_half, T>::value)
|
|
{
|
|
TestDef<T> tests[] = { { '+', half_plus, type2name[sizeof(T)] },
|
|
{ '-', half_minus, type2name[sizeof(T)] },
|
|
{ '*', half_mult, type2name[sizeof(T)] } };
|
|
for (auto &test : tests) err |= test_fpmath<T>(test);
|
|
}
|
|
else
|
|
{
|
|
TestDef<T> tests[] = {
|
|
{ '+', std::plus<T>(), type2name[sizeof(T)] },
|
|
{ '-', std::minus<T>(), type2name[sizeof(T)] },
|
|
{ '*', std::multiplies<T>(), type2name[sizeof(T)] }
|
|
};
|
|
for (auto &test : tests) err |= test_fpmath<T>(test);
|
|
}
|
|
|
|
return err;
|
|
}
|
|
|
|
template <typename T> bool skip_type()
|
|
{
|
|
if (std::is_same<double, T>::value && !fp64Support)
|
|
return true;
|
|
else if (std::is_same<cl_half, T>::value && !fp16Support)
|
|
return true;
|
|
return false;
|
|
}
|
|
|
|
template <std::size_t Cnt = 0, typename Type>
|
|
void iterate_type(const Type &t)
|
|
{
|
|
bool doTest = !skip_type<Type>();
|
|
|
|
if (doTest)
|
|
{
|
|
if (test_fpmath_common<Type>())
|
|
{
|
|
throw std::runtime_error("test_fpmath_common failed\n");
|
|
}
|
|
}
|
|
}
|
|
|
|
template <std::size_t Cnt = 0, typename... Tp>
|
|
inline typename std::enable_if<Cnt == sizeof...(Tp), void>::type
|
|
for_each_elem(
|
|
const std::tuple<Tp...> &) // Unused arguments are given no names.
|
|
{}
|
|
|
|
template <std::size_t Cnt = 0, typename... Tp>
|
|
inline typename std::enable_if < Cnt<sizeof...(Tp), void>::type
|
|
for_each_elem(const std::tuple<Tp...> &t)
|
|
{
|
|
iterate_type<Cnt>(std::get<Cnt>(t));
|
|
for_each_elem<Cnt + 1, Tp...>(t);
|
|
}
|
|
|
|
protected:
|
|
TypeIter it;
|
|
|
|
cl_context context;
|
|
cl_command_queue queue;
|
|
|
|
cl_device_fp_config fpConfigHalf;
|
|
cl_device_fp_config fpConfigFloat;
|
|
|
|
bool fp16Support;
|
|
bool fp64Support;
|
|
|
|
int num_elements;
|
|
std::map<size_t, std::string> type2name;
|
|
};
|
|
|
|
} // anonymous namespace
|
|
|
|
REGISTER_TEST(fpmath)
|
|
{
|
|
try
|
|
{
|
|
TypesIterator(device, context, queue, num_elements);
|
|
} catch (const std::runtime_error &e)
|
|
{
|
|
log_error("%s", e.what());
|
|
return TEST_FAIL;
|
|
}
|
|
|
|
return TEST_PASS;
|
|
}
|