mirror of
https://github.com/KhronosGroup/OpenCL-CTS.git
synced 2026-03-19 14:09:03 +00:00
d8733efc0f
The maintenance of the conformance tests is moving to Github. This commit contains all the changes that have been done in Gitlab since the first public release of the conformance tests. Signed-off-by: Kevin Petit <kevin.petit@arm.com>
365 lines
13 KiB
C++
365 lines
13 KiB
C++
//
|
|
// Copyright (c) 2017 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 "testBase.h"
|
|
#include "../../test_common/harness/conversions.h"
|
|
#include "../../test_common/harness/typeWrappers.h"
|
|
#include <math.h>
|
|
#include <float.h>
|
|
|
|
extern "C" { extern cl_uint gRandomSeed;};
|
|
|
|
#define TEST_SIZE 512
|
|
|
|
const char *equivTestKernelPattern_float =
|
|
"__kernel void sample_test(__global float%s *sourceA, __global float%s *sourceB, __global int%s *destValues, __global int%s *destValuesB)\n"
|
|
"{\n"
|
|
" int tid = get_global_id(0);\n"
|
|
" destValues[tid] = %s( sourceA[tid], sourceB[tid] );\n"
|
|
" destValuesB[tid] = sourceA[tid] %s sourceB[tid];\n"
|
|
"\n"
|
|
"}\n";
|
|
|
|
const char *equivTestKernelPatternLessGreater_float =
|
|
"__kernel void sample_test(__global float%s *sourceA, __global float%s *sourceB, __global int%s *destValues, __global int%s *destValuesB)\n"
|
|
"{\n"
|
|
" int tid = get_global_id(0);\n"
|
|
" destValues[tid] = %s( sourceA[tid], sourceB[tid] );\n"
|
|
" destValuesB[tid] = (sourceA[tid] < sourceB[tid]) | (sourceA[tid] > sourceB[tid]);\n"
|
|
"\n"
|
|
"}\n";
|
|
|
|
|
|
const char *equivTestKernelPattern_float3 =
|
|
"__kernel void sample_test(__global float%s *sourceA, __global float%s *sourceB, __global int%s *destValues, __global int%s *destValuesB)\n"
|
|
"{\n"
|
|
" int tid = get_global_id(0);\n"
|
|
" float3 sampA = vload3(tid, (__global float *)sourceA);\n"
|
|
" float3 sampB = vload3(tid, (__global float *)sourceB);\n"
|
|
" vstore3(%s( sampA, sampB ), tid, (__global int *)destValues);\n"
|
|
" vstore3(( sampA %s sampB ), tid, (__global int *)destValuesB);\n"
|
|
"\n"
|
|
"}\n";
|
|
|
|
const char *equivTestKernelPatternLessGreater_float3 =
|
|
"__kernel void sample_test(__global float%s *sourceA, __global float%s *sourceB, __global int%s *destValues, __global int%s *destValuesB)\n"
|
|
"{\n"
|
|
" int tid = get_global_id(0);\n"
|
|
" float3 sampA = vload3(tid, (__global float *)sourceA);\n"
|
|
" float3 sampB = vload3(tid, (__global float *)sourceB);\n"
|
|
" vstore3(%s( sampA, sampB ), tid, (__global int *)destValues);\n"
|
|
" vstore3(( sampA < sampB ) | (sampA > sampB), tid, (__global int *)destValuesB);\n"
|
|
"\n"
|
|
"}\n";
|
|
|
|
typedef bool (*equivVerifyFn)( float inDataA, float inDataB );
|
|
extern "C" { extern int gInfNanSupport; };
|
|
|
|
int IsFloatInfinity(float x)
|
|
{
|
|
return isinf(x);
|
|
}
|
|
|
|
int IsFloatNaN(float x)
|
|
{
|
|
return isnan(x);
|
|
}
|
|
|
|
void verify_equiv_values_float( unsigned int vecSize, float *inDataA, float *inDataB, int *outData, equivVerifyFn verifyFn )
|
|
{
|
|
unsigned int i;
|
|
int trueResult;
|
|
bool result;
|
|
|
|
trueResult = ( vecSize == 1 ) ? 1 : -1;
|
|
for( i = 0; i < vecSize; i++ )
|
|
{
|
|
result = verifyFn( inDataA[ i ], inDataB[ i ] );
|
|
outData[ i ] = result ? trueResult : 0;
|
|
}
|
|
}
|
|
|
|
void generate_equiv_test_data_float( float *outData, unsigned int vecSize, bool alpha, MTdata d )
|
|
{
|
|
unsigned int i;
|
|
|
|
generate_random_data( kFloat, vecSize * TEST_SIZE, d, outData );
|
|
|
|
// Fill the first few vectors with NAN in each vector element (or the second set if we're alpha, so we can test either case)
|
|
if( alpha )
|
|
outData += vecSize * vecSize;
|
|
for( i = 0; i < vecSize; i++ )
|
|
{
|
|
outData[ 0 ] = NAN;
|
|
outData += vecSize + 1;
|
|
}
|
|
// Make sure the third set is filled regardless, to test the case where both have NANs
|
|
if( !alpha )
|
|
outData += vecSize * vecSize;
|
|
for( i = 0; i < vecSize; i++ )
|
|
{
|
|
outData[ 0 ] = NAN;
|
|
outData += vecSize + 1;
|
|
}
|
|
}
|
|
|
|
int test_equiv_kernel_float(cl_context context, cl_command_queue queue, const char *fnName, const char *opName,
|
|
unsigned int vecSize, equivVerifyFn verifyFn, MTdata d )
|
|
{
|
|
clProgramWrapper program;
|
|
clKernelWrapper kernel;
|
|
clMemWrapper streams[4];
|
|
float inDataA[TEST_SIZE * 16], inDataB[ TEST_SIZE * 16 ];
|
|
int outData[TEST_SIZE * 16], expected[16];
|
|
int error, i, j;
|
|
size_t threads[1], localThreads[1];
|
|
char kernelSource[10240];
|
|
char *programPtr;
|
|
char sizeName[4];
|
|
|
|
|
|
/* Create the source */
|
|
if( vecSize == 1 )
|
|
sizeName[ 0 ] = 0;
|
|
else
|
|
sprintf( sizeName, "%d", vecSize );
|
|
|
|
|
|
if(DENSE_PACK_VECS && vecSize == 3) {
|
|
if (strcmp(fnName, "islessgreater")) {
|
|
sprintf( kernelSource, equivTestKernelPattern_float3, sizeName, sizeName, sizeName, sizeName, fnName, opName );
|
|
} else {
|
|
sprintf( kernelSource, equivTestKernelPatternLessGreater_float3, sizeName, sizeName, sizeName, sizeName, fnName );
|
|
}
|
|
} else {
|
|
if (strcmp(fnName, "islessgreater")) {
|
|
sprintf( kernelSource, equivTestKernelPattern_float, sizeName, sizeName, sizeName, sizeName, fnName, opName );
|
|
} else {
|
|
sprintf( kernelSource, equivTestKernelPatternLessGreater_float, sizeName, sizeName, sizeName, sizeName, fnName );
|
|
}
|
|
}
|
|
|
|
/* Create kernels */
|
|
programPtr = kernelSource;
|
|
if( create_single_kernel_helper( context, &program, &kernel, 1, (const char **)&programPtr, "sample_test" ) )
|
|
{
|
|
return -1;
|
|
}
|
|
|
|
/* Generate some streams */
|
|
generate_equiv_test_data_float( inDataA, vecSize, true, d );
|
|
generate_equiv_test_data_float( inDataB, vecSize, false, d );
|
|
|
|
streams[0] = clCreateBuffer( context, (cl_mem_flags)(CL_MEM_COPY_HOST_PTR), sizeof( cl_float ) * vecSize * TEST_SIZE, &inDataA, &error);
|
|
if( streams[0] == NULL )
|
|
{
|
|
print_error( error, "Creating input array A failed!\n");
|
|
return -1;
|
|
}
|
|
streams[1] = clCreateBuffer( context, (cl_mem_flags)(CL_MEM_COPY_HOST_PTR), sizeof( cl_float ) * vecSize * TEST_SIZE, &inDataB, &error);
|
|
if( streams[1] == NULL )
|
|
{
|
|
print_error( error, "Creating input array A failed!\n");
|
|
return -1;
|
|
}
|
|
streams[2] = clCreateBuffer( context, CL_MEM_READ_WRITE, sizeof( cl_int ) * vecSize * TEST_SIZE, NULL, &error);
|
|
if( streams[2] == NULL )
|
|
{
|
|
print_error( error, "Creating output array failed!\n");
|
|
return -1;
|
|
}
|
|
streams[3] = clCreateBuffer( context, CL_MEM_READ_WRITE, sizeof( cl_int ) * vecSize * TEST_SIZE, NULL, &error);
|
|
if( streams[3] == NULL )
|
|
{
|
|
print_error( error, "Creating output array failed!\n");
|
|
return -1;
|
|
}
|
|
|
|
|
|
/* Assign streams and execute */
|
|
error = clSetKernelArg( kernel, 0, sizeof( streams[0] ), &streams[0] );
|
|
test_error( error, "Unable to set indexed kernel arguments" );
|
|
error = clSetKernelArg( kernel, 1, sizeof( streams[1] ), &streams[1] );
|
|
test_error( error, "Unable to set indexed kernel arguments" );
|
|
error = clSetKernelArg( kernel, 2, sizeof( streams[2] ), &streams[2] );
|
|
test_error( error, "Unable to set indexed kernel arguments" );
|
|
error = clSetKernelArg( kernel, 3, sizeof( streams[3] ), &streams[3] );
|
|
test_error( error, "Unable to set indexed kernel arguments" );
|
|
|
|
|
|
/* Run the kernel */
|
|
threads[0] = TEST_SIZE;
|
|
|
|
error = get_max_common_work_group_size( context, kernel, threads[0], &localThreads[0] );
|
|
test_error( error, "Unable to get work group size to use" );
|
|
|
|
error = clEnqueueNDRangeKernel( queue, kernel, 1, NULL, threads, localThreads, 0, NULL, NULL );
|
|
test_error( error, "Unable to execute test kernel" );
|
|
|
|
/* Now get the results */
|
|
error = clEnqueueReadBuffer( queue, streams[2], true, 0, sizeof( int ) * TEST_SIZE * vecSize, outData, 0, NULL, NULL );
|
|
test_error( error, "Unable to read output array!" );
|
|
|
|
/* And verify! */
|
|
for( i = 0; i < TEST_SIZE; i++ )
|
|
{
|
|
verify_equiv_values_float( vecSize, &inDataA[ i * vecSize ], &inDataB[ i * vecSize ], expected, verifyFn);
|
|
|
|
for( j = 0; j < (int)vecSize; j++ )
|
|
{
|
|
if( expected[ j ] != outData[ i * vecSize + j ] )
|
|
{
|
|
log_error( "ERROR: Data sample %d:%d at size %d does not validate! Expected %d, got %d, source %f,%f\n",
|
|
i, j, vecSize, expected[ j ], outData[ i * vecSize + j ], inDataA[i*vecSize + j], inDataB[i*vecSize + j] );
|
|
return -1;
|
|
}
|
|
}
|
|
}
|
|
|
|
/* Now get the results */
|
|
error = clEnqueueReadBuffer( queue, streams[3], true, 0, sizeof( int ) * TEST_SIZE * vecSize, outData, 0, NULL, NULL );
|
|
test_error( error, "Unable to read output array!" );
|
|
|
|
/* And verify! */
|
|
int fail = 0;
|
|
for( i = 0; i < TEST_SIZE; i++ )
|
|
{
|
|
verify_equiv_values_float( vecSize, &inDataA[ i * vecSize ], &inDataB[ i * vecSize ], expected, verifyFn);
|
|
|
|
for( j = 0; j < (int)vecSize; j++ )
|
|
{
|
|
if( expected[ j ] != outData[ i * vecSize + j ] )
|
|
{
|
|
if (gInfNanSupport == 0)
|
|
{
|
|
if (IsFloatNaN(inDataA[i*vecSize + j]) || IsFloatNaN (inDataB[i*vecSize + j]))
|
|
{
|
|
fail = 0;
|
|
}
|
|
else
|
|
fail = 1;
|
|
}
|
|
if (fail)
|
|
{
|
|
log_error( "ERROR: Data sample %d:%d at size %d does not validate! Expected %d, got %d, source %f,%f\n",
|
|
i, j, vecSize, expected[ j ], outData[ i * vecSize + j ], inDataA[i*vecSize + j], inDataB[i*vecSize + j] );
|
|
return -1;
|
|
}
|
|
}
|
|
}
|
|
}
|
|
|
|
return 0;
|
|
}
|
|
|
|
int test_equiv_kernel_set_float(cl_context context, cl_command_queue queue, const char *fnName, const char *opName, equivVerifyFn verifyFn, MTdata d )
|
|
{
|
|
unsigned int vecSizes[] = { 1, 2, 3, 4, 8, 16, 0 };
|
|
unsigned int index;
|
|
int retVal = 0;
|
|
|
|
|
|
for( index = 0; vecSizes[ index ] != 0; index++ )
|
|
{
|
|
// Test!
|
|
if( test_equiv_kernel_float(context, queue, fnName, opName, vecSizes[ index ], verifyFn, d ) != 0 )
|
|
{
|
|
log_error( " Vector float%d FAILED\n", vecSizes[ index ] );
|
|
retVal = -1;
|
|
}
|
|
}
|
|
|
|
return retVal;
|
|
}
|
|
|
|
bool isequal_verify_fn_float( float valueA, float valueB )
|
|
{
|
|
return valueA == valueB;
|
|
}
|
|
|
|
int test_relational_isequal_float(cl_device_id device, cl_context context, cl_command_queue queue, int numElements )
|
|
{
|
|
RandomSeed seed( gRandomSeed );
|
|
return test_equiv_kernel_set_float( context, queue, "isequal", "==", isequal_verify_fn_float, seed );
|
|
}
|
|
|
|
bool isnotequal_verify_fn_float( float valueA, float valueB )
|
|
{
|
|
return valueA != valueB;
|
|
}
|
|
|
|
int test_relational_isnotequal_float(cl_device_id device, cl_context context, cl_command_queue queue, int numElements )
|
|
{
|
|
RandomSeed seed( gRandomSeed );
|
|
return test_equiv_kernel_set_float( context, queue, "isnotequal", "!=", isnotequal_verify_fn_float, seed );
|
|
}
|
|
|
|
bool isgreater_verify_fn_float( float valueA, float valueB )
|
|
{
|
|
return valueA > valueB;
|
|
}
|
|
|
|
int test_relational_isgreater_float(cl_device_id device, cl_context context, cl_command_queue queue, int numElements )
|
|
{
|
|
RandomSeed seed( gRandomSeed );
|
|
return test_equiv_kernel_set_float( context, queue, "isgreater", ">", isgreater_verify_fn_float, seed );
|
|
}
|
|
|
|
bool isgreaterequal_verify_fn_float( float valueA, float valueB )
|
|
{
|
|
return valueA >= valueB;
|
|
}
|
|
|
|
int test_relational_isgreaterequal_float(cl_device_id device, cl_context context, cl_command_queue queue, int numElements )
|
|
{
|
|
RandomSeed seed( gRandomSeed );
|
|
return test_equiv_kernel_set_float( context, queue, "isgreaterequal", ">=", isgreaterequal_verify_fn_float, seed );
|
|
}
|
|
|
|
bool isless_verify_fn_float( float valueA, float valueB )
|
|
{
|
|
return valueA < valueB;
|
|
}
|
|
|
|
int test_relational_isless_float(cl_device_id device, cl_context context, cl_command_queue queue, int numElements )
|
|
{
|
|
RandomSeed seed( gRandomSeed );
|
|
return test_equiv_kernel_set_float( context, queue, "isless", "<", isless_verify_fn_float, seed );
|
|
}
|
|
|
|
bool islessequal_verify_fn_float( float valueA, float valueB )
|
|
{
|
|
return valueA <= valueB;
|
|
}
|
|
|
|
int test_relational_islessequal_float(cl_device_id device, cl_context context, cl_command_queue queue, int numElements )
|
|
{
|
|
RandomSeed seed( gRandomSeed );
|
|
return test_equiv_kernel_set_float( context, queue, "islessequal", "<=", islessequal_verify_fn_float, seed );
|
|
}
|
|
|
|
bool islessgreater_verify_fn_float( float valueA, float valueB )
|
|
{
|
|
return ( valueA < valueB ) || ( valueA > valueB );
|
|
}
|
|
|
|
int test_relational_islessgreater_float(cl_device_id device, cl_context context, cl_command_queue queue, int numElements )
|
|
{
|
|
RandomSeed seed( gRandomSeed );
|
|
return test_equiv_kernel_set_float( context, queue, "islessgreater", "<>", islessgreater_verify_fn_float, seed );
|
|
}
|
|
|
|
|