mirror of
https://github.com/KhronosGroup/OpenCL-CTS.git
synced 2026-03-19 06:09:01 +00:00
Added cl_khr_fp16 extension support for test_explicit_s2v from basic (#1713)
* Added cl_khr_fp16 support for test_explicit_s2v from basic (issue #142, basic) * Cosmetic corrections * cosmetic fix * Added correction to distinguish signed and unsigned char types for ARM architecture tests * Added missing pieces of convertion procedure to support half * Corrected condition to verify if additional pragma is necessary (issue #142, basic) * Add NaN check for half to float conversion * check-format fixes * Add NaN check for all float types Use std::isnan for float/double types. Change-Id: I005bddccaa3f8490ac59b2aa431ed315733ad143 * Fix Ubuntu build error with isnan macro definition Change-Id: I671ed826a9631fbbc66d0aa9b674ab00124c7967 * Check format fixes * NAN define not needed anymore --------- Co-authored-by: Vasu Penugonda <vpenugon@qti.qualcomm.com> Co-authored-by: Sreelakshmi Haridas <sharidas@quicinc.com>
This commit is contained in:
Marcin Hajder
@@ -92,16 +92,8 @@ test_definition test_list[] = {
|
||||
ADD_TEST(image_param),
|
||||
ADD_TEST(image_multipass_integer_coord),
|
||||
ADD_TEST(image_multipass_float_coord),
|
||||
ADD_TEST(explicit_s2v_char),
|
||||
ADD_TEST(explicit_s2v_uchar),
|
||||
ADD_TEST(explicit_s2v_short),
|
||||
ADD_TEST(explicit_s2v_ushort),
|
||||
ADD_TEST(explicit_s2v_int),
|
||||
ADD_TEST(explicit_s2v_uint),
|
||||
ADD_TEST(explicit_s2v_long),
|
||||
ADD_TEST(explicit_s2v_ulong),
|
||||
ADD_TEST(explicit_s2v_float),
|
||||
ADD_TEST(explicit_s2v_double),
|
||||
|
||||
ADD_TEST(explicit_s2v),
|
||||
|
||||
ADD_TEST(enqueue_map_buffer),
|
||||
ADD_TEST(enqueue_map_image),
|
||||
|
||||
@@ -91,16 +91,8 @@ extern int test_vstore_global(cl_device_id deviceID, cl_context context, cl
|
||||
extern int test_vstore_local(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
|
||||
extern int test_vstore_private(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
|
||||
|
||||
extern int test_explicit_s2v_char(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
|
||||
extern int test_explicit_s2v_uchar(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
|
||||
extern int test_explicit_s2v_short(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
|
||||
extern int test_explicit_s2v_ushort(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
|
||||
extern int test_explicit_s2v_int(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
|
||||
extern int test_explicit_s2v_uint(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
|
||||
extern int test_explicit_s2v_long(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
|
||||
extern int test_explicit_s2v_ulong(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
|
||||
extern int test_explicit_s2v_float(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
|
||||
extern int test_explicit_s2v_double(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
|
||||
extern int test_explicit_s2v(cl_device_id deviceID, cl_context context,
|
||||
cl_command_queue queue, int num_elements);
|
||||
|
||||
extern int test_enqueue_map_buffer(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
|
||||
extern int test_enqueue_map_image(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
|
||||
|
||||
@@ -1,6 +1,6 @@
|
||||
//
|
||||
// Copyright (c) 2017 The Khronos Group Inc.
|
||||
//
|
||||
// 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
|
||||
@@ -13,6 +13,8 @@
|
||||
// See the License for the specific language governing permissions and
|
||||
// limitations under the License.
|
||||
//
|
||||
#include <cmath>
|
||||
using std::isnan;
|
||||
#include "harness/compat.h"
|
||||
|
||||
#include <stdio.h>
|
||||
@@ -20,131 +22,112 @@
|
||||
#include <string.h>
|
||||
#include <sys/types.h>
|
||||
#include <sys/stat.h>
|
||||
|
||||
#include <vector>
|
||||
|
||||
#include "procs.h"
|
||||
#include "harness/conversions.h"
|
||||
#include "harness/typeWrappers.h"
|
||||
|
||||
namespace {
|
||||
|
||||
// clang-format off
|
||||
|
||||
#define DECLARE_S2V_IDENT_KERNEL(srctype,dsttype,size) \
|
||||
"__kernel void test_conversion(__global " srctype " *sourceValues, __global " dsttype #size " *destValues )\n" \
|
||||
"{\n" \
|
||||
" int tid = get_global_id(0);\n" \
|
||||
" " srctype " src = sourceValues[tid];\n" \
|
||||
"\n" \
|
||||
" destValues[tid] = (" dsttype #size ")src;\n" \
|
||||
"\n" \
|
||||
"{\n" \
|
||||
" int tid = get_global_id(0);\n" \
|
||||
" " srctype " src = sourceValues[tid];\n" \
|
||||
"\n" \
|
||||
" destValues[tid] = (" dsttype #size ")src;\n" \
|
||||
"\n" \
|
||||
"}\n"
|
||||
|
||||
#define DECLARE_S2V_IDENT_KERNELS(srctype,dsttype) \
|
||||
{ \
|
||||
DECLARE_S2V_IDENT_KERNEL(srctype,#dsttype,2), \
|
||||
DECLARE_S2V_IDENT_KERNEL(srctype,#dsttype,4), \
|
||||
DECLARE_S2V_IDENT_KERNEL(srctype,#dsttype,8), \
|
||||
DECLARE_S2V_IDENT_KERNEL(srctype,#dsttype,16) \
|
||||
}
|
||||
#define DECLARE_S2V_IDENT_KERNELS(srctype, dsttype) \
|
||||
{ \
|
||||
DECLARE_S2V_IDENT_KERNEL(srctype, #dsttype, 2), \
|
||||
DECLARE_S2V_IDENT_KERNEL(srctype, #dsttype, 4), \
|
||||
DECLARE_S2V_IDENT_KERNEL(srctype, #dsttype, 8), \
|
||||
DECLARE_S2V_IDENT_KERNEL(srctype, #dsttype, 16) \
|
||||
}
|
||||
|
||||
#define DECLARE_EMPTY { NULL, NULL, NULL, NULL, NULL }
|
||||
#define DECLARE_EMPTY \
|
||||
{ \
|
||||
NULL, NULL, NULL, NULL, NULL \
|
||||
}
|
||||
|
||||
/* Note: the next four arrays all must match in order and size to the ExplicitTypes enum in conversions.h!!! */
|
||||
/* Note: the next four arrays all must match in order and size to the
|
||||
* ExplicitTypes enum in conversions.h!!! */
|
||||
|
||||
#define DECLARE_S2V_IDENT_KERNELS_SET(srctype) \
|
||||
{ \
|
||||
DECLARE_S2V_IDENT_KERNELS(#srctype,bool), \
|
||||
DECLARE_S2V_IDENT_KERNELS(#srctype,char), \
|
||||
DECLARE_S2V_IDENT_KERNELS(#srctype,uchar), \
|
||||
DECLARE_S2V_IDENT_KERNELS(#srctype,unsigned char), \
|
||||
DECLARE_S2V_IDENT_KERNELS(#srctype,short), \
|
||||
DECLARE_S2V_IDENT_KERNELS(#srctype,ushort), \
|
||||
DECLARE_S2V_IDENT_KERNELS(#srctype,unsigned short), \
|
||||
DECLARE_S2V_IDENT_KERNELS(#srctype,int), \
|
||||
DECLARE_S2V_IDENT_KERNELS(#srctype,uint), \
|
||||
DECLARE_S2V_IDENT_KERNELS(#srctype,unsigned int), \
|
||||
DECLARE_S2V_IDENT_KERNELS(#srctype,long), \
|
||||
DECLARE_S2V_IDENT_KERNELS(#srctype,ulong), \
|
||||
DECLARE_S2V_IDENT_KERNELS(#srctype,unsigned long), \
|
||||
DECLARE_S2V_IDENT_KERNELS(#srctype,float), \
|
||||
DECLARE_EMPTY \
|
||||
}
|
||||
#define DECLARE_S2V_IDENT_KERNELS_SET(srctype) \
|
||||
{ \
|
||||
DECLARE_S2V_IDENT_KERNELS(#srctype, char), \
|
||||
DECLARE_S2V_IDENT_KERNELS(#srctype, uchar), \
|
||||
DECLARE_S2V_IDENT_KERNELS(#srctype, short), \
|
||||
DECLARE_S2V_IDENT_KERNELS(#srctype, ushort), \
|
||||
DECLARE_S2V_IDENT_KERNELS(#srctype, int), \
|
||||
DECLARE_S2V_IDENT_KERNELS(#srctype, uint), \
|
||||
DECLARE_S2V_IDENT_KERNELS(#srctype, long), \
|
||||
DECLARE_S2V_IDENT_KERNELS(#srctype, ulong), \
|
||||
DECLARE_S2V_IDENT_KERNELS(#srctype, float), \
|
||||
DECLARE_S2V_IDENT_KERNELS(#srctype, half), \
|
||||
DECLARE_S2V_IDENT_KERNELS(#srctype, double) \
|
||||
}
|
||||
|
||||
#define DECLARE_EMPTY_SET \
|
||||
{ \
|
||||
DECLARE_EMPTY, \
|
||||
DECLARE_EMPTY, \
|
||||
DECLARE_EMPTY, \
|
||||
DECLARE_EMPTY, \
|
||||
DECLARE_EMPTY, \
|
||||
DECLARE_EMPTY, \
|
||||
DECLARE_EMPTY, \
|
||||
DECLARE_EMPTY, \
|
||||
DECLARE_EMPTY, \
|
||||
DECLARE_EMPTY, \
|
||||
DECLARE_EMPTY, \
|
||||
DECLARE_EMPTY, \
|
||||
DECLARE_EMPTY, \
|
||||
DECLARE_EMPTY, \
|
||||
DECLARE_EMPTY \
|
||||
}
|
||||
#define DECLARE_EMPTY_SET \
|
||||
{ \
|
||||
DECLARE_EMPTY, DECLARE_EMPTY, DECLARE_EMPTY, DECLARE_EMPTY, \
|
||||
DECLARE_EMPTY, DECLARE_EMPTY, DECLARE_EMPTY, DECLARE_EMPTY, \
|
||||
DECLARE_EMPTY, DECLARE_EMPTY, DECLARE_EMPTY, DECLARE_EMPTY, \
|
||||
DECLARE_EMPTY, DECLARE_EMPTY, DECLARE_EMPTY \
|
||||
}
|
||||
|
||||
#define NUM_VEC_TYPES 11
|
||||
|
||||
/* The overall array */
|
||||
const char * kernel_explicit_s2v_set[kNumExplicitTypes][kNumExplicitTypes][5] = {
|
||||
DECLARE_S2V_IDENT_KERNELS_SET(bool),
|
||||
const char * kernel_explicit_s2v_set[NUM_VEC_TYPES][NUM_VEC_TYPES][5] = {
|
||||
DECLARE_S2V_IDENT_KERNELS_SET(char),
|
||||
DECLARE_S2V_IDENT_KERNELS_SET(uchar),
|
||||
DECLARE_S2V_IDENT_KERNELS_SET(unsigned char),
|
||||
DECLARE_S2V_IDENT_KERNELS_SET(short),
|
||||
DECLARE_S2V_IDENT_KERNELS_SET(ushort),
|
||||
DECLARE_S2V_IDENT_KERNELS_SET(unsigned short),
|
||||
DECLARE_S2V_IDENT_KERNELS_SET(int),
|
||||
DECLARE_S2V_IDENT_KERNELS_SET(uint),
|
||||
DECLARE_S2V_IDENT_KERNELS_SET(unsigned int),
|
||||
DECLARE_S2V_IDENT_KERNELS_SET(long),
|
||||
DECLARE_S2V_IDENT_KERNELS_SET(ulong),
|
||||
DECLARE_S2V_IDENT_KERNELS_SET(unsigned long),
|
||||
DECLARE_S2V_IDENT_KERNELS_SET(float),
|
||||
DECLARE_EMPTY_SET
|
||||
DECLARE_S2V_IDENT_KERNELS_SET(half),
|
||||
DECLARE_S2V_IDENT_KERNELS_SET(double)
|
||||
};
|
||||
|
||||
int test_explicit_s2v_function(cl_device_id deviceID, cl_context context, cl_command_queue queue, const char *programSrc,
|
||||
ExplicitType srcType, unsigned int count, ExplicitType destType, unsigned int vecSize, void *inputData )
|
||||
// clang-format on
|
||||
|
||||
bool IsHalfNaN(cl_half v)
|
||||
{
|
||||
// Extract FP16 exponent and mantissa
|
||||
uint16_t h_exp = (((cl_half)v) >> (CL_HALF_MANT_DIG - 1)) & 0x1F;
|
||||
uint16_t h_mant = ((cl_half)v) & 0x3FF;
|
||||
|
||||
// NaN test
|
||||
return (h_exp == 0x1F && h_mant != 0);
|
||||
}
|
||||
|
||||
int test_explicit_s2v_function(cl_context context, cl_command_queue queue,
|
||||
cl_kernel kernel, ExplicitType srcType,
|
||||
unsigned int count, ExplicitType destType,
|
||||
unsigned int vecSize, void *inputData)
|
||||
{
|
||||
clProgramWrapper program;
|
||||
clKernelWrapper kernel;
|
||||
int error;
|
||||
clMemWrapper streams[2];
|
||||
void *outData;
|
||||
unsigned char convertedData[ 8 ]; /* Max type size is 8 bytes */
|
||||
size_t threadSize[3], groupSize[3];
|
||||
unsigned char convertedData[8]; /* Max type size is 8 bytes */
|
||||
unsigned int i, s;
|
||||
unsigned char *inPtr, *outPtr;
|
||||
size_t paramSize, destTypeSize;
|
||||
|
||||
const char* finalProgramSrc[2] = {
|
||||
"", // optional pragma
|
||||
programSrc
|
||||
};
|
||||
|
||||
if (srcType == kDouble || destType == kDouble) {
|
||||
finalProgramSrc[0] = "#pragma OPENCL EXTENSION cl_khr_fp64 : enable\n";
|
||||
}
|
||||
|
||||
|
||||
if( programSrc == NULL )
|
||||
return 0;
|
||||
|
||||
paramSize = get_explicit_type_size( srcType );
|
||||
destTypeSize = get_explicit_type_size( destType );
|
||||
|
||||
size_t destStride = destTypeSize * vecSize;
|
||||
|
||||
outData = malloc( destStride * count );
|
||||
|
||||
if( create_single_kernel_helper( context, &program, &kernel, 2, finalProgramSrc, "test_conversion" ) )
|
||||
{
|
||||
log_info( "****** %s%s *******\n", finalProgramSrc[0], finalProgramSrc[1] );
|
||||
return -1;
|
||||
}
|
||||
std::vector<char> outData(destStride * count);
|
||||
|
||||
streams[0] = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR,
|
||||
paramSize * count, inputData, &error);
|
||||
@@ -170,11 +153,13 @@ int test_explicit_s2v_function(cl_device_id deviceID, cl_context context, cl_com
|
||||
|
||||
/* Now verify the results. Each value should have been duplicated four times, and we should be able to just
|
||||
do a memcpy instead of relying on the actual type of data */
|
||||
error = clEnqueueReadBuffer( queue, streams[1], CL_TRUE, 0, destStride * count, outData, 0, NULL, NULL );
|
||||
error =
|
||||
clEnqueueReadBuffer(queue, streams[1], CL_TRUE, 0, destStride * count,
|
||||
outData.data(), 0, NULL, NULL);
|
||||
test_error( error, "Unable to read output values!" );
|
||||
|
||||
inPtr = (unsigned char *)inputData;
|
||||
outPtr = (unsigned char *)outData;
|
||||
outPtr = (unsigned char *)outData.data();
|
||||
|
||||
for( i = 0; i < count; i++ )
|
||||
{
|
||||
@@ -186,6 +171,28 @@ int test_explicit_s2v_function(cl_device_id deviceID, cl_context context, cl_com
|
||||
{
|
||||
if( memcmp( convertedData, outPtr + destTypeSize * s, destTypeSize ) != 0 )
|
||||
{
|
||||
bool isSrcNaN =
|
||||
(((srcType == kHalf)
|
||||
&& IsHalfNaN(*reinterpret_cast<cl_half *>(inPtr)))
|
||||
|| ((srcType == kFloat)
|
||||
&& isnan(*reinterpret_cast<cl_float *>(inPtr)))
|
||||
|| ((srcType == kDouble)
|
||||
&& isnan(*reinterpret_cast<cl_double *>(inPtr))));
|
||||
bool isDestNaN = (((destType == kHalf)
|
||||
&& IsHalfNaN(*reinterpret_cast<cl_half *>(
|
||||
outPtr + destTypeSize * s)))
|
||||
|| ((destType == kFloat)
|
||||
&& isnan(*reinterpret_cast<cl_float *>(
|
||||
outPtr + destTypeSize * s)))
|
||||
|| ((destType == kDouble)
|
||||
&& isnan(*reinterpret_cast<cl_double *>(
|
||||
outPtr + destTypeSize * s))));
|
||||
|
||||
if (isSrcNaN && isDestNaN)
|
||||
{
|
||||
continue;
|
||||
}
|
||||
|
||||
unsigned int *p = (unsigned int *)outPtr;
|
||||
log_error( "ERROR: Output value %d:%d does not validate for size %d:%d!\n", i, s, vecSize, (int)destTypeSize );
|
||||
log_error( " Input: 0x%0*x\n", (int)( paramSize * 2 ), *(unsigned int *)inPtr & ( 0xffffffff >> ( 32 - paramSize * 8 ) ) );
|
||||
@@ -196,179 +203,185 @@ int test_explicit_s2v_function(cl_device_id deviceID, cl_context context, cl_com
|
||||
inPtr += paramSize;
|
||||
outPtr += destStride;
|
||||
}
|
||||
|
||||
free( outData );
|
||||
|
||||
return 0;
|
||||
}
|
||||
|
||||
int test_explicit_s2v_function_set(cl_device_id deviceID, cl_context context, cl_command_queue queue, ExplicitType srcType,
|
||||
unsigned int count, void *inputData )
|
||||
struct TypesIterator
|
||||
{
|
||||
unsigned int sizes[] = { 2, 4, 8, 16, 0 };
|
||||
int i, dstType, failed = 0;
|
||||
using TypeIter =
|
||||
std::tuple<cl_char, cl_uchar, cl_short, cl_ushort, cl_int, cl_uint,
|
||||
cl_long, cl_ulong, cl_float, cl_half, cl_double>;
|
||||
|
||||
|
||||
for( dstType = kBool; dstType < kNumExplicitTypes; dstType++ )
|
||||
TypesIterator(cl_device_id deviceID, cl_context context,
|
||||
cl_command_queue queue)
|
||||
: dstType(0), srcType(0), context(context), queue(queue)
|
||||
{
|
||||
if( dstType == kDouble && !is_extension_available( deviceID, "cl_khr_fp64" ) )
|
||||
continue;
|
||||
vecTypes = { kChar, kUChar, kShort, kUShort, kInt, kUInt,
|
||||
kLong, kULong, kFloat, kHalf, kDouble };
|
||||
fp16Support = is_extension_available(deviceID, "cl_khr_fp16");
|
||||
fp64Support = is_extension_available(deviceID, "cl_khr_fp64");
|
||||
|
||||
if (( dstType == kLong || dstType == kULong ) && !gHasLong )
|
||||
continue;
|
||||
for_each_src_elem(it);
|
||||
}
|
||||
|
||||
for( i = 0; sizes[i] != 0; i++ )
|
||||
bool skip_type(ExplicitType type)
|
||||
{
|
||||
if ((type == kLong || type == kULong) && !gHasLong)
|
||||
return true;
|
||||
else if (type == kDouble && !fp64Support)
|
||||
return true;
|
||||
else if (type == kHalf && !fp16Support)
|
||||
return true;
|
||||
else if (strchr(get_explicit_type_name(type), ' ') != 0)
|
||||
return true;
|
||||
return false;
|
||||
}
|
||||
|
||||
template <std::size_t Src = 0, typename SrcType>
|
||||
void iterate_src_type(const SrcType &t)
|
||||
{
|
||||
bool doTest = !skip_type(vecTypes[srcType]);
|
||||
if (doTest)
|
||||
{
|
||||
if( dstType != srcType )
|
||||
continue;
|
||||
if( strchr( get_explicit_type_name( (ExplicitType)srcType ), ' ' ) != NULL ||
|
||||
strchr( get_explicit_type_name( (ExplicitType)dstType ), ' ' ) != NULL )
|
||||
continue;
|
||||
SrcType inputData[sample_count];
|
||||
RandomSeed seed(gRandomSeed);
|
||||
generate_random_data(vecTypes[srcType], 128, seed, inputData);
|
||||
|
||||
if( test_explicit_s2v_function( deviceID, context, queue, kernel_explicit_s2v_set[ srcType ][ dstType ][ i ],
|
||||
srcType, count, (ExplicitType)dstType, sizes[ i ], inputData ) != 0 )
|
||||
for_each_dst_elem<0, Src, SrcType>(it, inputData);
|
||||
}
|
||||
|
||||
srcType++;
|
||||
dstType = 0;
|
||||
}
|
||||
|
||||
// crucial to keep it in-sync with ExplicitType
|
||||
bool isExplicitTypeFloating(ExplicitType type) { return (type >= kFloat); }
|
||||
|
||||
template <std::size_t Dst, std::size_t Src, typename SrcType,
|
||||
typename DstType>
|
||||
void iterate_dst_type(const DstType &t, SrcType *inputData)
|
||||
{
|
||||
bool doTest = !skip_type(vecTypes[dstType]);
|
||||
|
||||
doTest = doTest
|
||||
&& ((isExplicitTypeFloating(vecTypes[srcType])
|
||||
&& isExplicitTypeFloating(vecTypes[dstType]))
|
||||
|| (!isExplicitTypeFloating(vecTypes[srcType])
|
||||
&& !isExplicitTypeFloating(vecTypes[dstType])));
|
||||
|
||||
if (doTest)
|
||||
test_explicit_s2v_function_set<SrcType, DstType>(
|
||||
vecTypes[srcType], vecTypes[dstType], inputData);
|
||||
dstType++;
|
||||
}
|
||||
|
||||
template <std::size_t Out = 0, typename... Tp>
|
||||
inline typename std::enable_if<Out == sizeof...(Tp), void>::type
|
||||
for_each_src_elem(
|
||||
const std::tuple<Tp...> &) // Unused arguments are given no names.
|
||||
{}
|
||||
|
||||
template <std::size_t Out = 0, typename... Tp>
|
||||
inline typename std::enable_if < Out<sizeof...(Tp), void>::type
|
||||
for_each_src_elem(const std::tuple<Tp...> &t)
|
||||
{
|
||||
iterate_src_type<Out>(std::get<Out>(t));
|
||||
for_each_src_elem<Out + 1, Tp...>(t);
|
||||
}
|
||||
|
||||
template <std::size_t In = 0, std::size_t Out, typename SrcType,
|
||||
typename... Tp>
|
||||
inline typename std::enable_if<In == sizeof...(Tp), void>::type
|
||||
for_each_dst_elem(const std::tuple<Tp...> &, SrcType *)
|
||||
{}
|
||||
|
||||
template <std::size_t In = 0, std::size_t Out, typename SrcType,
|
||||
typename... Tp>
|
||||
inline typename std::enable_if < In<sizeof...(Tp), void>::type
|
||||
for_each_dst_elem(const std::tuple<Tp...> &t, SrcType *inputData)
|
||||
{
|
||||
iterate_dst_type<In, Out, SrcType>(std::get<In>(t), inputData);
|
||||
for_each_dst_elem<In + 1, Out, SrcType, Tp...>(t, inputData);
|
||||
}
|
||||
|
||||
template <typename SrcType, typename DstType>
|
||||
void test_explicit_s2v_function_set(ExplicitType srcT, ExplicitType dstT,
|
||||
SrcType *inputData)
|
||||
{
|
||||
unsigned int sizes[] = { 2, 4, 8, 16, 0 };
|
||||
|
||||
for (int i = 0; sizes[i] != 0; i++)
|
||||
{
|
||||
clProgramWrapper program;
|
||||
clKernelWrapper kernel;
|
||||
|
||||
char pragma[256] = { 0 };
|
||||
const char *finalProgramSrc[2] = {
|
||||
pragma, // optional pragma
|
||||
kernel_explicit_s2v_set[srcType][dstType][i]
|
||||
};
|
||||
|
||||
std::stringstream sstr;
|
||||
if (srcT == kDouble || dstT == kDouble)
|
||||
sstr << "#pragma OPENCL EXTENSION cl_khr_fp64 : enable\n";
|
||||
|
||||
if (srcT == kHalf || dstT == kHalf)
|
||||
sstr << "#pragma OPENCL EXTENSION cl_khr_fp16 : enable\n";
|
||||
|
||||
snprintf(pragma, sizeof(pragma), "%s", sstr.str().c_str());
|
||||
|
||||
if (create_single_kernel_helper(context, &program, &kernel, 2,
|
||||
finalProgramSrc, "test_conversion"))
|
||||
{
|
||||
log_error( "ERROR: Explicit cast of scalar %s to vector %s%d FAILED; skipping other %s vector tests\n",
|
||||
get_explicit_type_name(srcType), get_explicit_type_name((ExplicitType)dstType), sizes[i], get_explicit_type_name((ExplicitType)dstType) );
|
||||
failed = -1;
|
||||
break;
|
||||
log_info("****** %s%s *******\n", finalProgramSrc[0],
|
||||
finalProgramSrc[1]);
|
||||
throw std::runtime_error(
|
||||
"create_single_kernel_helper failed\n");
|
||||
}
|
||||
|
||||
if (test_explicit_s2v_function(context, queue, kernel, srcT,
|
||||
sample_count, dstT, sizes[i],
|
||||
inputData)
|
||||
!= 0)
|
||||
{
|
||||
log_error("ERROR: Explicit cast of scalar %s to vector %s%d "
|
||||
"FAILED; skipping other %s vector tests\n",
|
||||
get_explicit_type_name(srcT),
|
||||
get_explicit_type_name(dstT), sizes[i],
|
||||
get_explicit_type_name(dstT));
|
||||
throw std::runtime_error("test_explicit_s2v_function failed\n");
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
return failed;
|
||||
}
|
||||
protected:
|
||||
bool fp16Support;
|
||||
bool fp64Support;
|
||||
|
||||
int test_explicit_s2v_char(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
|
||||
TypeIter it;
|
||||
unsigned int dstType, srcType;
|
||||
cl_context context;
|
||||
cl_command_queue queue;
|
||||
|
||||
std::vector<ExplicitType> vecTypes;
|
||||
|
||||
constexpr static unsigned int sample_count =
|
||||
128; // hardcoded in original test
|
||||
};
|
||||
|
||||
} // anonymous namespace
|
||||
|
||||
int test_explicit_s2v(cl_device_id deviceID, cl_context context,
|
||||
cl_command_queue queue, int num_elements)
|
||||
{
|
||||
char data[128];
|
||||
RandomSeed seed(gRandomSeed);
|
||||
|
||||
generate_random_data( kChar, 128, seed, data );
|
||||
|
||||
return test_explicit_s2v_function_set( deviceID, context, queue, kChar, 128, data );
|
||||
}
|
||||
|
||||
int test_explicit_s2v_uchar(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
|
||||
{
|
||||
unsigned char data[128];
|
||||
RandomSeed seed(gRandomSeed);
|
||||
|
||||
generate_random_data( kUChar, 128, seed, data );
|
||||
|
||||
if( test_explicit_s2v_function_set( deviceID, context, queue, kUChar, 128, data ) != 0 )
|
||||
return -1;
|
||||
if( test_explicit_s2v_function_set( deviceID, context, queue, kUnsignedChar, 128, data ) != 0 )
|
||||
return -1;
|
||||
return 0;
|
||||
}
|
||||
|
||||
int test_explicit_s2v_short(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
|
||||
{
|
||||
short data[128];
|
||||
RandomSeed seed(gRandomSeed);
|
||||
|
||||
generate_random_data( kShort, 128, seed, data );
|
||||
|
||||
if( test_explicit_s2v_function_set( deviceID, context, queue, kShort, 128, data ) != 0 )
|
||||
return -1;
|
||||
return 0;
|
||||
}
|
||||
|
||||
int test_explicit_s2v_ushort(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
|
||||
{
|
||||
unsigned short data[128];
|
||||
RandomSeed seed(gRandomSeed);
|
||||
|
||||
generate_random_data( kUShort, 128, seed, data );
|
||||
|
||||
if( test_explicit_s2v_function_set( deviceID, context, queue, kUShort, 128, data ) != 0 )
|
||||
return -1;
|
||||
if( test_explicit_s2v_function_set( deviceID, context, queue, kUnsignedShort, 128, data ) != 0 )
|
||||
return -1;
|
||||
return 0;
|
||||
}
|
||||
|
||||
int test_explicit_s2v_int(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
|
||||
{
|
||||
int data[128];
|
||||
RandomSeed seed(gRandomSeed);
|
||||
|
||||
generate_random_data( kInt, 128, seed, data );
|
||||
|
||||
if( test_explicit_s2v_function_set( deviceID, context, queue, kInt, 128, data ) != 0 )
|
||||
return -1;
|
||||
return 0;
|
||||
}
|
||||
|
||||
int test_explicit_s2v_uint(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
|
||||
{
|
||||
unsigned int data[128];
|
||||
RandomSeed seed(gRandomSeed);
|
||||
|
||||
generate_random_data( kUInt, 128, seed, data );
|
||||
|
||||
if( test_explicit_s2v_function_set( deviceID, context, queue, kUInt, 128, data ) != 0 )
|
||||
return -1;
|
||||
if( test_explicit_s2v_function_set( deviceID, context, queue, kUnsignedInt, 128, data ) != 0 )
|
||||
return -1;
|
||||
return 0;
|
||||
}
|
||||
|
||||
int test_explicit_s2v_long(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
|
||||
{
|
||||
cl_long data[128];
|
||||
RandomSeed seed(gRandomSeed);
|
||||
|
||||
generate_random_data( kLong, 128, seed, data );
|
||||
|
||||
if( test_explicit_s2v_function_set( deviceID, context, queue, kLong, 128, data ) != 0 )
|
||||
return -1;
|
||||
return 0;
|
||||
}
|
||||
|
||||
int test_explicit_s2v_ulong(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
|
||||
{
|
||||
cl_ulong data[128];
|
||||
RandomSeed seed(gRandomSeed);
|
||||
|
||||
generate_random_data( kULong, 128, seed, data );
|
||||
|
||||
if( test_explicit_s2v_function_set( deviceID, context, queue, kULong, 128, data ) != 0 )
|
||||
return -1;
|
||||
if( test_explicit_s2v_function_set( deviceID, context, queue, kUnsignedLong, 128, data ) != 0 )
|
||||
return -1;
|
||||
return 0;
|
||||
}
|
||||
|
||||
int test_explicit_s2v_float(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
|
||||
{
|
||||
float data[128];
|
||||
RandomSeed seed(gRandomSeed);
|
||||
|
||||
generate_random_data( kFloat, 128, seed, data );
|
||||
|
||||
if( test_explicit_s2v_function_set( deviceID, context, queue, kFloat, 128, data ) != 0 )
|
||||
return -1;
|
||||
return 0;
|
||||
}
|
||||
|
||||
|
||||
int test_explicit_s2v_double(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
|
||||
{
|
||||
double data[128];
|
||||
RandomSeed seed(gRandomSeed);
|
||||
|
||||
if( !is_extension_available( deviceID, "cl_khr_fp64" ) ) {
|
||||
log_info("Extension cl_khr_fp64 not supported. Skipping test.\n");
|
||||
return 0;
|
||||
try
|
||||
{
|
||||
TypesIterator(deviceID, context, queue);
|
||||
} catch (const std::runtime_error &e)
|
||||
{
|
||||
log_error("%s", e.what());
|
||||
return TEST_FAIL;
|
||||
}
|
||||
|
||||
generate_random_data( kDouble, 128, seed, data );
|
||||
|
||||
if( test_explicit_s2v_function_set( deviceID, context, queue, kDouble, 128, data ) != 0 )
|
||||
return -1;
|
||||
return 0;
|
||||
return TEST_PASS;
|
||||
}
|
||||
|
||||
|
||||
|
||||
@@ -33,7 +33,11 @@
|
||||
|
||||
#include "procs.h"
|
||||
|
||||
static const char *fp_kernel_code = R"(
|
||||
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)
|
||||
{
|
||||
@@ -42,8 +46,6 @@ __kernel void test_fp(__global TYPE *srcA, __global TYPE *srcB, __global TYPE *d
|
||||
dst[tid] = srcA[tid] OP srcB[tid];
|
||||
})";
|
||||
|
||||
extern cl_half_rounding_mode halfRoundingMode;
|
||||
|
||||
#define HFF(num) cl_half_from_float(num, halfRoundingMode)
|
||||
#define HTF(num) cl_half_to_float(num)
|
||||
|
||||
@@ -370,6 +372,8 @@ protected:
|
||||
std::map<size_t, std::string> type2name;
|
||||
};
|
||||
|
||||
} // anonymous namespace
|
||||
|
||||
int test_fpmath(cl_device_id device, cl_context context, cl_command_queue queue,
|
||||
int num_elements)
|
||||
{
|
||||
|
||||
Reference in New Issue
Block a user