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>
429 lines
15 KiB
C++
429 lines
15 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"
|
|
#if (defined( __APPLE__ ) || defined( __linux__ ))
|
|
#include <unistd.h>
|
|
#endif
|
|
|
|
const char *preprocessor_test_kernel[] = {
|
|
"__kernel void sample_test(__global int *dst)\n"
|
|
"{\n"
|
|
" dst[0] = TEST_MACRO;\n"
|
|
"\n"
|
|
"}\n" };
|
|
|
|
const char *preprocessor_existence_test_kernel[] = {
|
|
"__kernel void sample_test(__global int *dst)\n"
|
|
"{\n"
|
|
"#ifdef TEST_MACRO\n"
|
|
" dst[0] = 42;\n"
|
|
"#else\n"
|
|
" dst[0] = 24;\n"
|
|
"#endif\n"
|
|
"\n"
|
|
"}\n" };
|
|
|
|
const char *include_test_kernel[] = {
|
|
"#include \"./testIncludeFile.h\"\n"
|
|
"__kernel void sample_test(__global int *dst)\n"
|
|
"{\n"
|
|
" dst[0] = HEADER_FOUND;\n"
|
|
"\n"
|
|
"}\n" };
|
|
|
|
const char *options_test_kernel[] = {
|
|
"__kernel void sample_test(__global float *src, __global int *dst)\n"
|
|
"{\n"
|
|
" size_t tid = get_global_id(0);\n"
|
|
" dst[tid] = src[tid];\n"
|
|
"}\n" };
|
|
|
|
const char *optimization_options[] = {
|
|
"-cl-single-precision-constant",
|
|
"-cl-denorms-are-zero",
|
|
"-cl-opt-disable",
|
|
"-cl-strict-aliasing",
|
|
"-cl-mad-enable",
|
|
"-cl-no-signed-zeros",
|
|
"-cl-unsafe-math-optimizations",
|
|
"-cl-finite-math-only",
|
|
"-cl-fast-relaxed-math",
|
|
"-w",
|
|
"-Werror",
|
|
#if defined( __APPLE__ )
|
|
"-cl-opt-enable",
|
|
"-cl-auto-vectorize-enable"
|
|
#endif
|
|
};
|
|
|
|
cl_int get_result_from_program( cl_context context, cl_command_queue queue, cl_program program, cl_int *outValue )
|
|
{
|
|
cl_int error;
|
|
clKernelWrapper kernel = clCreateKernel( program, "sample_test", &error );
|
|
test_error( error, "Unable to create kernel from program" );
|
|
|
|
clMemWrapper outStream;
|
|
outStream = clCreateBuffer( context, (cl_mem_flags)(CL_MEM_READ_WRITE), sizeof(cl_int), NULL, &error );
|
|
test_error( error, "Unable to create test buffer" );
|
|
|
|
error = clSetKernelArg( kernel, 0, sizeof( outStream ), &outStream );
|
|
test_error( error, "Unable to set kernel argument" );
|
|
|
|
size_t threads[1] = { 1 };
|
|
|
|
error = clEnqueueNDRangeKernel( queue, kernel, 1, NULL, threads, NULL, 0, NULL, NULL );
|
|
test_error( error, "Unable to execute test kernel" );
|
|
|
|
error = clEnqueueReadBuffer( queue, outStream, true, 0, sizeof( cl_int ), outValue, 0, NULL, NULL );
|
|
test_error( error, "Unable to read output array!" );
|
|
|
|
return CL_SUCCESS;
|
|
}
|
|
|
|
int test_options_optimizations(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
|
|
{
|
|
int error;
|
|
cl_build_status status;
|
|
|
|
for(size_t i = 0; i < sizeof(optimization_options) / (sizeof(char*)); i++) {
|
|
|
|
clProgramWrapper program = clCreateProgramWithSource( context, 1, options_test_kernel, NULL, &error );
|
|
if( program == NULL || error != CL_SUCCESS )
|
|
{
|
|
log_error( "ERROR: Unable to create reference program!\n" );
|
|
return -1;
|
|
}
|
|
|
|
/* Build with the macro defined */
|
|
log_info("Testing optimization option '%s'\n", optimization_options[i]);
|
|
error = clBuildProgram( program, 1, &deviceID, optimization_options[i], NULL, NULL );
|
|
test_error( error, "Test program did not properly build" );
|
|
|
|
error = clGetProgramBuildInfo( program, deviceID, CL_PROGRAM_BUILD_STATUS, sizeof( status ), &status, NULL );
|
|
test_error( error, "Unable to get program build status" );
|
|
|
|
if( (int)status != CL_BUILD_SUCCESS )
|
|
{
|
|
log_info("Building with optimization option '%s' failed to compile!\n", optimization_options[i]);
|
|
print_error( error, "Failed to build with optimization defined")
|
|
return -1;
|
|
}
|
|
}
|
|
return 0;
|
|
}
|
|
|
|
int test_options_build_macro(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
|
|
{
|
|
int error;
|
|
clProgramWrapper program;
|
|
cl_build_status status;
|
|
|
|
|
|
program = clCreateProgramWithSource( context, 1, preprocessor_test_kernel, NULL, &error );
|
|
if( program == NULL || error != CL_SUCCESS )
|
|
{
|
|
log_error( "ERROR: Unable to create reference program!\n" );
|
|
return -1;
|
|
}
|
|
|
|
/* Build with the macro defined */
|
|
error = clBuildProgram( program, 1, &deviceID, "-DTEST_MACRO=1 ", NULL, NULL );
|
|
test_error( error, "Test program did not properly build" );
|
|
|
|
error = clGetProgramBuildInfo( program, deviceID, CL_PROGRAM_BUILD_STATUS, sizeof( status ), &status, NULL );
|
|
test_error( error, "Unable to get program build status" );
|
|
|
|
if( (int)status != CL_BUILD_SUCCESS )
|
|
{
|
|
print_error( error, "Failed to build with macro defined" );
|
|
return -1;
|
|
}
|
|
|
|
|
|
// Go ahead and run the program to verify results
|
|
cl_int firstResult, secondResult;
|
|
|
|
error = get_result_from_program( context, queue, program, &firstResult );
|
|
test_error( error, "Unable to get result from first program" );
|
|
|
|
if( firstResult != 1 )
|
|
{
|
|
log_error( "ERROR: Result from first program did not validate! (Expected 1, got %d)\n", firstResult );
|
|
return -1;
|
|
}
|
|
|
|
// Rebuild with a different value for the define macro, to make sure caching behaves properly
|
|
error = clBuildProgram( program, 1, &deviceID, "-DTEST_MACRO=5 ", NULL, NULL );
|
|
test_error( error, "Test program did not properly rebuild" );
|
|
|
|
error = get_result_from_program( context, queue, program, &secondResult );
|
|
test_error( error, "Unable to get result from second program" );
|
|
|
|
if( secondResult != 5 )
|
|
{
|
|
if( secondResult == firstResult )
|
|
log_error( "ERROR: Program result did not change with device macro change (program was not recompiled)!\n" );
|
|
else
|
|
log_error( "ERROR: Result from second program did not validate! (Expected 5, got %d)\n", secondResult );
|
|
return -1;
|
|
}
|
|
|
|
return 0;
|
|
}
|
|
|
|
int test_options_build_macro_existence(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
|
|
{
|
|
int error;
|
|
clProgramWrapper program;
|
|
|
|
|
|
// In this case, the program should still run without the macro, but it should give a different result
|
|
program = clCreateProgramWithSource( context, 1, preprocessor_existence_test_kernel, NULL, &error );
|
|
if( program == NULL || error != CL_SUCCESS )
|
|
{
|
|
log_error( "ERROR: Unable to create reference program!\n" );
|
|
return -1;
|
|
}
|
|
|
|
/* Build without the macro defined */
|
|
error = clBuildProgram( program, 1, &deviceID, NULL, NULL, NULL );
|
|
test_error( error, "Test program did not properly build" );
|
|
|
|
// Go ahead and run the program to verify results
|
|
cl_int firstResult, secondResult;
|
|
|
|
error = get_result_from_program( context, queue, program, &firstResult );
|
|
test_error( error, "Unable to get result from first program" );
|
|
|
|
if( firstResult != 24 )
|
|
{
|
|
log_error( "ERROR: Result from first program did not validate! (Expected 24, got %d)\n", firstResult );
|
|
return -1;
|
|
}
|
|
|
|
// Now compile again with the macro defined and verify a change in results
|
|
error = clBuildProgram( program, 1, &deviceID, "-DTEST_MACRO", NULL, NULL );
|
|
test_error( error, "Test program did not properly build" );
|
|
|
|
error = get_result_from_program( context, queue, program, &secondResult );
|
|
test_error( error, "Unable to get result from second program" );
|
|
|
|
if( secondResult != 42 )
|
|
{
|
|
if( secondResult == firstResult )
|
|
log_error( "ERROR: Program result did not change with device macro addition (program was not recompiled)!\n" );
|
|
else
|
|
log_error( "ERROR: Result from second program did not validate! (Expected 42, got %d)\n", secondResult );
|
|
return -1;
|
|
}
|
|
|
|
return 0;
|
|
}
|
|
|
|
int test_options_include_directory(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
|
|
{
|
|
int error;
|
|
|
|
char include_dir[4096] = {0};
|
|
const char *sep = "/";
|
|
|
|
clProgramWrapper program;
|
|
cl_build_status status;
|
|
|
|
#if defined( __APPLE__ )
|
|
char *path = 0;
|
|
path = getwd(NULL);
|
|
#elif defined( __linux__ )
|
|
char *path = 0;
|
|
path = get_current_dir_name();
|
|
#elif defined( WIN32 )
|
|
char *path = (char*)malloc(MAX_PATH);
|
|
if (!path) {
|
|
return -1;
|
|
}
|
|
GetCurrentDirectoryA(MAX_PATH, path);
|
|
#else
|
|
#error Please set the cwd for your platform
|
|
#endif
|
|
|
|
/* Try compiling the program first without the directory included Should fail. */
|
|
program = clCreateProgramWithSource( context, 1, include_test_kernel, NULL, &error );
|
|
if( program == NULL || error != CL_SUCCESS )
|
|
{
|
|
log_error( "ERROR: Unable to create reference program!\n" );
|
|
return -1;
|
|
}
|
|
|
|
/* Build with the include directory defined */
|
|
sprintf(include_dir, "-I %s%sincludeTestDirectory", path, sep);
|
|
|
|
// log_info("%s\n", include_dir);
|
|
error = clBuildProgram( program, 1, &deviceID, include_dir, NULL, NULL );
|
|
test_error( error, "Test program did not properly build" );
|
|
|
|
error = clGetProgramBuildInfo( program, deviceID, CL_PROGRAM_BUILD_STATUS, sizeof( status ), &status, NULL );
|
|
test_error( error, "Unable to get program build status" );
|
|
|
|
if( (int)status != CL_BUILD_SUCCESS )
|
|
{
|
|
print_error( error, "Failed to build with include directory" );
|
|
return -1;
|
|
}
|
|
|
|
// Go ahead and run the program to verify results
|
|
cl_int firstResult, secondResult;
|
|
|
|
error = get_result_from_program( context, queue, program, &firstResult );
|
|
test_error( error, "Unable to get result from first program" );
|
|
|
|
if( firstResult != 12 )
|
|
{
|
|
log_error( "ERROR: Result from first program did not validate! (Expected 12, got %d)\n", firstResult );
|
|
return -1;
|
|
}
|
|
|
|
// Rebuild with a different include directory
|
|
sprintf(include_dir, "-I %s%ssecondIncludeTestDirectory", path, sep);
|
|
error = clBuildProgram( program, 1, &deviceID, include_dir, NULL, NULL );
|
|
test_error( error, "Test program did not properly rebuild" );
|
|
|
|
error = get_result_from_program( context, queue, program, &secondResult );
|
|
test_error( error, "Unable to get result from second program" );
|
|
|
|
if( secondResult != 42 )
|
|
{
|
|
if( secondResult == firstResult )
|
|
log_error( "ERROR: Program result did not change with include path change (program was not recompiled)!\n" );
|
|
else
|
|
log_error( "ERROR: Result from second program did not validate! (Expected 42, got %d)\n", secondResult );
|
|
return -1;
|
|
}
|
|
|
|
free(path);
|
|
return 0;
|
|
}
|
|
|
|
const char *denorm_test_kernel[] = {
|
|
"__kernel void sample_test( float src1, float src2, __global float *dst)\n"
|
|
"{\n"
|
|
" dst[ 0 ] = src1 + src2;\n"
|
|
"\n"
|
|
"}\n" };
|
|
|
|
cl_int get_float_result_from_program( cl_context context, cl_command_queue queue, cl_program program, cl_float inA, cl_float inB, cl_float *outValue )
|
|
{
|
|
cl_int error;
|
|
|
|
clKernelWrapper kernel = clCreateKernel( program, "sample_test", &error );
|
|
test_error( error, "Unable to create kernel from program" );
|
|
|
|
clMemWrapper outStream = clCreateBuffer( context, (cl_mem_flags)(CL_MEM_READ_WRITE), sizeof(cl_float), NULL, &error );
|
|
test_error( error, "Unable to create test buffer" );
|
|
|
|
error = clSetKernelArg( kernel, 0, sizeof( cl_float ), &inA );
|
|
test_error( error, "Unable to set kernel argument" );
|
|
|
|
error = clSetKernelArg( kernel, 1, sizeof( cl_float ), &inB );
|
|
test_error( error, "Unable to set kernel argument" );
|
|
|
|
error = clSetKernelArg( kernel, 2, sizeof( outStream ), &outStream );
|
|
test_error( error, "Unable to set kernel argument" );
|
|
|
|
size_t threads[1] = { 1 };
|
|
|
|
error = clEnqueueNDRangeKernel( queue, kernel, 1, NULL, threads, NULL, 0, NULL, NULL );
|
|
test_error( error, "Unable to execute test kernel" );
|
|
|
|
error = clEnqueueReadBuffer( queue, outStream, true, 0, sizeof( cl_float ), outValue, 0, NULL, NULL );
|
|
test_error( error, "Unable to read output array!" );
|
|
|
|
return CL_SUCCESS;
|
|
}
|
|
|
|
int test_options_denorm_cache(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
|
|
{
|
|
int error;
|
|
|
|
clProgramWrapper program;
|
|
cl_build_status status;
|
|
|
|
|
|
// If denorms aren't even supported, testing this flag is pointless
|
|
cl_device_fp_config floatCaps = 0;
|
|
error = clGetDeviceInfo( deviceID, CL_DEVICE_SINGLE_FP_CONFIG, sizeof(floatCaps), &floatCaps, NULL);
|
|
test_error( error, "Unable to get device FP config" );
|
|
if( ( floatCaps & CL_FP_DENORM ) == 0 )
|
|
{
|
|
log_info( "Device does not support denormalized single-precision floats; skipping test.\n" );
|
|
return 0;
|
|
}
|
|
|
|
program = clCreateProgramWithSource( context, 1, denorm_test_kernel, NULL, &error );
|
|
test_error( error, "Unable to create test program" );
|
|
|
|
// Build first WITH the denorm flush flag
|
|
error = clBuildProgram( program, 1, &deviceID, "-cl-denorms-are-zero", NULL, NULL );
|
|
test_error( error, "Test program did not properly build" );
|
|
|
|
error = clGetProgramBuildInfo( program, deviceID, CL_PROGRAM_BUILD_STATUS, sizeof( status ), &status, NULL );
|
|
test_error( error, "Unable to get program build status" );
|
|
|
|
if( (int)status != CL_BUILD_SUCCESS )
|
|
{
|
|
print_error( error, "Failed to build with include directory" );
|
|
return -1;
|
|
}
|
|
|
|
// Note: the following in floating point is a subnormal number, thus adding two of them together
|
|
// should give us a subnormalized result. If denormals are flushed to zero, however, it'll give us zero instead
|
|
uint32_t intSubnormal = 0x00000001;
|
|
cl_float *input = (cl_float *)&intSubnormal;
|
|
cl_float firstResult, secondResult;
|
|
|
|
error = get_float_result_from_program( context, queue, program, *input, *input, &firstResult );
|
|
test_error( error, "Unable to get result from first program" );
|
|
|
|
// Note: since -cl-denorms-are-zero is a HINT, not a requirement, the result we got could
|
|
// either be subnormal (hint ignored) or zero (hint respected). Since either is technically
|
|
// valid, there isn't anything we can to do validate results for now
|
|
|
|
// Rebuild without flushing flag set
|
|
error = clBuildProgram( program, 1, &deviceID, NULL, NULL, NULL );
|
|
test_error( error, "Test program did not properly rebuild" );
|
|
|
|
error = get_float_result_from_program( context, queue, program, *input, *input, &secondResult );
|
|
test_error( error, "Unable to get result from second program" );
|
|
|
|
// Now, there are three possiblities here:
|
|
// 1. The denorms-are-zero hint is not respected, in which case the first and second result will be identical
|
|
// 2. The hint is respected, and the program was properly rebuilt, in which case the first result will be zero and the second non-zero
|
|
// 3. The hint is respected, but the program was not properly rebuilt, in which case both results will be zero
|
|
// 3 is the only error condition we need to look for
|
|
uint32_t *fPtr = (uint32_t *)&firstResult;
|
|
uint32_t *sPtr = (uint32_t *)&secondResult;
|
|
|
|
if( ( *fPtr == 0 ) && ( *sPtr == 0 ) )
|
|
{
|
|
log_error( "ERROR: Program result didn't change when -cl-denorms-are-zero flag was removed.\n"
|
|
"First result (should be zero): 0x%08x, Second result (should be non-zero): 0x%08x\n",
|
|
*fPtr, *sPtr );
|
|
return -1;
|
|
}
|
|
|
|
return 0;
|
|
}
|
|
|