test_conformance/api/test_kernel_arg_multi_setup.cpp - external/github.com/KhronosGroup/OpenCL-CTS - Git at Google

 //
 // 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 "harness/conversions.h"

 // This test is designed to stress passing multiple vector parameters to kernels and verifying access between them all

 const char *multi_arg_kernel_source_pattern =
 "__kernel void sample_test(__global %s *src1, __global %s *src2, __global %s *src3, __global %s *dst1, __global %s *dst2, __global %s *dst3 )\n"
 "{\n"
 "    int tid = get_global_id(0);\n"
 "    dst1[tid] = src1[tid];\n"
 "    dst2[tid] = src2[tid];\n"
 "    dst3[tid] = src3[tid];\n"
 "}\n";

 #define MAX_ERROR_TOLERANCE 0.0005f

 int test_multi_arg_set(cl_device_id device, cl_context context, cl_command_queue queue,
                        ExplicitType vec1Type, int vec1Size,
                        ExplicitType vec2Type, int vec2Size,
                        ExplicitType vec3Type, int vec3Size, MTdata d)
 {
     clProgramWrapper program;
     clKernelWrapper kernel;
     int error, i, j;
     clMemWrapper streams[ 6 ];
     size_t threads[1], localThreads[1];
     char programSrc[ 10248 ], vec1Name[ 64 ], vec2Name[ 64 ], vec3Name[ 64 ];
     char sizeNames[][ 4 ] = { "", "2", "3", "4", "", "", "", "8" };
     const char *ptr;
     void *initData[3], *resultData[3];


     // Create the program source
     sprintf( vec1Name, "%s%s", get_explicit_type_name( vec1Type ), sizeNames[ vec1Size - 1 ] );
     sprintf( vec2Name, "%s%s", get_explicit_type_name( vec2Type ), sizeNames[ vec2Size - 1 ] );
     sprintf( vec3Name, "%s%s", get_explicit_type_name( vec3Type ), sizeNames[ vec3Size - 1 ] );

     sprintf( programSrc, multi_arg_kernel_source_pattern,
             vec1Name, vec2Name, vec3Name, vec1Name, vec2Name, vec3Name,
             vec1Size, vec1Size, vec2Size, vec2Size, vec3Size, vec3Size );
     ptr = programSrc;

     // Create our testing kernel
     error = create_single_kernel_helper( context, &program, &kernel, 1, &ptr, "sample_test" );
     test_error( error, "Unable to create testing kernel" );

     // Get thread dimensions
     threads[0] = 1024;
     error = get_max_common_work_group_size( context, kernel, threads[0], &localThreads[0] );
     test_error( error, "Unable to get work group size for kernel" );

     // Create input streams
     initData[ 0 ] = create_random_data( vec1Type, d, (unsigned int)threads[ 0 ] * vec1Size );
     streams[0] =
         clCreateBuffer(context, CL_MEM_COPY_HOST_PTR,
                        get_explicit_type_size(vec1Type) * threads[0] * vec1Size,
                        initData[0], &error);
     test_error( error, "Unable to create testing stream" );

     initData[ 1 ] = create_random_data( vec2Type, d, (unsigned int)threads[ 0 ] * vec2Size );
     streams[1] =
         clCreateBuffer(context, CL_MEM_COPY_HOST_PTR,
                        get_explicit_type_size(vec2Type) * threads[0] * vec2Size,
                        initData[1], &error);
     test_error( error, "Unable to create testing stream" );

     initData[ 2 ] = create_random_data( vec3Type, d, (unsigned int)threads[ 0 ] * vec3Size );
     streams[2] =
         clCreateBuffer(context, CL_MEM_COPY_HOST_PTR,
                        get_explicit_type_size(vec3Type) * threads[0] * vec3Size,
                        initData[2], &error);
     test_error( error, "Unable to create testing stream" );

     streams[3] = clCreateBuffer(
         context, CL_MEM_READ_WRITE,
         get_explicit_type_size(vec1Type) * threads[0] * vec1Size, NULL, &error);
     test_error( error, "Unable to create testing stream" );

     streams[4] = clCreateBuffer(
         context, CL_MEM_READ_WRITE,
         get_explicit_type_size(vec2Type) * threads[0] * vec2Size, NULL, &error);
     test_error( error, "Unable to create testing stream" );

     streams[5] = clCreateBuffer(
         context, CL_MEM_READ_WRITE,
         get_explicit_type_size(vec3Type) * threads[0] * vec3Size, NULL, &error);
     test_error( error, "Unable to create testing stream" );

     // Set the arguments
     error = 0;
     for( i = 0; i < 6; i++ )
         error |= clSetKernelArg( kernel, i, sizeof( cl_mem ), &streams[ i ] );
     test_error( error, "Unable to set arguments for kernel" );

     // Execute!
     error = clEnqueueNDRangeKernel( queue, kernel, 1, NULL, threads, localThreads, 0, NULL, NULL );
     test_error( error, "Unable to execute kernel" );

     // Read results
     resultData[0] = malloc( get_explicit_type_size( vec1Type ) * vec1Size * threads[0] );
     resultData[1] = malloc( get_explicit_type_size( vec2Type ) * vec2Size * threads[0] );
     resultData[2] = malloc( get_explicit_type_size( vec3Type ) * vec3Size * threads[0] );
     error = clEnqueueReadBuffer( queue, streams[ 3 ], CL_TRUE, 0, get_explicit_type_size( vec1Type ) * vec1Size * threads[ 0 ], resultData[0], 0, NULL, NULL );
     error |= clEnqueueReadBuffer( queue, streams[ 4 ], CL_TRUE, 0, get_explicit_type_size( vec2Type ) * vec2Size * threads[ 0 ], resultData[1], 0, NULL, NULL );
     error |= clEnqueueReadBuffer( queue, streams[ 5 ], CL_TRUE, 0, get_explicit_type_size( vec3Type ) * vec3Size * threads[ 0 ], resultData[2], 0, NULL, NULL );
     test_error( error, "Unable to read result stream" );

     // Verify
     char *ptr1 = (char *)initData[ 0 ], *ptr2 = (char *)resultData[ 0 ];
     size_t span = get_explicit_type_size( vec1Type );
     for( i = 0; i < (int)threads[0]; i++ )
     {
         for( j = 0; j < vec1Size; j++ )
         {
             if( memcmp( ptr1 + span * j , ptr2 + span * j, span ) != 0 )
             {
                 log_error( "ERROR: Value did not validate for component %d of item %d of stream 0!\n", j, i );
                 free( initData[ 0 ] );
                 free( initData[ 1 ] );
                 free( initData[ 2 ] );
                 free( resultData[ 0 ] );
                 free( resultData[ 1 ] );
                 free( resultData[ 2 ] );
                 return -1;
             }
         }
         ptr1 += span * vec1Size;
         ptr2 += span * vec1Size;
     }

     ptr1 = (char *)initData[ 1 ];
     ptr2 = (char *)resultData[ 1 ];
     span = get_explicit_type_size( vec2Type );
     for( i = 0; i < (int)threads[0]; i++ )
     {
         for( j = 0; j < vec2Size; j++ )
         {
             if( memcmp( ptr1 + span * j , ptr2 + span * j, span ) != 0 )
             {
                 log_error( "ERROR: Value did not validate for component %d of item %d of stream 1!\n", j, i );
                 free( initData[ 0 ] );
                 free( initData[ 1 ] );
                 free( initData[ 2 ] );
                 free( resultData[ 0 ] );
                 free( resultData[ 1 ] );
                 free( resultData[ 2 ] );
                 return -1;
             }
         }
         ptr1 += span * vec2Size;
         ptr2 += span * vec2Size;
     }

     ptr1 = (char *)initData[ 2 ];
     ptr2 = (char *)resultData[ 2 ];
     span = get_explicit_type_size( vec3Type );
     for( i = 0; i < (int)threads[0]; i++ )
     {
         for( j = 0; j < vec3Size; j++ )
         {
             if( memcmp( ptr1 + span * j , ptr2 + span * j, span ) != 0 )
             {
                 log_error( "ERROR: Value did not validate for component %d of item %d of stream 2!\n", j, i );
                 free( initData[ 0 ] );
                 free( initData[ 1 ] );
                 free( initData[ 2 ] );
                 free( resultData[ 0 ] );
                 free( resultData[ 1 ] );
                 free( resultData[ 2 ] );
                 return -1;
             }
         }
         ptr1 += span * vec3Size;
         ptr2 += span * vec3Size;
     }

     // If we got here, everything verified successfully
     free( initData[ 0 ] );
     free( initData[ 1 ] );
     free( initData[ 2 ] );
     free( resultData[ 0 ] );
     free( resultData[ 1 ] );
     free( resultData[ 2 ] );

     return 0;
 }

 int test_kernel_arg_multi_setup_exhaustive(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
 {
     // Loop through every combination of input and output types
     ExplicitType types[] = { kChar, kShort, kInt, kFloat, kNumExplicitTypes };
     int type1, type2, type3;
     int size1, size2, size3;
     RandomSeed seed( gRandomSeed );

     log_info( "\n" ); // for formatting

     for( type1 = 0; types[ type1 ] != kNumExplicitTypes; type1++ )
     {
         for( type2 = 0; types[ type2 ] != kNumExplicitTypes; type2++ )
         {
             for( type3 = 0; types[ type3 ] != kNumExplicitTypes; type3++ )
             {
                 log_info( "\n\ttesting %s, %s, %s...", get_explicit_type_name( types[ type1 ] ), get_explicit_type_name( types[ type2 ] ), get_explicit_type_name( types[ type3 ] ) );

                 // Loop through every combination of vector size
                 for( size1 = 2; size1 <= 8; size1 <<= 1 )
                 {
                     for( size2 = 2; size2 <= 8; size2 <<= 1 )
                     {
                         for( size3 = 2; size3 <= 8; size3 <<= 1 )
                         {
                             log_info(".");
                             fflush( stdout);
                             if( test_multi_arg_set( device, context, queue,
                                                    types[ type1 ], size1,
                                                    types[ type2 ], size2,
                                                    types[ type3 ], size3, seed ) )
                                 return -1;
                         }
                     }
                 }
             }
         }
     }
     log_info( "\n" );
     return 0;
 }

 int test_kernel_arg_multi_setup_random(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
 {
     // Loop through a selection of combinations
     ExplicitType types[] = { kChar, kShort, kInt, kFloat, kNumExplicitTypes };
     int type1, type2, type3;
     int size1, size2, size3;
     RandomSeed seed( gRandomSeed );

     num_elements = 3*3*3*4;
     log_info( "Testing %d random configurations\n", num_elements );

     // Loop through every combination of vector size
     for( size1 = 2; size1 <= 8; size1 <<= 1 )
     {
         for( size2 = 2; size2 <= 8; size2 <<= 1 )
         {
             for( size3 = 2; size3 <= 8; size3 <<= 1 )
             {
                 // Loop through 4 type combinations for each size combination
                 int n;
                 for (n=0; n<4; n++) {
                     type1 = (int)get_random_float(0,4, seed);
                     type2 = (int)get_random_float(0,4, seed);
                     type3 = (int)get_random_float(0,4, seed);


                     log_info( "\ttesting %s%d, %s%d, %s%d...\n",
                              get_explicit_type_name( types[ type1 ] ), size1,
                              get_explicit_type_name( types[ type2 ] ), size2,
                              get_explicit_type_name( types[ type3 ] ), size3 );

                     if( test_multi_arg_set( device, context, queue,
                                            types[ type1 ], size1,
                                            types[ type2 ], size2,
                                            types[ type3 ], size3, seed ) )
                         return -1;
                 }
             }
         }
     }
     return 0;
 }
	//
	// 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 "harness/conversions.h"

	// This test is designed to stress passing multiple vector parameters to kernels and verifying access between them all

	const char *multi_arg_kernel_source_pattern =
	"__kernel void sample_test(__global %s src1, __global %s src2, __global %s src3, __global %s dst1, __global %s dst2, __global %s dst3 )\n"
	"{\n"
	" int tid = get_global_id(0);\n"
	" dst1[tid] = src1[tid];\n"
	" dst2[tid] = src2[tid];\n"
	" dst3[tid] = src3[tid];\n"
	"}\n";

	#define MAX_ERROR_TOLERANCE 0.0005f

	int test_multi_arg_set(cl_device_id device, cl_context context, cl_command_queue queue,
	ExplicitType vec1Type, int vec1Size,
	ExplicitType vec2Type, int vec2Size,
	ExplicitType vec3Type, int vec3Size, MTdata d)
	{
	clProgramWrapper program;
	clKernelWrapper kernel;
	int error, i, j;
	clMemWrapper streams[ 6 ];
	size_t threads[1], localThreads[1];
	char programSrc[ 10248 ], vec1Name[ 64 ], vec2Name[ 64 ], vec3Name[ 64 ];
	char sizeNames[][ 4 ] = { "", "2", "3", "4", "", "", "", "8" };
	const char *ptr;
	void initData[3], resultData[3];


	// Create the program source
	sprintf( vec1Name, "%s%s", get_explicit_type_name( vec1Type ), sizeNames[ vec1Size - 1 ] );
	sprintf( vec2Name, "%s%s", get_explicit_type_name( vec2Type ), sizeNames[ vec2Size - 1 ] );
	sprintf( vec3Name, "%s%s", get_explicit_type_name( vec3Type ), sizeNames[ vec3Size - 1 ] );

	sprintf( programSrc, multi_arg_kernel_source_pattern,
	vec1Name, vec2Name, vec3Name, vec1Name, vec2Name, vec3Name,
	vec1Size, vec1Size, vec2Size, vec2Size, vec3Size, vec3Size );
	ptr = programSrc;

	// Create our testing kernel
	error = create_single_kernel_helper( context, &program, &kernel, 1, &ptr, "sample_test" );
	test_error( error, "Unable to create testing kernel" );

	// Get thread dimensions
	threads[0] = 1024;
	error = get_max_common_work_group_size( context, kernel, threads[0], &localThreads[0] );
	test_error( error, "Unable to get work group size for kernel" );

	// Create input streams
	initData[ 0 ] = create_random_data( vec1Type, d, (unsigned int)threads[ 0 ] * vec1Size );
	streams[0] =
	clCreateBuffer(context, CL_MEM_COPY_HOST_PTR,
	get_explicit_type_size(vec1Type) * threads[0] * vec1Size,
	initData[0], &error);
	test_error( error, "Unable to create testing stream" );

	initData[ 1 ] = create_random_data( vec2Type, d, (unsigned int)threads[ 0 ] * vec2Size );
	streams[1] =
	clCreateBuffer(context, CL_MEM_COPY_HOST_PTR,
	get_explicit_type_size(vec2Type) * threads[0] * vec2Size,
	initData[1], &error);
	test_error( error, "Unable to create testing stream" );

	initData[ 2 ] = create_random_data( vec3Type, d, (unsigned int)threads[ 0 ] * vec3Size );
	streams[2] =
	clCreateBuffer(context, CL_MEM_COPY_HOST_PTR,
	get_explicit_type_size(vec3Type) * threads[0] * vec3Size,
	initData[2], &error);
	test_error( error, "Unable to create testing stream" );

	streams[3] = clCreateBuffer(
	context, CL_MEM_READ_WRITE,
	get_explicit_type_size(vec1Type) * threads[0] * vec1Size, NULL, &error);
	test_error( error, "Unable to create testing stream" );

	streams[4] = clCreateBuffer(
	context, CL_MEM_READ_WRITE,
	get_explicit_type_size(vec2Type) * threads[0] * vec2Size, NULL, &error);
	test_error( error, "Unable to create testing stream" );

	streams[5] = clCreateBuffer(
	context, CL_MEM_READ_WRITE,
	get_explicit_type_size(vec3Type) * threads[0] * vec3Size, NULL, &error);
	test_error( error, "Unable to create testing stream" );

	// Set the arguments
	error = 0;
	for( i = 0; i < 6; i++ )
	error \|= clSetKernelArg( kernel, i, sizeof( cl_mem ), &streams[ i ] );
	test_error( error, "Unable to set arguments for kernel" );

	// Execute!
	error = clEnqueueNDRangeKernel( queue, kernel, 1, NULL, threads, localThreads, 0, NULL, NULL );
	test_error( error, "Unable to execute kernel" );

	// Read results
	resultData[0] = malloc( get_explicit_type_size( vec1Type ) * vec1Size * threads[0] );
	resultData[1] = malloc( get_explicit_type_size( vec2Type ) * vec2Size * threads[0] );
	resultData[2] = malloc( get_explicit_type_size( vec3Type ) * vec3Size * threads[0] );
	error = clEnqueueReadBuffer( queue, streams[ 3 ], CL_TRUE, 0, get_explicit_type_size( vec1Type ) * vec1Size * threads[ 0 ], resultData[0], 0, NULL, NULL );
	error \|= clEnqueueReadBuffer( queue, streams[ 4 ], CL_TRUE, 0, get_explicit_type_size( vec2Type ) * vec2Size * threads[ 0 ], resultData[1], 0, NULL, NULL );
	error \|= clEnqueueReadBuffer( queue, streams[ 5 ], CL_TRUE, 0, get_explicit_type_size( vec3Type ) * vec3Size * threads[ 0 ], resultData[2], 0, NULL, NULL );
	test_error( error, "Unable to read result stream" );

	// Verify
	char ptr1 = (char )initData[ 0 ], ptr2 = (char )resultData[ 0 ];
	size_t span = get_explicit_type_size( vec1Type );
	for( i = 0; i < (int)threads[0]; i++ )
	{
	for( j = 0; j < vec1Size; j++ )
	{
	if( memcmp( ptr1 + span * j , ptr2 + span * j, span ) != 0 )
	{
	log_error( "ERROR: Value did not validate for component %d of item %d of stream 0!\n", j, i );
	free( initData[ 0 ] );
	free( initData[ 1 ] );
	free( initData[ 2 ] );
	free( resultData[ 0 ] );
	free( resultData[ 1 ] );
	free( resultData[ 2 ] );
	return -1;
	}
	}
	ptr1 += span * vec1Size;
	ptr2 += span * vec1Size;
	}

	ptr1 = (char *)initData[ 1 ];
	ptr2 = (char *)resultData[ 1 ];
	span = get_explicit_type_size( vec2Type );
	for( i = 0; i < (int)threads[0]; i++ )
	{
	for( j = 0; j < vec2Size; j++ )
	{
	if( memcmp( ptr1 + span * j , ptr2 + span * j, span ) != 0 )
	{
	log_error( "ERROR: Value did not validate for component %d of item %d of stream 1!\n", j, i );
	free( initData[ 0 ] );
	free( initData[ 1 ] );
	free( initData[ 2 ] );
	free( resultData[ 0 ] );
	free( resultData[ 1 ] );
	free( resultData[ 2 ] );
	return -1;
	}
	}
	ptr1 += span * vec2Size;
	ptr2 += span * vec2Size;
	}

	ptr1 = (char *)initData[ 2 ];
	ptr2 = (char *)resultData[ 2 ];
	span = get_explicit_type_size( vec3Type );
	for( i = 0; i < (int)threads[0]; i++ )
	{
	for( j = 0; j < vec3Size; j++ )
	{
	if( memcmp( ptr1 + span * j , ptr2 + span * j, span ) != 0 )
	{
	log_error( "ERROR: Value did not validate for component %d of item %d of stream 2!\n", j, i );
	free( initData[ 0 ] );
	free( initData[ 1 ] );
	free( initData[ 2 ] );
	free( resultData[ 0 ] );
	free( resultData[ 1 ] );
	free( resultData[ 2 ] );
	return -1;
	}
	}
	ptr1 += span * vec3Size;
	ptr2 += span * vec3Size;
	}

	// If we got here, everything verified successfully
	free( initData[ 0 ] );
	free( initData[ 1 ] );
	free( initData[ 2 ] );
	free( resultData[ 0 ] );
	free( resultData[ 1 ] );
	free( resultData[ 2 ] );

	return 0;
	}

	int test_kernel_arg_multi_setup_exhaustive(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
	{
	// Loop through every combination of input and output types
	ExplicitType types[] = { kChar, kShort, kInt, kFloat, kNumExplicitTypes };
	int type1, type2, type3;
	int size1, size2, size3;
	RandomSeed seed( gRandomSeed );

	log_info( "\n" ); // for formatting

	for( type1 = 0; types[ type1 ] != kNumExplicitTypes; type1++ )
	{
	for( type2 = 0; types[ type2 ] != kNumExplicitTypes; type2++ )
	{
	for( type3 = 0; types[ type3 ] != kNumExplicitTypes; type3++ )
	{
	log_info( "\n\ttesting %s, %s, %s...", get_explicit_type_name( types[ type1 ] ), get_explicit_type_name( types[ type2 ] ), get_explicit_type_name( types[ type3 ] ) );

	// Loop through every combination of vector size
	for( size1 = 2; size1 <= 8; size1 <<= 1 )
	{
	for( size2 = 2; size2 <= 8; size2 <<= 1 )
	{
	for( size3 = 2; size3 <= 8; size3 <<= 1 )
	{
	log_info(".");
	fflush( stdout);
	if( test_multi_arg_set( device, context, queue,
	types[ type1 ], size1,
	types[ type2 ], size2,
	types[ type3 ], size3, seed ) )
	return -1;
	}
	}
	}
	}
	}
	}
	log_info( "\n" );
	return 0;
	}

	int test_kernel_arg_multi_setup_random(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
	{
	// Loop through a selection of combinations
	ExplicitType types[] = { kChar, kShort, kInt, kFloat, kNumExplicitTypes };
	int type1, type2, type3;
	int size1, size2, size3;
	RandomSeed seed( gRandomSeed );

	num_elements = 333*4;
	log_info( "Testing %d random configurations\n", num_elements );

	// Loop through every combination of vector size
	for( size1 = 2; size1 <= 8; size1 <<= 1 )
	{
	for( size2 = 2; size2 <= 8; size2 <<= 1 )
	{
	for( size3 = 2; size3 <= 8; size3 <<= 1 )
	{
	// Loop through 4 type combinations for each size combination
	int n;
	for (n=0; n<4; n++) {
	type1 = (int)get_random_float(0,4, seed);
	type2 = (int)get_random_float(0,4, seed);
	type3 = (int)get_random_float(0,4, seed);


	log_info( "\ttesting %s%d, %s%d, %s%d...\n",
	get_explicit_type_name( types[ type1 ] ), size1,
	get_explicit_type_name( types[ type2 ] ), size2,
	get_explicit_type_name( types[ type3 ] ), size3 );

	if( test_multi_arg_set( device, context, queue,
	types[ type1 ], size1,
	types[ type2 ], size2,
	types[ type3 ], size3, seed ) )
	return -1;
	}
	}
	}
	}
	return 0;
	}