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

Deduplicate test_basic int2float/float2int (#1537)

Browse Source 
Merge int2float,float2int.

Signed-off-by: John Kesapides <john.kesapides@arm.com>
This commit is contained in:
John Kesapides
2023-05-24 16:55:25 +01:00
committed by GitHub
parent 957e3b3985
commit 3e8898ffeb
2 changed files with 95 additions and 98 deletions
Download Patch File Download Diff File Expand all files Collapse all files

2
test_conformance/basic/CMakeLists.txt
View File

@@ -11,7 +11,7 @@ set(${MODULE_NAME}_SOURCES
test_multireadimageonefmt.cpp test_multireadimagemultifmt.cpp test_multireadimageonefmt.cpp test_multireadimagemultifmt.cpp
test_imagedim.cpp test_imagedim.cpp
test_vloadstore.cpp test_vloadstore.cpp
test_int2float.cpp test_float2int.cpp test_int2float.cpp
test_createkernelsinprogram.cpp test_createkernelsinprogram.cpp
test_hostptr.cpp test_hostptr.cpp
test_explicit_s2v.cpp test_explicit_s2v.cpp

191
test_conformance/basic/test_int2float.cpp
View File

@@ -1,6 +1,6 @@
// //
// Copyright (c) 2017 The Khronos Group Inc. // Copyright (c) 2017 The Khronos Group Inc.
// //
// Licensed under the Apache License, Version 2.0 (the "License"); // Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License. // you may not use this file except in compliance with the License.
// You may obtain a copy of the License at // You may obtain a copy of the License at
@@ -21,123 +21,120 @@
#include <sys/types.h> #include <sys/types.h>
#include <sys/stat.h> #include <sys/stat.h>
#include <algorithm>
#include <vector>
#include "procs.h" #include "procs.h"
const char *int2float_kernel_code = namespace {
"__kernel void test_int2float(__global int *src, __global float *dst)\n" const char *int2float_kernel_code = R"(
"{\n" __kernel void test_X2Y(__global TYPE_X *src, __global TYPE_Y *dst)
" int tid = get_global_id(0);\n"
"\n"
" dst[tid] = (float)src[tid];\n"
"\n"
"}\n";
int
verify_int2float(cl_int *inptr, cl_float *outptr, int n)
{ {
int i; int tid = get_global_id(0);
for (i=0; i<n; i++) dst[tid] = (TYPE_Y)src[tid];
{
if (outptr[i] != (float)inptr[i])
{
log_error("INT2FLOAT test failed\n");
return -1;
}
}
log_info("INT2FLOAT test passed\n"); })";
return 0;
template <typename T> const char *Type2str() { return ""; }
template <> const char *Type2str<cl_int>() { return "int"; }
template <> const char *Type2str<cl_float>() { return "float"; }
template <typename T> void generate_random_inputs(std::vector<T> &v)
{
RandomSeed seed(gRandomSeed);
auto random_generator = [&seed]() {
return get_random_float(-MAKE_HEX_FLOAT(0x1.0p31f, 0x1, 31),
MAKE_HEX_FLOAT(0x1.0p31f, 0x1, 31), seed);
};
std::generate(v.begin(), v.end(), random_generator);
} }
int template <typename Tx, typename Ty> bool equal_value(Tx a, Ty b)
test_int2float(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
{ {
cl_mem streams[2]; return a == (Tx)b;
cl_int *input_ptr; }
cl_float *output_ptr;
cl_program program; template <typename Tx, typename Ty>
cl_kernel kernel; int verify_X2Y(std::vector<Tx> input, std::vector<Ty> output,
size_t threads[1]; const char *test_name)
int err; {
int i;
MTdata d; if (!std::equal(output.begin(), output.end(), input.begin(),
equal_value<Tx, Ty>))
{
log_error("%s test failed\n", test_name);
return -1;
}
log_info("%s test passed\n", test_name);
return 0;
}
template <typename Tx, typename Ty>
int test_X2Y(cl_device_id device, cl_context context, cl_command_queue queue,
int num_elements, const char *test_name)
{
clMemWrapper streams[2];
clProgramWrapper program;
clKernelWrapper kernel;
int err;
std::vector<Tx> input(num_elements);
std::vector<Ty> output(num_elements);
input_ptr = (cl_int*)malloc(sizeof(cl_int) * num_elements);
output_ptr = (cl_float*)malloc(sizeof(cl_float) * num_elements);
streams[0] = clCreateBuffer(context, CL_MEM_READ_WRITE, streams[0] = clCreateBuffer(context, CL_MEM_READ_WRITE,
sizeof(cl_int) * num_elements, NULL, NULL); sizeof(Tx) * num_elements, nullptr, &err);
if (!streams[0]) test_error(err, "clCreateBuffer failed.");
{
log_error("clCreateBuffer failed\n");
return -1;
}
streams[1] = clCreateBuffer(context, CL_MEM_READ_WRITE, streams[1] = clCreateBuffer(context, CL_MEM_READ_WRITE,
sizeof(cl_float) * num_elements, NULL, NULL); sizeof(Ty) * num_elements, nullptr, &err);
if (!streams[1]) test_error(err, "clCreateBuffer failed.");
{
log_error("clCreateBuffer failed\n");
return -1;
}
d = init_genrand( gRandomSeed ); generate_random_inputs(input);
for (i=0; i<num_elements; i++)
input_ptr[i] = (cl_int)get_random_float(-MAKE_HEX_FLOAT( 0x1.0p31f, 0x1, 31), MAKE_HEX_FLOAT( 0x1.0p31f, 0x1, 31), d);
free_mtdata(d); d = NULL;
err = clEnqueueWriteBuffer(queue, streams[0], CL_TRUE, 0, sizeof(cl_int)*num_elements, (void *)input_ptr, 0, NULL, NULL); err = clEnqueueWriteBuffer(queue, streams[0], CL_TRUE, 0,
if (err != CL_SUCCESS) sizeof(Tx) * num_elements, input.data(), 0,
{ nullptr, nullptr);
log_error("clWriteArray failed\n"); test_error(err, "clEnqueueWriteBuffer failed.");
return -1;
}
err = create_single_kernel_helper(context, &program, &kernel, 1, &int2float_kernel_code, "test_int2float"); std::string build_options;
if (err != CL_SUCCESS) build_options.append("-DTYPE_X=").append(Type2str<Tx>());
{ build_options.append(" -DTYPE_Y=").append(Type2str<Ty>());
log_error("create_single_kernel_helper failed\n"); err = create_single_kernel_helper(context, &program, &kernel, 1,
return -1; &int2float_kernel_code, "test_X2Y",
} build_options.c_str());
test_error(err, "create_single_kernel_helper failed.");
err = clSetKernelArg(kernel, 0, sizeof streams[0], &streams[0]); err = clSetKernelArg(kernel, 0, sizeof streams[0], &streams[0]);
err |= clSetKernelArg(kernel, 1, sizeof streams[1], &streams[1]); err |= clSetKernelArg(kernel, 1, sizeof streams[1], &streams[1]);
if (err != CL_SUCCESS) test_error(err, "clSetKernelArg failed.");
{
log_error("clSetKernelArgs failed\n");
return -1;
}
threads[0] = (size_t)num_elements; size_t threads[] = { (size_t)num_elements };
err = clEnqueueNDRangeKernel( queue, kernel, 1, NULL, threads, NULL, 0, NULL, NULL ); err = clEnqueueNDRangeKernel(queue, kernel, 1, nullptr, threads, nullptr, 0,
if (err != CL_SUCCESS) nullptr, nullptr);
{ test_error(err, "clEnqueueNDRangeKernel failed.");
log_error("clEnqueueNDRangeKernel failed\n");
return -1;
}
err = clEnqueueReadBuffer( queue, streams[1], true, 0, sizeof(cl_float)*num_elements, (void *)output_ptr, 0, NULL, NULL ); err = clEnqueueReadBuffer(queue, streams[1], CL_TRUE, 0,
if (err != CL_SUCCESS) sizeof(Ty) * num_elements, output.data(), 0,
{ nullptr, nullptr);
log_error("clEnqueueReadBuffer failed\n"); test_error(err, "clEnqueueReadBuffer failed.");
return -1;
}
err = verify_int2float(input_ptr, output_ptr, num_elements); err = verify_X2Y(input, output, test_name);
// cleanup
clReleaseMemObject(streams[0]);
clReleaseMemObject(streams[1]);
clReleaseKernel(kernel);
clReleaseProgram(program);
free(input_ptr);
free(output_ptr);
return err; return err;
} }
}
int test_int2float(cl_device_id device, cl_context context,
cl_command_queue queue, int num_elements)
{
return test_X2Y<cl_int, cl_float>(device, context, queue, num_elements,
"INT2FLOAT");
}
int test_float2int(cl_device_id device, cl_context context,
cl_command_queue queue, int num_elements)
{
return test_X2Y<cl_float, cl_int>(device, context, queue, num_elements,
"FLOAT2INT");
}