blob: 39b54b4f173208ed3e07e75b583b2ba0b95b8696 [file] [log] [blame]
//
// 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 "procs.h"
#include <ctype.h>
const char *work_offset_test[] = {
"__kernel void test( __global int * outputID_A, \n"
" __global int * outputID_B, __global int * outputID_C )\n"
"{\n"
" size_t id0 = get_local_id( 0 ) + get_group_id( 0 ) * get_local_size( 0 );\n"
" size_t id1 = get_local_id( 1 ) + get_group_id( 1 ) * get_local_size( 1 );\n"
" size_t id2 = get_local_id( 2 ) + get_group_id( 2 ) * get_local_size( 2 );\n"
" size_t id = ( id2 * get_global_size( 0 ) * get_global_size( 1 ) ) + ( id1 * get_global_size( 0 ) ) + id0;\n"
"\n"
" outputID_A[ id ] = get_global_id( 0 );\n"
" outputID_B[ id ] = get_global_id( 1 );\n"
" outputID_C[ id ] = get_global_id( 2 );\n"
"}\n"
};
#define MAX_TEST_ITEMS 16 * 16 * 16
#define NUM_TESTS 16
#define MAX_OFFSET 256
#define CHECK_RANGE( v, m, c ) \
if( ( v >= (cl_int)m ) || ( v < 0 ) ) \
{ \
log_error( "ERROR: ouputID_%c[%lu]: %d is < 0 or >= %lu\n", c, i, v, m ); \
return -1; \
}
int check_results( size_t threads[], size_t offsets[], cl_int outputA[], cl_int outputB[], cl_int outputC[] )
{
size_t offsettedSizes[ 3 ] = { threads[ 0 ] + offsets[ 0 ], threads[ 1 ] + offsets[ 1 ], threads[ 2 ] + offsets[ 2 ] };
size_t limit = threads[ 0 ] * threads[ 1 ] * threads[ 2 ];
static char counts[ MAX_OFFSET + 32 ][ MAX_OFFSET + 16 ][ MAX_OFFSET + 16 ];
memset( counts, 0, sizeof( counts ) );
for( size_t i = 0; i < limit; i++ )
{
// Check ranges first
CHECK_RANGE( outputA[ i ], offsettedSizes[ 0 ], 'A' )
CHECK_RANGE( outputB[ i ], offsettedSizes[ 1 ], 'B' )
CHECK_RANGE( outputC[ i ], offsettedSizes[ 2 ], 'C' )
// Now set the value in the map
counts[ outputA[ i ] ][ outputB[ i ] ][ outputC[ i ] ]++;
}
// Now check the map
int missed = 0, multiple = 0, errored = 0, corrected = 0;
for( size_t x = 0; x < offsettedSizes[ 0 ]; x++ )
{
for( size_t y = 0; y < offsettedSizes[ 1 ]; y++ )
{
for( size_t z = 0; z < offsettedSizes[ 2 ]; z++ )
{
const char * limitMsg = " (further errors of this type suppressed)";
if( ( x >= offsets[ 0 ] ) && ( y >= offsets[ 1 ] ) && ( z >= offsets[ 2 ] ) )
{
if( counts[ x ][ y ][ z ] < 1 )
{
if( missed < 3 )
log_error( "ERROR: Map value (%ld,%ld,%ld) was missed%s\n", x, y, z, ( missed == 2 ) ? limitMsg : "" );
missed++;
}
else if( counts[ x ][ y ][ z ] > 1 )
{
if( multiple < 3 )
log_error( "ERROR: Map value (%ld,%ld,%ld) was returned multiple times%s\n", x, y, z, ( multiple == 2 ) ? limitMsg : "" );
multiple++;
}
}
else
{
if( counts[ x ][ y ][ z ] > 0 )
{
if( errored < 3 )
log_error( "ERROR: Map value (%ld,%ld,%ld) was erroneously returned%s\n", x, y, z, ( errored == 2 ) ? limitMsg : "" );
errored++;
}
}
}
}
}
if( missed || multiple || errored )
{
size_t diffs[3] = { ( offsets[ 0 ] > threads[ 0 ] ? 0 : threads[ 0 ] - offsets[ 0 ] ),
( offsets[ 1 ] > threads[ 1 ] ? 0 : threads[ 1 ] - offsets[ 1 ] ),
( offsets[ 2 ] > threads[ 2 ] ? 0 : threads[ 2 ] - offsets[ 2 ] ) };
int diff = (int)( ( threads[ 0 ] - diffs[ 0 ] ) * ( threads[ 1 ] - diffs[ 1 ] ) * ( threads[ 2 ] - diffs[ 2 ] ) );
if( ( multiple == 0 ) && ( missed == diff ) && ( errored == diff ) )
log_error( "ERROR: Global work offset values are not being respected by get_global_id()\n" );
else
log_error( "ERROR: Global work offset values did not function as expected (%d missed, %d reported multiple times, %d erroneously hit)\n",
missed, multiple, errored );
}
return ( missed | multiple | errored | corrected );
}
int test_global_work_offsets(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
{
clProgramWrapper program;
clKernelWrapper kernel;
clMemWrapper streams[ 7 ];
int error;
size_t threads[] = {1,1,1}, localThreads[] = {1,1,1}, offsets[] = {0,0,0};
cl_int outputA[ MAX_TEST_ITEMS ], outputB[ MAX_TEST_ITEMS ], outputC[ MAX_TEST_ITEMS ];
// Create the kernel
if( create_single_kernel_helper( context, &program, &kernel, 1, work_offset_test, "test" ) != 0 )
{
return -1;
}
//// Create some output streams
// Use just one output array to init them all (no need to init every single stack storage here)
memset( outputA, 0xff, sizeof( outputA ) );
for( int i = 0; i < 3; i++ )
{
streams[i] =
clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR,
sizeof(outputA), outputA, &error);
test_error( error, "Unable to create output array" );
}
// Run a few different times
MTdata seed = init_genrand( gRandomSeed );
for( int test = 0; test < NUM_TESTS; test++ )
{
// Choose a random combination of thread size, but in total less than MAX_TEST_ITEMS
threads[ 0 ] = random_in_range( 1, 32, seed );
threads[ 1 ] = random_in_range( 1, 16, seed );
threads[ 2 ] = random_in_range( 1, MAX_TEST_ITEMS / (int)( threads[ 0 ] * threads[ 1 ] ), seed );
// Make sure we get the local thread count right
error = get_max_common_3D_work_group_size( context, kernel, threads, localThreads );
test_error( error, "Unable to determine local work group sizes" );
// Randomize some offsets
for( int j = 0; j < 3; j++ )
offsets[ j ] = random_in_range( 0, MAX_OFFSET, seed );
log_info( "\tTesting %ld,%ld,%ld (%ld,%ld,%ld) with offsets (%ld,%ld,%ld)...\n",
threads[ 0 ], threads[ 1 ], threads[ 2 ], localThreads[ 0 ], localThreads[ 1 ], localThreads[ 2 ],
offsets[ 0 ], offsets[ 1 ], offsets[ 2 ] );
// Now set up and run
for( int i = 0; i < 3; i++ )
{
error = clSetKernelArg( kernel, i, sizeof( streams[i] ), &streams[i] );
test_error( error, "Unable to set indexed kernel arguments" );
}
error = clEnqueueNDRangeKernel( queue, kernel, 3, offsets, threads, localThreads, 0, NULL, NULL );
test_error( error, "Kernel execution failed" );
// Read our results back now
cl_int * resultBuffers[] = { outputA, outputB, outputC };
for( int i = 0; i < 3; i++ )
{
error = clEnqueueReadBuffer( queue, streams[ i ], CL_TRUE, 0, sizeof( outputA ), resultBuffers[ i ], 0, NULL, NULL );
test_error( error, "Unable to get result data" );
}
// Now we need to check the results. The outputs should have one entry for each possible ID,
// but they won't be in order, so we need to construct a count map to determine what we got
if( check_results( threads, offsets, outputA, outputB, outputC ) )
{
log_error( "\t(Test failed for global dim %ld,%ld,%ld, local dim %ld,%ld,%ld, offsets %ld,%ld,%ld)\n",
threads[ 0 ], threads[ 1 ], threads[ 2 ], localThreads[ 0 ], localThreads[ 1 ], localThreads[ 2 ],
offsets[ 0 ], offsets[ 1 ], offsets[ 2 ] );
return -1;
}
}
free_mtdata(seed);
// All done!
return 0;
}
const char *get_offset_test[] = {
"__kernel void test( __global int * outOffsets )\n"
"{\n"
" // We use local ID here so we don't have to worry about offsets\n"
" // Also note that these should be the same for ALL threads, so we won't worry about contention\n"
" outOffsets[ 0 ] = (int)get_global_offset( 0 );\n"
" outOffsets[ 1 ] = (int)get_global_offset( 1 );\n"
" outOffsets[ 2 ] = (int)get_global_offset( 2 );\n"
"}\n"
};
int test_get_global_offset(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
{
clProgramWrapper program;
clKernelWrapper kernel;
clMemWrapper streams[ 1 ];
int error;
size_t threads[] = {1,1,1}, localThreads[] = {1,1,1}, offsets[] = {0,0,0};
cl_int outOffsets[ 3 ];
// Create the kernel
if( create_single_kernel_helper( context, &program, &kernel, 1, get_offset_test, "test" ) != 0 )
{
return -1;
}
// Create some output streams, and storage for a single control ID
memset( outOffsets, 0xff, sizeof( outOffsets ) );
streams[0] =
clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR,
sizeof(outOffsets), outOffsets, &error);
test_error( error, "Unable to create control ID buffer" );
// Run a few different times
MTdata seed = init_genrand( gRandomSeed );
for( int test = 0; test < NUM_TESTS; test++ )
{
// Choose a random combination of thread size, but in total less than MAX_TEST_ITEMS
threads[ 0 ] = random_in_range( 1, 32, seed );
threads[ 1 ] = random_in_range( 1, 16, seed );
threads[ 2 ] = random_in_range( 1, MAX_TEST_ITEMS / (int)( threads[ 0 ] * threads[ 1 ] ), seed );
// Make sure we get the local thread count right
error = get_max_common_3D_work_group_size( context, kernel, threads, localThreads );
test_error( error, "Unable to determine local work group sizes" );
// Randomize some offsets
for( int j = 0; j < 3; j++ )
offsets[ j ] = random_in_range( 0, MAX_OFFSET, seed );
log_info( "\tTesting %ld,%ld,%ld (%ld,%ld,%ld) with offsets (%ld,%ld,%ld)...\n",
threads[ 0 ], threads[ 1 ], threads[ 2 ], localThreads[ 0 ], localThreads[ 1 ], localThreads[ 2 ],
offsets[ 0 ], offsets[ 1 ], offsets[ 2 ] );
// Now set up and run
error = clSetKernelArg( kernel, 0, sizeof( streams[0] ), &streams[0] );
test_error( error, "Unable to set indexed kernel arguments" );
error = clEnqueueNDRangeKernel( queue, kernel, 3, offsets, threads, localThreads, 0, NULL, NULL );
test_error( error, "Kernel execution failed" );
// Read our results back now
error = clEnqueueReadBuffer( queue, streams[ 0 ], CL_TRUE, 0, sizeof( outOffsets ), outOffsets, 0, NULL, NULL );
test_error( error, "Unable to get result data" );
// And check!
int errors = 0;
for( int j = 0; j < 3; j++ )
{
if( outOffsets[ j ] != (cl_int)offsets[ j ] )
{
log_error( "ERROR: get_global_offset( %d ) did not return expected value (expected %ld, got %d)\n", j, offsets[ j ], outOffsets[ j ] );
errors++;
}
}
if( errors > 0 )
return errors;
}
free_mtdata(seed);
// All done!
return 0;
}