mirror of
https://github.com/KhronosGroup/OpenCL-CTS.git
synced 2026-03-19 06:09:01 +00:00
Rename test .c sources to .cpp where necessary (#604)
Remove hacks to force language from CMake files. Closes KhronosGroup/OpenCL-CTS#25
This commit is contained in:
James Price
285
test_conformance/basic/test_image_param.cpp
Normal file
285
test_conformance/basic/test_image_param.cpp
Normal file
@@ -0,0 +1,285 @@
|
||||
//
|
||||
// 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 "harness/compat.h"
|
||||
|
||||
#include <stdio.h>
|
||||
#include <stdlib.h>
|
||||
#include <string.h>
|
||||
#include <sys/types.h>
|
||||
#include <sys/stat.h>
|
||||
|
||||
|
||||
#include "procs.h"
|
||||
#include "harness/typeWrappers.h"
|
||||
#include "harness/imageHelpers.h"
|
||||
#include "harness/conversions.h"
|
||||
|
||||
|
||||
static const char *param_kernel[] = {
|
||||
"__kernel void test_fn(read_only image2d_t srcimg, sampler_t sampler, __global float4 *results )\n"
|
||||
"{\n"
|
||||
" int tid_x = get_global_id(0);\n"
|
||||
" int tid_y = get_global_id(1);\n"
|
||||
" results[ tid_y * get_image_width( srcimg ) + tid_x ] = read_imagef(srcimg, sampler, (int2)(tid_x, tid_y));\n"
|
||||
"\n"
|
||||
"}\n" };
|
||||
|
||||
int validate_results( size_t width, size_t height, cl_image_format &format, char *inputData, cl_float *actualResults )
|
||||
{
|
||||
for( size_t i = 0; i < width * height; i++ )
|
||||
{
|
||||
cl_float expected[ 4 ], tolerance;
|
||||
|
||||
switch( format.image_channel_data_type )
|
||||
{
|
||||
case CL_UNORM_INT8:
|
||||
{
|
||||
cl_uchar *p = (cl_uchar *)inputData;
|
||||
expected[ 0 ] = p[ 0 ] / 255.f;
|
||||
expected[ 1 ] = p[ 1 ] / 255.f;
|
||||
expected[ 2 ] = p[ 2 ] / 255.f;
|
||||
expected[ 3 ] = p[ 3 ] / 255.f;
|
||||
tolerance = 1.f / 255.f;
|
||||
break;
|
||||
}
|
||||
case CL_SNORM_INT8:
|
||||
{
|
||||
cl_char *p = (cl_char *)inputData;
|
||||
expected[ 0 ] = fmaxf( p[ 0 ] / 127.f, -1.f );
|
||||
expected[ 1 ] = fmaxf( p[ 1 ] / 127.f, -1.f );
|
||||
expected[ 2 ] = fmaxf( p[ 2 ] / 127.f, -1.f );
|
||||
expected[ 3 ] = fmaxf( p[ 3 ] / 127.f, -1.f );
|
||||
tolerance = 1.f / 127.f;
|
||||
break;
|
||||
}
|
||||
case CL_UNSIGNED_INT8:
|
||||
{
|
||||
cl_uchar *p = (cl_uchar *)inputData;
|
||||
expected[ 0 ] = p[ 0 ];
|
||||
expected[ 1 ] = p[ 1 ];
|
||||
expected[ 2 ] = p[ 2 ];
|
||||
expected[ 3 ] = p[ 3 ];
|
||||
tolerance = 1.f / 127.f;
|
||||
break;
|
||||
}
|
||||
case CL_SIGNED_INT8:
|
||||
{
|
||||
cl_short *p = (cl_short *)inputData;
|
||||
expected[ 0 ] = p[ 0 ];
|
||||
expected[ 1 ] = p[ 1 ];
|
||||
expected[ 2 ] = p[ 2 ];
|
||||
expected[ 3 ] = p[ 3 ];
|
||||
tolerance = 1.f / 127.f;
|
||||
break;
|
||||
}
|
||||
case CL_UNORM_INT16:
|
||||
{
|
||||
cl_ushort *p = (cl_ushort *)inputData;
|
||||
expected[ 0 ] = p[ 0 ] / 65535.f;
|
||||
expected[ 1 ] = p[ 1 ] / 65535.f;
|
||||
expected[ 2 ] = p[ 2 ] / 65535.f;
|
||||
expected[ 3 ] = p[ 3 ] / 65535.f;
|
||||
tolerance = 1.f / 65535.f;
|
||||
break;
|
||||
}
|
||||
case CL_UNSIGNED_INT32:
|
||||
{
|
||||
cl_uint *p = (cl_uint *)inputData;
|
||||
expected[ 0 ] = p[ 0 ];
|
||||
expected[ 1 ] = p[ 1 ];
|
||||
expected[ 2 ] = p[ 2 ];
|
||||
expected[ 3 ] = p[ 3 ];
|
||||
tolerance = 0.0001f;
|
||||
break;
|
||||
}
|
||||
case CL_FLOAT:
|
||||
{
|
||||
cl_float *p = (cl_float *)inputData;
|
||||
expected[ 0 ] = p[ 0 ];
|
||||
expected[ 1 ] = p[ 1 ];
|
||||
expected[ 2 ] = p[ 2 ];
|
||||
expected[ 3 ] = p[ 3 ];
|
||||
tolerance = 0.0001f;
|
||||
break;
|
||||
}
|
||||
default:
|
||||
// Should never get here
|
||||
break;
|
||||
}
|
||||
|
||||
if( format.image_channel_order == CL_BGRA )
|
||||
{
|
||||
cl_float tmp = expected[ 0 ];
|
||||
expected[ 0 ] = expected[ 2 ];
|
||||
expected[ 2 ] = tmp;
|
||||
}
|
||||
|
||||
// Within an error tolerance, make sure the results match
|
||||
cl_float error1 = fabsf( expected[ 0 ] - actualResults[ 0 ] );
|
||||
cl_float error2 = fabsf( expected[ 1 ] - actualResults[ 1 ] );
|
||||
cl_float error3 = fabsf( expected[ 2 ] - actualResults[ 2 ] );
|
||||
cl_float error4 = fabsf( expected[ 3 ] - actualResults[ 3 ] );
|
||||
|
||||
if( error1 > tolerance || error2 > tolerance || error3 > tolerance || error4 > tolerance )
|
||||
{
|
||||
log_error( "ERROR: Sample %d did not validate against expected results for %d x %d %s:%s image\n", (int)i, (int)width, (int)height,
|
||||
GetChannelOrderName( format.image_channel_order ), GetChannelTypeName( format.image_channel_data_type ) );
|
||||
log_error( " Expected: %f %f %f %f\n", (float)expected[ 0 ], (float)expected[ 1 ], (float)expected[ 2 ], (float)expected[ 3 ] );
|
||||
log_error( " Actual: %f %f %f %f\n", (float)actualResults[ 0 ], (float)actualResults[ 1 ], (float)actualResults[ 2 ], (float)actualResults[ 3 ] );
|
||||
|
||||
// Check real quick a special case error here
|
||||
cl_float error1 = fabsf( expected[ 3 ] - actualResults[ 0 ] );
|
||||
cl_float error2 = fabsf( expected[ 2 ] - actualResults[ 1 ] );
|
||||
cl_float error3 = fabsf( expected[ 1 ] - actualResults[ 2 ] );
|
||||
cl_float error4 = fabsf( expected[ 0 ] - actualResults[ 3 ] );
|
||||
if( error1 <= tolerance && error2 <= tolerance && error3 <= tolerance && error4 <= tolerance )
|
||||
{
|
||||
log_error( "\t(Kernel did not respect change in channel order)\n" );
|
||||
}
|
||||
return -1;
|
||||
}
|
||||
|
||||
// Increment and go
|
||||
actualResults += 4;
|
||||
inputData += get_format_type_size( &format ) * 4;
|
||||
}
|
||||
|
||||
return 0;
|
||||
}
|
||||
|
||||
int test_image_param(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
|
||||
{
|
||||
size_t sizes[] = { 64, 100, 128, 250, 512 };
|
||||
cl_image_format formats[] = { { CL_RGBA, CL_UNORM_INT8 }, { CL_RGBA, CL_UNORM_INT16 }, { CL_RGBA, CL_FLOAT }, { CL_BGRA, CL_UNORM_INT8 } };
|
||||
cl_image_format *supported_formats;
|
||||
ExplicitType types[] = { kUChar, kUShort, kFloat, kUChar };
|
||||
int error;
|
||||
size_t i, j, idx;
|
||||
size_t threads[ 2 ];
|
||||
MTdata d;
|
||||
int supportsBGRA = 0;
|
||||
cl_uint numSupportedFormats = 0;
|
||||
|
||||
const size_t numSizes = sizeof( sizes ) / sizeof( sizes[ 0 ] );
|
||||
const size_t numFormats = sizeof( formats ) / sizeof( formats[ 0 ] );
|
||||
const size_t numAttempts = numSizes * numFormats;
|
||||
|
||||
|
||||
clProgramWrapper program;
|
||||
clKernelWrapper kernel;
|
||||
clMemWrapper streams[ numAttempts ][ 2 ];
|
||||
BufferOwningPtr<char> inputs[ numAttempts ];
|
||||
|
||||
PASSIVE_REQUIRE_IMAGE_SUPPORT( device )
|
||||
|
||||
if(gIsEmbedded)
|
||||
{
|
||||
/* Get the supported image formats to see if BGRA is supported */
|
||||
clGetSupportedImageFormats (context, CL_MEM_READ_WRITE, CL_MEM_OBJECT_IMAGE2D, 0, NULL, &numSupportedFormats);
|
||||
supported_formats = (cl_image_format *) malloc(sizeof(cl_image_format) * numSupportedFormats);
|
||||
clGetSupportedImageFormats (context, CL_MEM_READ_WRITE, CL_MEM_OBJECT_IMAGE2D, numFormats, supported_formats, NULL);
|
||||
|
||||
for(i = 0; i < numSupportedFormats; i++)
|
||||
{
|
||||
if(supported_formats[i].image_channel_order == CL_BGRA)
|
||||
{
|
||||
supportsBGRA = 1;
|
||||
break;
|
||||
}
|
||||
}
|
||||
}
|
||||
else
|
||||
{
|
||||
supportsBGRA = 1;
|
||||
}
|
||||
|
||||
d = init_genrand( gRandomSeed );
|
||||
for( i = 0, idx = 0; i < numSizes; i++ )
|
||||
{
|
||||
for( j = 0; j < numFormats; j++, idx++ )
|
||||
{
|
||||
if(formats[j].image_channel_order == CL_BGRA && !supportsBGRA)
|
||||
continue;
|
||||
|
||||
// For each attempt, we create a pair: an input image, whose parameters keep changing, and an output buffer
|
||||
// that we can read values from. The output buffer will remain consistent to ensure that any changes we
|
||||
// witness are due to the image changes
|
||||
inputs[ idx ].reset(create_random_data( types[ j ], d, sizes[ i ] * sizes[ i ] * 4 ));
|
||||
|
||||
streams[ idx ][ 0 ] = create_image_2d( context, CL_MEM_COPY_HOST_PTR, &formats[ j ], sizes[ i ], sizes[ i ], 0, inputs[ idx ], &error );
|
||||
{
|
||||
char err_str[256];
|
||||
sprintf(err_str, "Unable to create input image for format %s order %s" ,
|
||||
GetChannelOrderName( formats[j].image_channel_order ),
|
||||
GetChannelTypeName( formats[j].image_channel_data_type ));
|
||||
test_error( error, err_str);
|
||||
}
|
||||
|
||||
streams[ idx ][ 1 ] = clCreateBuffer( context, CL_MEM_READ_WRITE, sizes[ i ] * sizes[ i ] * 4 * sizeof( cl_float ), NULL, &error );
|
||||
test_error( error, "Unable to create output buffer" );
|
||||
}
|
||||
}
|
||||
free_mtdata(d); d = NULL;
|
||||
|
||||
// Create a single kernel to use for all the tests
|
||||
error = create_single_kernel_helper( context, &program, &kernel, 1, param_kernel, "test_fn" );
|
||||
test_error( error, "Unable to create testing kernel" );
|
||||
|
||||
// Also create a sampler to use for all the runs
|
||||
clSamplerWrapper sampler = clCreateSampler(context, CL_FALSE, CL_ADDRESS_CLAMP_TO_EDGE, CL_FILTER_NEAREST, &error );
|
||||
test_error( error, "clCreateSampler failed" );
|
||||
|
||||
// Set up the arguments for each and queue
|
||||
for( i = 0, idx = 0; i < numSizes; i++ )
|
||||
{
|
||||
for( j = 0; j < numFormats; j++, idx++ )
|
||||
{
|
||||
if(formats[j].image_channel_order == CL_BGRA && !supportsBGRA)
|
||||
continue;
|
||||
|
||||
error = clSetKernelArg( kernel, 0, sizeof( streams[ idx ][ 0 ] ), &streams[ idx ][ 0 ] );
|
||||
error |= clSetKernelArg( kernel, 1, sizeof( sampler ), &sampler );
|
||||
error |= clSetKernelArg( kernel, 2, sizeof( streams[ idx ][ 1 ] ), &streams[ idx ][ 1 ]);
|
||||
test_error( error, "Unable to set kernel arguments" );
|
||||
|
||||
threads[ 0 ] = threads[ 1 ] = (size_t)sizes[ i ];
|
||||
|
||||
error = clEnqueueNDRangeKernel( queue, kernel, 2, NULL, threads, NULL, 0, NULL, NULL );
|
||||
test_error( error, "clEnqueueNDRangeKernel failed" );
|
||||
}
|
||||
}
|
||||
|
||||
// Now go through each combo and validate the results
|
||||
for( i = 0, idx = 0; i < numSizes; i++ )
|
||||
{
|
||||
for( j = 0; j < numFormats; j++, idx++ )
|
||||
{
|
||||
if(formats[j].image_channel_order == CL_BGRA && !supportsBGRA)
|
||||
continue;
|
||||
|
||||
BufferOwningPtr<cl_float> output(malloc(sizeof(cl_float) * sizes[ i ] * sizes[ i ] * 4 ));
|
||||
|
||||
error = clEnqueueReadBuffer( queue, streams[ idx ][ 1 ], CL_TRUE, 0, sizes[ i ] * sizes[ i ] * 4 * sizeof( cl_float ), output, 0, NULL, NULL );
|
||||
test_error( error, "Unable to read results" );
|
||||
|
||||
error = validate_results( sizes[ i ], sizes[ i ], formats[ j ], inputs[ idx ], output );
|
||||
if( error )
|
||||
return -1;
|
||||
}
|
||||
}
|
||||
|
||||
return 0;
|
||||
}
|
||||
Reference in New Issue
Block a user