blob: f3ee4312246ed1a535d8eef017284492b0908fc1 [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 "testBase.h"
#if defined(_WIN32)
#include <time.h>
#elif defined(__linux__) || defined(__APPLE__)
#include <sys/time.h>
#include <unistd.h>
#endif
#include "harness/conversions.h"
#define MAX_LINE_SIZE_IN_PROGRAM 1024
#define MAX_LOG_SIZE_IN_PROGRAM 2048
const char *sample_kernel_start =
"__kernel void sample_test(__global float *src, __global int *dst)\n"
"{\n"
" float temp;\n"
" int tid = get_global_id(0);\n";
const char *sample_kernel_end = "}\n";
const char *sample_kernel_lines[] = {
"dst[tid] = src[tid];\n",
"dst[tid] = src[tid] * 3.f;\n",
"temp = src[tid] / 4.f;\n",
"dst[tid] = dot(temp,src[tid]);\n",
"dst[tid] = dst[tid] + temp;\n" };
/* I compile and link therefore I am. Robert Ioffe */
/* The following kernels are used in testing Improved Compilation and Linking feature */
const char *simple_kernel =
"__kernel void\n"
"CopyBuffer(\n"
" __global float* src,\n"
" __global float* dst )\n"
"{\n"
" int id = (int)get_global_id(0);\n"
" dst[id] = src[id];\n"
"}\n";
const char *simple_kernel_with_defines =
"__kernel void\n"
"CopyBuffer(\n"
" __global float* src,\n"
" __global float* dst )\n"
"{\n"
" int id = (int)get_global_id(0);\n"
" float temp = src[id] - 42;\n"
" dst[id] = FIRST + temp + SECOND;\n"
"}\n";
const char *simple_kernel_template =
"__kernel void\n"
"CopyBuffer%d(\n"
" __global float* src,\n"
" __global float* dst )\n"
"{\n"
" int id = (int)get_global_id(0);\n"
" dst[id] = src[id];\n"
"}\n";
const char *composite_kernel_start =
"__kernel void\n"
"CompositeKernel(\n"
" __global float* src,\n"
" __global float* dst )\n"
"{\n";
const char *composite_kernel_end = "}\n";
const char *composite_kernel_template =
" CopyBuffer%d(src, dst);\n";
const char *composite_kernel_extern_template =
"extern __kernel void\n"
"CopyBuffer%d(\n"
" __global float* src,\n"
" __global float* dst );\n";
const char *another_simple_kernel =
"extern __kernel void\n"
"CopyBuffer(\n"
" __global float* src,\n"
" __global float* dst );\n"
"__kernel void\n"
"AnotherCopyBuffer(\n"
" __global float* src,\n"
" __global float* dst )\n"
"{\n"
" CopyBuffer(src, dst);\n"
"}\n";
const char* simple_header =
"extern __kernel void\n"
"CopyBuffer(\n"
" __global float* src,\n"
" __global float* dst );\n";
const char* simple_header_name = "simple_header.h";
const char* another_simple_kernel_with_header =
"#include \"simple_header.h\"\n"
"__kernel void\n"
"AnotherCopyBuffer(\n"
" __global float* src,\n"
" __global float* dst )\n"
"{\n"
" CopyBuffer(src, dst);\n"
"}\n";
const char* header_name_templates[4] = { "simple_header%d.h",
"foo/simple_header%d.h",
"foo/bar/simple_header%d.h",
"foo/bar/baz/simple_header%d.h"};
const char* include_header_name_templates[4] = { "#include \"simple_header%d.h\"\n",
"#include \"foo/simple_header%d.h\"\n",
"#include \"foo/bar/simple_header%d.h\"\n",
"#include \"foo/bar/baz/simple_header%d.h\"\n"};
const char* compile_extern_var = "extern constant float foo;\n";
const char* compile_extern_struct = "extern constant struct bar bart;\n";
const char* compile_extern_function = "extern int baz(int, int);\n";
const char* compile_static_var = "static constant float foo = 2.78;\n";
const char* compile_static_struct = "static constant struct bar {float x, y, z, r; int color; } foo = {3.14159};\n";
const char* compile_static_function = "static int foo(int x, int y) { return x*x + y*y; }\n";
const char* compile_regular_var = "constant float foo = 4.0f;\n";
const char* compile_regular_struct = "constant struct bar {float x, y, z, r; int color; } foo = {0.f, 0.f, 0.f, 0.f, 0};\n";
const char* compile_regular_function = "int foo(int x, int y) { return x*x + y*y; }\n";
const char* link_static_var_access = // use with compile_static_var
"extern constant float foo;\n"
"float access_foo() { return foo; }\n";
const char* link_static_struct_access = // use with compile_static_struct
"extern constant struct bar{float x, y, z, r; int color; } foo;\n"
"struct bar access_foo() {return foo; }\n";
const char* link_static_function_access = // use with compile_static_function
"extern int foo(int, int);\n"
"int access_foo() { int blah = foo(3, 4); return blah + 5; }\n";
int test_large_single_compile(cl_context context, cl_device_id deviceID, unsigned int numLines)
{
int error;
cl_program program;
const char **lines;
unsigned int numChoices, i;
MTdata d;
/* First, allocate the array for our line pointers */
lines = (const char **)malloc( numLines * sizeof( const char * ) );
if (lines == NULL) {
log_error( "ERROR: Unable to allocate lines array with %d lines! (in %s:%d)\n", numLines, __FILE__, __LINE__);
return -1;
}
/* First and last lines are easy */
lines[ 0 ] = sample_kernel_start;
lines[ numLines - 1 ] = sample_kernel_end;
numChoices = sizeof( sample_kernel_lines ) / sizeof( sample_kernel_lines[ 0 ] );
/* Fill the rest with random lines to hopefully prevent much optimization */
d = init_genrand( gRandomSeed );
for( i = 1; i < numLines - 1; i++ )
{
lines[ i ] = sample_kernel_lines[ genrand_int32(d) % numChoices ];
}
free_mtdata(d); d = NULL;
/* Try to create a program with these lines */
error = create_single_kernel_helper_create_program(context, &program, numLines, lines);
if( program == NULL || error != CL_SUCCESS )
{
log_error( "ERROR: Unable to create long test program with %d lines! (%s in %s:%d)", numLines, IGetErrorString( error ), __FILE__, __LINE__ );
free( lines );
if (program != NULL)
{
error = clReleaseProgram( program );
test_error( error, "Unable to release a program object" );
}
return -1;
}
/* Build it */
error = clBuildProgram( program, 1, &deviceID, NULL, NULL, NULL );
test_error( error, "Unable to build a long program" );
/* All done! */
error = clReleaseProgram( program );
test_error( error, "Unable to release a program object" );
free( lines );
return 0;
}
int test_large_compile(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
{
unsigned int toTest[] = { 64, 128, 256, 512, 1024, 2048, 4096, 0 }; //8192, 16384, 32768, 0 };
unsigned int i;
log_info( "Testing large compiles...this might take awhile...\n" );
for( i = 0; toTest[ i ] != 0; i++ )
{
log_info( " %d...\n", toTest[ i ] );
#if defined(_WIN32)
clock_t start = clock();
#elif defined(__linux__) || defined(__APPLE__)
timeval time1, time2;
gettimeofday(&time1, NULL);
#endif
if( test_large_single_compile( context, deviceID, toTest[ i ] ) != 0 )
{
log_error( "ERROR: long program test failed for %d lines! (in %s:%d)\n", toTest[ i ], __FILE__, __LINE__);
return -1;
}
#if defined(_WIN32)
clock_t end = clock();
log_perf( (float)( end - start ) / (float)CLOCKS_PER_SEC, false, "clock() time in secs", "%d lines", toTest[i] );
#elif defined(__linux__) || defined(__APPLE__)
gettimeofday(&time2, NULL);
log_perf( (float)(float)(time2.tv_sec - time1.tv_sec) + 1.0e-6 * (time2.tv_usec - time1.tv_usec) , false, "wall time in secs", "%d lines", toTest[i] );
#endif
}
return 0;
}
static int verifyCopyBuffer(cl_context context, cl_command_queue queue, cl_kernel kernel);
#if defined(__APPLE__) || defined(__linux)
#define _strdup strdup
#endif
int test_large_multi_file_library(cl_context context, cl_device_id deviceID, cl_command_queue queue, unsigned int numLines)
{
int error;
cl_program program;
cl_program *simple_kernels;
const char **lines;
unsigned int i;
char buffer[MAX_LINE_SIZE_IN_PROGRAM];
simple_kernels = (cl_program*)malloc(numLines*sizeof(cl_program));
if (simple_kernels == NULL) {
log_error( "ERROR: Unable to allocate kernels array with %d kernels! (in %s:%d)\n", numLines, __FILE__, __LINE__);
return -1;
}
/* First, allocate the array for our line pointers */
lines = (const char **)malloc( (2*numLines + 2) * sizeof( const char * ) );
if (lines == NULL) {
free(simple_kernels);
log_error( "ERROR: Unable to allocate lines array with %d lines! (in %s:%d)\n", (2*numLines + 2), __FILE__, __LINE__ );
return -1;
}
for( i = 0; i < numLines; i++)
{
sprintf(buffer, composite_kernel_extern_template, i);
lines[i] = _strdup(buffer);
}
/* First and last lines are easy */
lines[ numLines ] = composite_kernel_start;
lines[ 2* numLines + 1] = composite_kernel_end;
/* Fill the rest with templated kernels */
for( i = numLines + 1; i < 2* numLines + 1; i++ )
{
sprintf(buffer, composite_kernel_template, i - numLines - 1);
lines[ i ] = _strdup(buffer);
}
/* Try to create a program with these lines */
error = create_single_kernel_helper_create_program(context, &program, 2 * numLines + 2, lines);
if( program == NULL || error != CL_SUCCESS )
{
log_error( "ERROR: Unable to create long test program with %d lines! (%s) (in %s:%d)\n", numLines, IGetErrorString( error ), __FILE__, __LINE__ );
free( simple_kernels );
for( i = 0; i < numLines; i++)
{
free( (void*)lines[i] );
free( (void*)lines[i+numLines+1] );
}
free( lines );
if (program != NULL)
{
error = clReleaseProgram( program );
test_error( error, "Unable to release program object" );
}
return -1;
}
/* Compile it */
error = clCompileProgram(program, 1, &deviceID, NULL, 0, NULL, NULL, NULL, NULL);
test_error( error, "Unable to compile a simple program" );
/* Create and compile templated kernels */
for( i = 0; i < numLines; i++)
{
sprintf(buffer, simple_kernel_template, i);
const char* kernel_source = _strdup(buffer);
simple_kernels[i] = clCreateProgramWithSource( context, 1, &kernel_source, NULL, &error );
if( simple_kernels[i] == NULL || error != CL_SUCCESS )
{
log_error( "ERROR: Unable to create long test program with %d lines! (%s) (in %s:%d)\n", numLines, IGetErrorString( error ), __FILE__, __LINE__ );
return -1;
}
/* Compile it */
error = clCompileProgram(simple_kernels[i], 1, &deviceID, NULL, 0, NULL, NULL, NULL, NULL);
test_error( error, "Unable to compile a simple program" );
free((void*)kernel_source);
}
/* Create library out of compiled templated kernels */
cl_program my_newly_minted_library = clLinkProgram(context, 1, &deviceID, "-create-library", numLines, simple_kernels, NULL, NULL, &error);
test_error( error, "Unable to create a multi-line library" );
/* Link the program that calls the kernels and the library that contains them */
cl_program programs[2] = { program, my_newly_minted_library };
cl_program my_newly_linked_program = clLinkProgram(context, 1, &deviceID, NULL, 2, programs, NULL, NULL, &error);
test_error( error, "Unable to link a program with a library" );
// Create the composite kernel
cl_kernel kernel = clCreateKernel(my_newly_linked_program, "CompositeKernel", &error);
test_error( error, "Unable to create a composite kernel" );
// Run the composite kernel and verify the results
error = verifyCopyBuffer(context, queue, kernel);
if (error != CL_SUCCESS)
return error;
/* All done! */
error = clReleaseProgram( program );
test_error( error, "Unable to release program object" );
for( i = 0; i < numLines; i++)
{
free( (void*)lines[i] );
free( (void*)lines[i+numLines+1] );
}
free( lines );
for(i = 0; i < numLines; i++)
{
error = clReleaseProgram( simple_kernels[i] );
test_error( error, "Unable to release program object" );
}
free( simple_kernels );
error = clReleaseKernel( kernel );
test_error( error, "Unable to release kernel object" );
error = clReleaseProgram( my_newly_minted_library );
test_error( error, "Unable to release program object" );
error = clReleaseProgram( my_newly_linked_program );
test_error( error, "Unable to release program object" );
return 0;
}
int test_multi_file_libraries(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
{
unsigned int toTest[] = { 2, 4, 8, 16, 32, 64, 128, 256, 0 }; // 512, 1024, 2048, 4096, 8192, 16384, 32768, 0 };
unsigned int i;
log_info( "Testing multi-file libraries ...this might take awhile...\n" );
for( i = 0; toTest[ i ] != 0; i++ )
{
log_info( " %d...\n", toTest[ i ] );
#if defined(_WIN32)
clock_t start = clock();
#elif defined(__linux__) || defined(__APPLE__)
timeval time1, time2;
gettimeofday(&time1, NULL);
#endif
if( test_large_multi_file_library( context, deviceID, queue, toTest[ i ] ) != 0 )
{
log_error( "ERROR: multi-file library program test failed for %d lines! (in %s:%d)\n\n", toTest[ i ], __FILE__, __LINE__ );
return -1;
}
#if defined(_WIN32)
clock_t end = clock();
log_perf( (float)( end - start ) / (float)CLOCKS_PER_SEC, false, "clock() time in secs", "%d lines", toTest[i] );
#elif defined(__linux__) || defined(__APPLE__)
gettimeofday(&time2, NULL);
log_perf( (float)(float)(time2.tv_sec - time1.tv_sec) + 1.0e-6 * (time2.tv_usec - time1.tv_usec) , false, "wall time in secs", "%d lines", toTest[i] );
#endif
}
return 0;
}
int test_large_multiple_embedded_headers(cl_context context, cl_device_id deviceID, cl_command_queue queue, unsigned int numLines)
{
int error;
cl_program program;
cl_program *simple_kernels;
cl_program *headers;
const char **header_names;
const char **lines;
unsigned int i;
char buffer[MAX_LINE_SIZE_IN_PROGRAM];
simple_kernels = (cl_program*)malloc(numLines*sizeof(cl_program));
if (simple_kernels == NULL) {
log_error( "ERROR: Unable to allocate simple_kernels array with %d lines! (in %s:%d)\n", numLines, __FILE__, __LINE__ );
return -1;
}
headers = (cl_program*)malloc(numLines*sizeof(cl_program));
if (headers == NULL) {
log_error( "ERROR: Unable to allocate headers array with %d lines! (in %s:%d)\n", numLines, __FILE__, __LINE__ );
return -1;
}
/* First, allocate the array for our line pointers */
header_names = (const char**)malloc( numLines*sizeof( const char * ) );
if (header_names == NULL) {
log_error( "ERROR: Unable to allocate header_names array with %d lines! (in %s:%d)\n", numLines, __FILE__, __LINE__ );
return -1;
}
lines = (const char **)malloc( (2*numLines + 2)*sizeof( const char * ) );
if (lines == NULL) {
log_error( "ERROR: Unable to allocate lines array with %d lines! (in %s:%d)\n", (2*numLines + 2), __FILE__, __LINE__ );
return -1;
}
for( i = 0; i < numLines; i++)
{
sprintf(buffer, include_header_name_templates[i % 4], i);
lines[i] = _strdup(buffer);
sprintf(buffer, header_name_templates[i % 4], i);
header_names[i] = _strdup(buffer);
sprintf(buffer, composite_kernel_extern_template, i);
const char* line = _strdup(buffer);
error = create_single_kernel_helper_create_program(context, &headers[i], 1, &line);
if( headers[i] == NULL || error != CL_SUCCESS )
{
log_error( "ERROR: Unable to create a simple header program! (%s in %s:%d)\n", IGetErrorString( error ), __FILE__, __LINE__);
return -1;
}
}
/* First and last lines are easy */
lines[ numLines ] = composite_kernel_start;
lines[ 2* numLines + 1 ] = composite_kernel_end;
/* Fill the rest with templated kernels */
for( i = numLines + 1; i < 2* numLines + 1; i++ )
{
sprintf(buffer, composite_kernel_template, i - numLines - 1);
lines[ i ] = _strdup(buffer);
}
/* Try to create a program with these lines */
error = create_single_kernel_helper_create_program(context, &program, 2 * numLines + 2, lines);
if( program == NULL || error != CL_SUCCESS )
{
log_error( "ERROR: Unable to create long test program with %d lines! (%s) (in %s:%d)\n", numLines, IGetErrorString( error ), __FILE__, __LINE__ );
return -1;
}
/* Compile it */
error = clCompileProgram(program, 1, &deviceID, NULL, numLines, headers, header_names, NULL, NULL);
test_error( error, "Unable to compile a simple program" );
/* Create and compile templated kernels */
for( i = 0; i < numLines; i++)
{
sprintf(buffer, simple_kernel_template, i);
const char* kernel_source = _strdup(buffer);
error = create_single_kernel_helper_create_program(context, &simple_kernels[i], 1, &kernel_source);
if( simple_kernels[i] == NULL || error != CL_SUCCESS )
{
log_error( "ERROR: Unable to create long test program with %d lines! (%s) (in %s:%d)\n", numLines, IGetErrorString( error ), __FILE__, __LINE__ );
return -1;
}
/* Compile it */
error = clCompileProgram(simple_kernels[i], 1, &deviceID, NULL, 0, NULL, NULL, NULL, NULL);
test_error( error, "Unable to compile a simple program" );
free((void*)kernel_source);
}
/* Create library out of compiled templated kernels */
cl_program my_newly_minted_library = clLinkProgram(context, 1, &deviceID, "-create-library", numLines, simple_kernels, NULL, NULL, &error);
test_error( error, "Unable to create a multi-line library" );
/* Link the program that calls the kernels and the library that contains them */
cl_program programs[2] = { program, my_newly_minted_library };
cl_program my_newly_linked_program = clLinkProgram(context, 1, &deviceID, NULL, 2, programs, NULL, NULL, &error);
test_error( error, "Unable to link a program with a library" );
// Create the composite kernel
cl_kernel kernel = clCreateKernel(my_newly_linked_program, "CompositeKernel", &error);
test_error( error, "Unable to create a composite kernel" );
// Run the composite kernel and verify the results
error = verifyCopyBuffer(context, queue, kernel);
if (error != CL_SUCCESS)
return error;
/* All done! */
error = clReleaseProgram( program );
test_error( error, "Unable to release program object" );
for( i = 0; i < numLines; i++)
{
free( (void*)lines[i] );
free( (void*)header_names[i] );
}
for( i = numLines + 1; i < 2* numLines + 1; i++ )
{
free( (void*)lines[i] );
}
free( lines );
free( header_names );
for(i = 0; i < numLines; i++)
{
error = clReleaseProgram( simple_kernels[i] );
test_error( error, "Unable to release program object" );
error = clReleaseProgram( headers[i] );
test_error( error, "Unable to release header program object" );
}
free( simple_kernels );
free( headers );
error = clReleaseKernel( kernel );
test_error( error, "Unable to release kernel object" );
error = clReleaseProgram( my_newly_minted_library );
test_error( error, "Unable to release program object" );
error = clReleaseProgram( my_newly_linked_program );
test_error( error, "Unable to release program object" );
return 0;
}
int test_multiple_embedded_headers(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
{
unsigned int toTest[] = { 2, 4, 8, 16, 32, 64, 128, 256, 0 }; // 512, 1024, 2048, 4096, 8192, 16384, 32768, 0 };
unsigned int i;
log_info( "Testing multiple embedded headers ...this might take awhile...\n" );
for( i = 0; toTest[ i ] != 0; i++ )
{
log_info( " %d...\n", toTest[ i ] );
#if defined(_WIN32)
clock_t start = clock();
#elif defined(__linux__) || defined(__APPLE__)
timeval time1, time2;
gettimeofday(&time1, NULL);
#endif
if( test_large_multiple_embedded_headers( context, deviceID, queue, toTest[ i ] ) != 0 )
{
log_error( "ERROR: multiple embedded headers program test failed for %d lines! (in %s:%d)\n", toTest[ i ], __FILE__, __LINE__ );
return -1;
}
#if defined(_WIN32)
clock_t end = clock();
log_perf( (float)( end - start ) / (float)CLOCKS_PER_SEC, false, "clock() time in secs", "%d lines", toTest[i] );
#elif defined(__linux__) || defined(__APPLE__)
gettimeofday(&time2, NULL);
log_perf( (float)(float)(time2.tv_sec - time1.tv_sec) + 1.0e-6 * (time2.tv_usec - time1.tv_usec) , false, "wall time in secs", "%d lines", toTest[i] );
#endif
}
return 0;
}
double logbase(double a, double base)
{
return log(a) / log(base);
}
int test_large_multiple_libraries(cl_context context, cl_device_id deviceID, cl_command_queue queue, unsigned int numLines)
{
int error;
cl_program *simple_kernels;
const char **lines;
unsigned int i;
char buffer[MAX_LINE_SIZE_IN_PROGRAM];
/* I want to create (log2(N)+1)/2 libraries */
unsigned int level = (unsigned int)(logbase(numLines, 2.0) + 1.000001)/2;
unsigned int numLibraries = (unsigned int)pow(2.0, level - 1.0);
unsigned int numFilesInLib = numLines/numLibraries;
cl_program *my_program_and_libraries = (cl_program*)malloc((1+numLibraries)*sizeof(cl_program));
if (my_program_and_libraries == NULL) {
log_error( "ERROR: Unable to allocate program array with %d programs! (in %s:%d)\n", (1+numLibraries), __FILE__, __LINE__);
return -1;
}
log_info("level - %d, numLibraries - %d, numFilesInLib - %d\n", level, numLibraries, numFilesInLib);
simple_kernels = (cl_program*)malloc(numLines*sizeof(cl_program));
if (simple_kernels == NULL) {
log_error( "ERROR: Unable to allocate kernels array with %d kernels! (in %s:%d)\n", numLines, __FILE__, __LINE__);
return -1;
}
/* First, allocate the array for our line pointers */
lines = (const char **)malloc( (2*numLines + 2) * sizeof( const char * ) );
if (lines == NULL) {
log_error( "ERROR: Unable to allocate lines array with %d lines! (in %s:%d)\n", (2*numLines + 2), __FILE__, __LINE__);
return -1;
}
for(i = 0; i < numLines; i++)
{
sprintf(buffer, composite_kernel_extern_template, i);
lines[i] = _strdup(buffer);
}
/* First and last lines are easy */
lines[ numLines ] = composite_kernel_start;
lines[ 2*numLines + 1] = composite_kernel_end;
/* Fill the rest with templated kernels */
for(i = numLines + 1; i < 2*numLines + 1; i++ )
{
sprintf(buffer, composite_kernel_template, i - numLines - 1);
lines[ i ] = _strdup(buffer);
}
/* Try to create a program with these lines */
error = create_single_kernel_helper_create_program(context, &my_program_and_libraries[0], 2 * numLines + 2, lines);
if( my_program_and_libraries[0] == NULL || error != CL_SUCCESS )
{
log_error( "ERROR: Unable to create long test program with %d lines! (%s in %s:%d)\n", numLines, IGetErrorString( error ), __FILE__, __LINE__ );
return -1;
}
/* Compile it */
error = clCompileProgram(my_program_and_libraries[0], 1, &deviceID, NULL, 0, NULL, NULL, NULL, NULL);
test_error( error, "Unable to compile a simple program" );
/* Create and compile templated kernels */
for(i = 0; i < numLines; i++)
{
sprintf(buffer, simple_kernel_template, i);
const char* kernel_source = _strdup(buffer);
error = create_single_kernel_helper_create_program(context, &simple_kernels[i], 1, &kernel_source);
if( simple_kernels[i] == NULL || error != CL_SUCCESS )
{
log_error( "ERROR: Unable to create long test program with %d lines! (%s in %s:%d)\n", numLines, IGetErrorString( error ), __FILE__, __LINE__ );
return -1;
}
/* Compile it */
error = clCompileProgram(simple_kernels[i], 1, &deviceID, NULL, 0, NULL, NULL, NULL, NULL);
test_error( error, "Unable to compile a simple program" );
free((void*)kernel_source);
}
/* Create library out of compiled templated kernels */
for(i = 0; i < numLibraries; i++) {
my_program_and_libraries[i+1] = clLinkProgram(context, 1, &deviceID, "-create-library", numFilesInLib, simple_kernels+i*numFilesInLib, NULL, NULL, &error);
test_error( error, "Unable to create a multi-line library" );
}
/* Link the program that calls the kernels and the library that contains them */
cl_program my_newly_linked_program = clLinkProgram(context, 1, &deviceID, NULL, numLibraries+1, my_program_and_libraries, NULL, NULL, &error);
test_error( error, "Unable to link a program with a library" );
// Create the composite kernel
cl_kernel kernel = clCreateKernel(my_newly_linked_program, "CompositeKernel", &error);
test_error( error, "Unable to create a composite kernel" );
// Run the composite kernel and verify the results
error = verifyCopyBuffer(context, queue, kernel);
if (error != CL_SUCCESS)
return error;
/* All done! */
for(i = 0; i <= numLibraries; i++) {
error = clReleaseProgram( my_program_and_libraries[i] );
test_error( error, "Unable to release program object" );
}
free( my_program_and_libraries );
for(i = 0; i < numLines; i++)
{
free( (void*)lines[i] );
}
for(i = numLines + 1; i < 2*numLines + 1; i++ )
{
free( (void*)lines[i] );
}
free( lines );
for(i = 0; i < numLines; i++)
{
error = clReleaseProgram( simple_kernels[i] );
test_error( error, "Unable to release program object" );
}
free( simple_kernels );
error = clReleaseKernel( kernel );
test_error( error, "Unable to release kernel object" );
error = clReleaseProgram( my_newly_linked_program );
test_error( error, "Unable to release program object" );
return 0;
}
int test_multiple_libraries(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
{
unsigned int toTest[] = { 2, 8, 32, 128, 256, 0 }; // 512, 2048, 8192, 32768, 0 };
unsigned int i;
log_info( "Testing multiple libraries ...this might take awhile...\n" );
for( i = 0; toTest[ i ] != 0; i++ )
{
log_info( " %d...\n", toTest[ i ] );
#if defined(_WIN32)
clock_t start = clock();
#elif defined(__linux__) || defined(__APPLE__)
timeval time1, time2;
gettimeofday(&time1, NULL);
#endif
if( test_large_multiple_libraries( context, deviceID, queue, toTest[ i ] ) != 0 )
{
log_error( "ERROR: multiple library program test failed for %d lines! (in %s:%d)\n\n", toTest[ i ], __FILE__, __LINE__ );
return -1;
}
#if defined(_WIN32)
clock_t end = clock();
log_perf( (float)( end - start ) / (float)CLOCKS_PER_SEC, false, "clock() time in secs", "%d lines", toTest[i] );
#elif defined(__linux__) || defined(__APPLE__)
gettimeofday(&time2, NULL);
log_perf( (float)(float)(time2.tv_sec - time1.tv_sec) + 1.0e-6 * (time2.tv_usec - time1.tv_usec) , false, "wall time in secs", "%d lines", toTest[i] );
#endif
}
return 0;
}
int test_large_multiple_files_multiple_libraries(cl_context context, cl_device_id deviceID, cl_command_queue queue, unsigned int numLines)
{
int error;
cl_program *simple_kernels;
const char **lines;
unsigned int i;
char buffer[MAX_LINE_SIZE_IN_PROGRAM];
/* I want to create (log2(N)+1)/4 libraries */
unsigned int level = (unsigned int)(logbase(numLines, 2.0) + 1.000001)/2;
unsigned int numLibraries = (unsigned int)pow(2.0, level - 2.0);
unsigned int numFilesInLib = numLines/(2*numLibraries);
cl_program *my_programs_and_libraries = (cl_program*)malloc((1+numLibraries+numLibraries*numFilesInLib)*sizeof(cl_program));
if (my_programs_and_libraries == NULL) {
log_error( "ERROR: Unable to allocate program array with %d programs! (in %s:%d)\n", (1+numLibraries+numLibraries*numFilesInLib), __FILE__, __LINE__ );
return -1;
}
log_info("level - %d, numLibraries - %d, numFilesInLib - %d\n", level, numLibraries, numFilesInLib);
simple_kernels = (cl_program*)malloc(numLines*sizeof(cl_program));
if (simple_kernels == NULL) {
log_error( "ERROR: Unable to allocate kernels array with %d kernels! (in %s:%d)\n", numLines, __FILE__, __LINE__ );
return -1;
}
/* First, allocate the array for our line pointers */
lines = (const char **)malloc( (2*numLines + 2) * sizeof( const char * ) );
if (lines == NULL) {
log_error( "ERROR: Unable to allocate lines array with %d lines! (in %s:%d)\n", (2*numLines + 2), __FILE__, __LINE__ );
return -1;
}
for(i = 0; i < numLines; i++)
{
sprintf(buffer, composite_kernel_extern_template, i);
lines[i] = _strdup(buffer);
}
/* First and last lines are easy */
lines[ numLines ] = composite_kernel_start;
lines[ 2*numLines + 1] = composite_kernel_end;
/* Fill the rest with templated kernels */
for(i = numLines + 1; i < 2*numLines + 1; i++ )
{
sprintf(buffer, composite_kernel_template, i - numLines - 1);
lines[ i ] = _strdup(buffer);
}
/* Try to create a program with these lines */
error = create_single_kernel_helper_create_program(context, &my_programs_and_libraries[0], 2 * numLines + 2, lines);
if( my_programs_and_libraries[0] == NULL || error != CL_SUCCESS )
{
log_error( "ERROR: Unable to create long test program with %d lines! (%s in %s:%d)\n", numLines, IGetErrorString( error ), __FILE__, __LINE__ );
return -1;
}
/* Compile it */
error = clCompileProgram(my_programs_and_libraries[0], 1, &deviceID, NULL, 0, NULL, NULL, NULL, NULL);
test_error( error, "Unable to compile a simple program" );
/* Create and compile templated kernels */
for(i = 0; i < numLines; i++)
{
sprintf(buffer, simple_kernel_template, i);
const char* kernel_source = _strdup(buffer);
error = create_single_kernel_helper_create_program(context, &simple_kernels[i], 1, &kernel_source);
if( simple_kernels[i] == NULL || error != CL_SUCCESS )
{
log_error( "ERROR: Unable to create long test program with %d lines! (%s in %s:%d)\n", numLines, IGetErrorString( error ), __FILE__, __LINE__ );
return -1;
}
/* Compile it */
error = clCompileProgram(simple_kernels[i], 1, &deviceID, NULL, 0, NULL, NULL, NULL, NULL);
test_error( error, "Unable to compile a simple program" );
free((void*)kernel_source);
}
/* Copy already compiled kernels */
for( i = 0; i < numLibraries*numFilesInLib; i++) {
my_programs_and_libraries[i+1] = simple_kernels[i];
}
/* Create library out of compiled templated kernels */
for( i = 0; i < numLibraries; i++) {
my_programs_and_libraries[i+1+numLibraries*numFilesInLib] = clLinkProgram(context, 1, &deviceID, "-create-library", numFilesInLib, simple_kernels+(i*numFilesInLib+numLibraries*numFilesInLib), NULL, NULL, &error);
test_error( error, "Unable to create a multi-line library" );
}
/* Link the program that calls the kernels and the library that contains them */
cl_program my_newly_linked_program = clLinkProgram(context, 1, &deviceID, NULL, numLibraries+1+numLibraries*numFilesInLib, my_programs_and_libraries, NULL, NULL, &error);
test_error( error, "Unable to link a program with a library" );
// Create the composite kernel
cl_kernel kernel = clCreateKernel(my_newly_linked_program, "CompositeKernel", &error);
test_error( error, "Unable to create a composite kernel" );
// Run the composite kernel and verify the results
error = verifyCopyBuffer(context, queue, kernel);
if (error != CL_SUCCESS)
return error;
/* All done! */
for(i = 0; i < numLibraries+1+numLibraries*numFilesInLib; i++) {
error = clReleaseProgram( my_programs_and_libraries[i] );
test_error( error, "Unable to release program object" );
}
free( my_programs_and_libraries );
for(i = 0; i < numLines; i++)
{
free( (void*)lines[i] );
}
for(i = numLines + 1; i < 2*numLines + 1; i++ )
{
free( (void*)lines[i] );
}
free( lines );
for(i = numLibraries*numFilesInLib; i < numLines; i++)
{
error = clReleaseProgram( simple_kernels[i] );
test_error( error, "Unable to release program object" );
}
free( simple_kernels );
error = clReleaseKernel( kernel );
test_error( error, "Unable to release kernel object" );
error = clReleaseProgram( my_newly_linked_program );
test_error( error, "Unable to release program object" );
return 0;
}
int test_multiple_files_multiple_libraries(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
{
unsigned int toTest[] = { 8, 32, 128, 256, 0 }; // 512, 2048, 8192, 32768, 0 };
unsigned int i;
log_info( "Testing multiple files and multiple libraries ...this might take awhile...\n" );
for( i = 0; toTest[ i ] != 0; i++ )
{
log_info( " %d...\n", toTest[ i ] );
#if defined(_WIN32)
clock_t start = clock();
#elif defined(__linux__) || defined(__APPLE__)
timeval time1, time2;
gettimeofday(&time1, NULL);
#endif
if( test_large_multiple_files_multiple_libraries( context, deviceID, queue, toTest[ i ] ) != 0 )
{
log_error( "ERROR: multiple files, multiple libraries program test failed for %d lines! (in %s:%d)\n\n", toTest[ i ], __FILE__, __LINE__ );
return -1;
}
#if defined(_WIN32)
clock_t end = clock();
log_perf( (float)( end - start ) / (float)CLOCKS_PER_SEC, false, "clock() time in secs", "%d lines", toTest[i] );
#elif defined(__linux__) || defined(__APPLE__)
gettimeofday(&time2, NULL);
log_perf( (float)(float)(time2.tv_sec - time1.tv_sec) + 1.0e-6 * (time2.tv_usec - time1.tv_usec) , false, "wall time in secs", "%d lines", toTest[i] );
#endif
}
return 0;
}
int test_large_multiple_files(cl_context context, cl_device_id deviceID, cl_command_queue queue, unsigned int numLines)
{
int error;
const char **lines;
unsigned int i;
char buffer[MAX_LINE_SIZE_IN_PROGRAM];
cl_program *my_programs = (cl_program*)malloc((1+numLines)*sizeof(cl_program));
if (my_programs == NULL) {
log_error( "ERROR: Unable to allocate my_programs array with %d programs! (in %s:%d)\n", (1+numLines), __FILE__, __LINE__);
return -1;
}
/* First, allocate the array for our line pointers */
lines = (const char **)malloc( (2*numLines + 2) * sizeof( const char * ) );
if (lines == NULL) {
log_error( "ERROR: Unable to allocate lines array with %d lines! (in %s:%d)\n", (2*numLines + 2), __FILE__, __LINE__);
return -1;
}
for(i = 0; i < numLines; i++)
{
sprintf(buffer, composite_kernel_extern_template, i);
lines[i] = _strdup(buffer);
}
/* First and last lines are easy */
lines[ numLines ] = composite_kernel_start;
lines[ 2* numLines + 1] = composite_kernel_end;
/* Fill the rest with templated kernels */
for(i = numLines + 1; i < 2*numLines + 1; i++ )
{
sprintf(buffer, composite_kernel_template, i - numLines - 1);
lines[ i ] = _strdup(buffer);
}
/* Try to create a program with these lines */
error = create_single_kernel_helper_create_program(context, &my_programs[0], 2 * numLines + 2, lines);
if( my_programs[0] == NULL || error != CL_SUCCESS )
{
log_error( "ERROR: Unable to create long test program with %d lines! (%s in %s:%d)\n", numLines, IGetErrorString( error ), __FILE__, __LINE__ );
return -1;
}
/* Compile it */
error = clCompileProgram(my_programs[0], 1, &deviceID, NULL, 0, NULL, NULL, NULL, NULL);
test_error( error, "Unable to compile a simple program" );
/* Create and compile templated kernels */
for( i = 0; i < numLines; i++)
{
sprintf(buffer, simple_kernel_template, i);
const char* kernel_source = _strdup(buffer);
error = create_single_kernel_helper_create_program(context, &my_programs[i + 1], 1, &kernel_source);
if( my_programs[i+1] == NULL || error != CL_SUCCESS )
{
log_error( "ERROR: Unable to create long test program with %d lines! (%s in %s:%d)\n", numLines, IGetErrorString( error ), __FILE__, __LINE__ );
return -1;
}
/* Compile it */
error = clCompileProgram(my_programs[i+1], 1, &deviceID, NULL, 0, NULL, NULL, NULL, NULL);
test_error( error, "Unable to compile a simple program" );
free((void*)kernel_source);
}
/* Link the program that calls the kernels and the library that contains them */
cl_program my_newly_linked_program = clLinkProgram(context, 1, &deviceID, NULL, 1+numLines, my_programs, NULL, NULL, &error);
test_error( error, "Unable to link a program with a library" );
// Create the composite kernel
cl_kernel kernel = clCreateKernel(my_newly_linked_program, "CompositeKernel", &error);
test_error( error, "Unable to create a composite kernel" );
// Run the composite kernel and verify the results
error = verifyCopyBuffer(context, queue, kernel);
if (error != CL_SUCCESS)
return error;
/* All done! */
for(i = 0; i < 1+numLines; i++) {
error = clReleaseProgram( my_programs[i] );
test_error( error, "Unable to release program object" );
}
free( my_programs );
for(i = 0; i < numLines; i++)
{
free( (void*)lines[i] );
}
for(i = numLines + 1; i < 2*numLines + 1; i++ )
{
free( (void*)lines[i] );
}
free( lines );
error = clReleaseKernel( kernel );
test_error( error, "Unable to release kernel object" );
error = clReleaseProgram( my_newly_linked_program );
test_error( error, "Unable to release program object" );
return 0;
}
int test_multiple_files(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
{
unsigned int toTest[] = { 8, 32, 128, 256, 0 }; // 512, 2048, 8192, 32768, 0 };
unsigned int i;
log_info( "Testing multiple files compilation and linking into a single executable ...this might take awhile...\n" );
for( i = 0; toTest[ i ] != 0; i++ )
{
log_info( " %d...\n", toTest[ i ] );
#if defined(_WIN32)
clock_t start = clock();
#elif defined(__linux__) || defined(__APPLE__)
timeval time1, time2;
gettimeofday(&time1, NULL);
#endif
if( test_large_multiple_files( context, deviceID, queue, toTest[ i ] ) != 0 )
{
log_error( "ERROR: multiple files program test failed for %d lines! (in %s:%d)\n\n", toTest[ i ], __FILE__, __LINE__ );
return -1;
}
#if defined(_WIN32)
clock_t end = clock();
log_perf( (float)( end - start ) / (float)CLOCKS_PER_SEC, false, "clock() time in secs", "%d lines", toTest[i] );
#elif defined(__linux__) || defined(__APPLE__)
gettimeofday(&time2, NULL);
log_perf( (float)(float)(time2.tv_sec - time1.tv_sec) + 1.0e-6 * (time2.tv_usec - time1.tv_usec) , false, "wall time in secs", "%d lines", toTest[i] );
#endif
}
return 0;
}
int test_simple_compile_only(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
{
int error;
cl_program program;
log_info("Testing a simple compilation only...\n");
error = create_single_kernel_helper_create_program(context, &program, 1, &simple_kernel);
if( program == NULL || error != CL_SUCCESS )
{
log_error( "ERROR: Unable to create a simple test program! (%s in %s:%d)\n", IGetErrorString( error ), __FILE__, __LINE__ );
return -1;
}
error = clCompileProgram(program, 1, &deviceID, NULL, 0, NULL, NULL, NULL, NULL);
test_error( error, "Unable to compile a simple program" );
/* All done! */
error = clReleaseProgram( program );
test_error( error, "Unable to release program object" );
return 0;
}
int test_simple_static_compile_only(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
{
int error;
cl_program program;
log_info("Testing a simple static compilations only...\n");
error = create_single_kernel_helper_create_program(context, &program, 1, &compile_static_var);
if( program == NULL || error != CL_SUCCESS )
{
log_error( "ERROR: Unable to create a simple static variable test program! (%s in %s:%d)\n", IGetErrorString( error ), __FILE__, __LINE__ );
return -1;
}
log_info("Compiling a static variable...\n");
error = clCompileProgram(program, 1, &deviceID, NULL, 0, NULL, NULL, NULL, NULL);
test_error( error, "Unable to compile a simple static variable program" );
/* All done! */
error = clReleaseProgram( program );
test_error( error, "Unable to release program object" );
error = create_single_kernel_helper_create_program(context, &program, 1, &compile_static_struct);
if( program == NULL || error != CL_SUCCESS )
{
log_error( "ERROR: Unable to create a simple static struct test program! (%s in %s:%d)\n", IGetErrorString( error ), __FILE__, __LINE__ );
return -1;
}
log_info("Compiling a static struct...\n");
error = clCompileProgram(program, 1, &deviceID, NULL, 0, NULL, NULL, NULL, NULL);
test_error( error, "Unable to compile a simple static variable program" );
/* All done! */
error = clReleaseProgram( program );
test_error( error, "Unable to release program object" );
error = create_single_kernel_helper_create_program(context, &program, 1, &compile_static_function);
if( program == NULL || error != CL_SUCCESS )
{
log_error( "ERROR: Unable to create a simple static function test program! (%s in %s:%d)\n", IGetErrorString( error ), __FILE__, __LINE__ );
return -1;
}
log_info("Compiling a static function...\n");
error = clCompileProgram(program, 1, &deviceID, NULL, 0, NULL, NULL, NULL, NULL);
test_error( error, "Unable to compile a simple static function program" );
/* All done! */
error = clReleaseProgram( program );
test_error( error, "Unable to release program object" );
return 0;
}
int test_simple_extern_compile_only(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
{
int error;
cl_program program;
log_info("Testing a simple extern compilations only...\n");
error = create_single_kernel_helper_create_program(context, &program, 1, &simple_header);
if( program == NULL || error != CL_SUCCESS )
{
log_error( "ERROR: Unable to create a simple extern kernel test program! (%s in %s:%d)\n", IGetErrorString( error ), __FILE__, __LINE__ );
return -1;
}
log_info("Compiling an extern kernel...\n");
error = clCompileProgram(program, 1, &deviceID, NULL, 0, NULL, NULL, NULL, NULL);
test_error( error, "Unable to compile a simple extern kernel program" );
/* All done! */
error = clReleaseProgram( program );
test_error( error, "Unable to release program object" );
error = create_single_kernel_helper_create_program(context, &program, 1, &compile_extern_var);
if( program == NULL || error != CL_SUCCESS )
{
log_error( "ERROR: Unable to create a simple extern variable test program! (%s in %s:%d)\n", IGetErrorString( error ), __FILE__, __LINE__ );
return -1;
}
log_info("Compiling an extern variable...\n");
error = clCompileProgram(program, 1, &deviceID, NULL, 0, NULL, NULL, NULL, NULL);
test_error( error, "Unable to compile a simple extern variable program" );
/* All done! */
error = clReleaseProgram( program );
test_error( error, "Unable to release program object" );
error = create_single_kernel_helper_create_program(context, &program, 1, &compile_extern_struct);
if( program == NULL || error != CL_SUCCESS )
{
log_error( "ERROR: Unable to create a simple extern struct test program! (%s in %s:%d)\n", IGetErrorString( error ), __FILE__, __LINE__ );
return -1;
}
log_info("Compiling an extern struct...\n");
error = clCompileProgram(program, 1, &deviceID, NULL, 0, NULL, NULL, NULL, NULL);
test_error( error, "Unable to compile a simple extern variable program" );
/* All done! */
error = clReleaseProgram( program );
test_error( error, "Unable to release program object" );
error = create_single_kernel_helper_create_program(context, &program, 1, &compile_extern_function);
if( program == NULL || error != CL_SUCCESS )
{
log_error( "ERROR: Unable to create a simple extern function test program! (%s in %s:%d)\n", IGetErrorString( error ), __FILE__, __LINE__ );
return -1;
}
log_info("Compiling an extern function...\n");
error = clCompileProgram(program, 1, &deviceID, NULL, 0, NULL, NULL, NULL, NULL);
test_error( error, "Unable to compile a simple extern function program" );
/* All done! */
error = clReleaseProgram( program );
test_error( error, "Unable to release program object" );
return 0;
}
struct simple_user_data {
const char* m_message;
cl_event m_event;
};
const char* once_upon_a_midnight_dreary = "Once upon a midnight dreary!";
static void CL_CALLBACK simple_compile_callback(cl_program program, void* user_data)
{
simple_user_data* simple_compile_user_data = (simple_user_data*)user_data;
log_info("in the simple_compile_callback: program %p just completed compiling with '%s'\n", program, simple_compile_user_data->m_message);
if (strcmp(once_upon_a_midnight_dreary, simple_compile_user_data->m_message) != 0)
{
log_error("ERROR: in the simple_compile_callback: Expected '%s' and got %s (in %s:%d)!\n", once_upon_a_midnight_dreary, simple_compile_user_data->m_message, __FILE__, __LINE__);
}
int error;
log_info("in the simple_compile_callback: program %p just completed compiling with '%p'\n", program, simple_compile_user_data->m_event);
error = clSetUserEventStatus(simple_compile_user_data->m_event, CL_COMPLETE);
if (error != CL_SUCCESS)
{
log_error( "ERROR: in the simple_compile_callback: Unable to set user event status to CL_COMPLETE! (%s in %s:%d)\n", IGetErrorString( error ), __FILE__, __LINE__ );
exit(-1);
}
log_info("in the simple_compile_callback: Successfully signaled compile_program_completion_event!\n");
}
int test_simple_compile_with_callback(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
{
int error;
cl_program program;
cl_event compile_program_completion_event;
log_info("Testing a simple compilation with callback...\n");
error = create_single_kernel_helper_create_program(context, &program, 1, &simple_kernel);
if( program == NULL || error != CL_SUCCESS )
{
log_error( "ERROR: Unable to create a simple test program! (%s in %s:%d)\n", IGetErrorString( error ), __FILE__, __LINE__ );
return -1;
}
compile_program_completion_event = clCreateUserEvent(context, &error);
test_error( error, "Unable to create a user event");
simple_user_data simple_compile_user_data = {once_upon_a_midnight_dreary, compile_program_completion_event};
error = clCompileProgram(program, 1, &deviceID, NULL, 0, NULL, NULL, simple_compile_callback, (void*)&simple_compile_user_data);
test_error( error, "Unable to compile a simple program with a callback" );
error = clWaitForEvents(1, &compile_program_completion_event);
test_error( error, "clWaitForEvents failed when waiting on compile_program_completion_event");
/* All done! */
error = clReleaseEvent(compile_program_completion_event);
test_error( error, "Unable to release event object" );
error = clReleaseProgram( program );
test_error( error, "Unable to release program object" );
return 0;
}
int test_simple_embedded_header_compile(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
{
int error;
cl_program program, header;
log_info("Testing a simple embedded header compile only...\n");
program = clCreateProgramWithSource(context, 1, &another_simple_kernel_with_header, NULL, &error);
if( program == NULL || error != CL_SUCCESS )
{
log_error( "ERROR: Unable to create a simple test program! (%s in %s:%d)\n", IGetErrorString( error ), __FILE__, __LINE__ );
return -1;
}
header = clCreateProgramWithSource(context, 1, &simple_header, NULL, &error);
if( header == NULL || error != CL_SUCCESS )
{
log_error( "ERROR: Unable to create a simple header program! (%s in %s:%d)\n", IGetErrorString( error ), __FILE__, __LINE__ );
return -1;
}
error = clCompileProgram(program, 1, &deviceID, NULL, 1, &header, &simple_header_name, NULL, NULL);
test_error( error, "Unable to compile a simple program with embedded header" );
/* All done! */
error = clReleaseProgram( program );
test_error( error, "Unable to release program object" );
error = clReleaseProgram( header );
test_error( error, "Unable to release program object" );
return 0;
}
int test_simple_link_only(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
{
int error;
cl_program program;
log_info("Testing a simple linking only...\n");
error = create_single_kernel_helper_create_program(context, &program, 1, &simple_kernel);
if( program == NULL || error != CL_SUCCESS )
{
log_error( "ERROR: Unable to create a simple test program! (%s in %s:%d)\n", IGetErrorString( error ), __FILE__, __LINE__ );
return -1;
}
error = clCompileProgram(program, 1, &deviceID, NULL, 0, NULL, NULL, NULL, NULL);
test_error( error, "Unable to compile a simple program" );
cl_program my_newly_linked_program = clLinkProgram(context, 1, &deviceID, NULL, 1, &program, NULL, NULL, &error);
test_error( error, "Unable to link a simple program" );
/* All done! */
error = clReleaseProgram( program );
test_error( error, "Unable to release program object" );
error = clReleaseProgram( my_newly_linked_program );
test_error( error, "Unable to release program object" );
return 0;
}
int test_two_file_regular_variable_access(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
{
int error;
cl_program program, second_program, my_newly_linked_program;
const char* sources[2] = {simple_kernel, compile_regular_var}; // here we want to avoid linking error due to lack of kernels
log_info("Compiling and linking two program objects, where one tries to access regular variable from another...\n");
error = create_single_kernel_helper_create_program(context, &program, 2, sources);
if( program == NULL || error != CL_SUCCESS )
{
log_error( "ERROR: Unable to create a test program with regular variable! (%s in %s:%d)\n", IGetErrorString( error ), __FILE__, __LINE__ );
return -1;
}
error = clCompileProgram(program, 1, &deviceID, NULL, 0, NULL, NULL, NULL, NULL);
test_error( error, "Unable to compile a simple program with regular function" );
error = create_single_kernel_helper_create_program(context, &second_program, 1, &link_static_var_access);
if( program == NULL || error != CL_SUCCESS )
{
log_error( "ERROR: Unable to create a test program that tries to access a regular variable! (%s in %s:%d)\n", IGetErrorString( error ), __FILE__, __LINE__ );
return -1;
}
error = clCompileProgram(second_program, 1, &deviceID, NULL, 0, NULL, NULL, NULL, NULL);
test_error( error, "Unable to compile a program that tries to access a regular variable" );
cl_program two_programs[2] = { program, second_program };
my_newly_linked_program = clLinkProgram(context, 1, &deviceID, NULL, 2, two_programs, NULL, NULL, &error);
test_error( error, "clLinkProgram: Expected a different error code while linking a program that tries to access a regular variable" );
/* All done! */
error = clReleaseProgram( program );
test_error( error, "Unable to release program object" );
error = clReleaseProgram( second_program );
test_error( error, "Unable to release program object" );
error = clReleaseProgram( my_newly_linked_program );
test_error( error, "Unable to release program object" );
return 0;
}
int test_two_file_regular_struct_access(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
{
int error;
cl_program program, second_program, my_newly_linked_program;
const char* sources[2] = {simple_kernel, compile_regular_struct}; // here we want to avoid linking error due to lack of kernels
log_info("Compiling and linking two program objects, where one tries to access regular struct from another...\n");
error = create_single_kernel_helper_create_program(context, &program, 2, sources);
if( program == NULL || error != CL_SUCCESS )
{
log_error( "ERROR: Unable to create a test program with regular struct! (%s in %s:%d)\n", IGetErrorString( error ), __FILE__, __LINE__ );
return -1;
}
error = clCompileProgram(program, 1, &deviceID, NULL, 0, NULL, NULL, NULL, NULL);
test_error( error, "Unable to compile a simple program with regular struct" );
error = create_single_kernel_helper_create_program(context, &second_program, 1, &link_static_struct_access);
if( second_program == NULL || error != CL_SUCCESS )
{
log_error( "ERROR: Unable to create a test program that tries to access a regular struct! (%s in %s:%d)\n", IGetErrorString( error ), __FILE__, __LINE__ );
return -1;
}
error = clCompileProgram(second_program, 1, &deviceID, NULL, 0, NULL, NULL, NULL, NULL);
test_error( error, "Unable to compile a program that tries to access a regular struct" );
cl_program two_programs[2] = { program, second_program };
my_newly_linked_program = clLinkProgram(context, 1, &deviceID, NULL, 2, two_programs, NULL, NULL, &error);
test_error( error, "clLinkProgram: Expected a different error code while linking a program that tries to access a regular struct" );
/* All done! */
error = clReleaseProgram( program );
test_error( error, "Unable to release program object" );
error = clReleaseProgram( second_program );
test_error( error, "Unable to release program object" );
error = clReleaseProgram( my_newly_linked_program );
test_error( error, "Unable to release program object" );
return 0;
}
int test_two_file_regular_function_access(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
{
int error;
cl_program program, second_program, my_newly_linked_program;
const char* sources[2] = {simple_kernel, compile_regular_function}; // here we want to avoid linking error due to lack of kernels
log_info("Compiling and linking two program objects, where one tries to access regular function from another...\n");
error = create_single_kernel_helper_create_program(context, &program, 2, sources);
if( program == NULL || error != CL_SUCCESS )
{
log_error( "ERROR: Unable to create a test program with regular function! (%s in %s:%d)\n", IGetErrorString( error ), __FILE__, __LINE__ );
return -1;
}
error = clCompileProgram(program, 1, &deviceID, NULL, 0, NULL, NULL, NULL, NULL);
test_error( error, "Unable to compile a simple program with regular function" );
error = create_single_kernel_helper_create_program(context, &second_program, 1, &link_static_function_access);
if( second_program == NULL || error != CL_SUCCESS )
{
log_error( "ERROR: Unable to create a test program that tries to access a regular function! (%s in %s:%d)\n", IGetErrorString( error ), __FILE__, __LINE__ );
return -1;
}
error = clCompileProgram(second_program, 1, &deviceID, NULL, 0, NULL, NULL, NULL, NULL);
test_error( error, "Unable to compile a program that tries to access a regular function" );
cl_program two_programs[2] = { program, second_program };
my_newly_linked_program = clLinkProgram(context, 1, &deviceID, NULL, 2, two_programs, NULL, NULL, &error);
test_error( error, "clLinkProgram: Expected a different error code while linking a program that tries to access a regular function" );
/* All done! */
error = clReleaseProgram( program );
test_error( error, "Unable to release program object" );
error = clReleaseProgram( second_program );
test_error( error, "Unable to release program object" );
error = clReleaseProgram( my_newly_linked_program );
test_error( error, "Unable to release program object" );
return 0;
}
int test_simple_embedded_header_link(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
{
int error;
cl_program program, header, simple_program;
log_info("Testing a simple embedded header link...\n");
program = clCreateProgramWithSource(context, 1, &another_simple_kernel_with_header, NULL, &error);
if( program == NULL || error != CL_SUCCESS )
{
log_error( "ERROR: Unable to create a simple test program! (%s in %s:%d)\n", IGetErrorString( error ), __FILE__, __LINE__ );
return -1;
}
header = clCreateProgramWithSource(context, 1, &simple_header, NULL, &error);
if( header == NULL || error != CL_SUCCESS )
{
log_error( "ERROR: Unable to create a simple header program! (%s in %s:%d)\n", IGetErrorString( error ), __FILE__, __LINE__ );
return -1;
}
error = clCompileProgram(program, 1, &deviceID, NULL, 1, &header, &simple_header_name, NULL, NULL);
test_error( error, "Unable to compile a simple program with embedded header" );
error = create_single_kernel_helper_create_program(context, &simple_program, 1, &simple_kernel);
if( simple_program == NULL || error != CL_SUCCESS )
{
log_error( "ERROR: Unable to create a simple test program! (%s in %s:%d)\n", IGetErrorString( error ), __FILE__, __LINE__ );
return -1;
}
error = clCompileProgram(simple_program, 1, &deviceID, NULL, 0, NULL, NULL, NULL, NULL);
test_error( error, "Unable to compile a simple program" );
cl_program two_programs[2] = { program, simple_program };
cl_program fully_linked_program = clLinkProgram(context, 1, &deviceID, "", 2, two_programs, NULL, NULL, &error);
test_error( error, "Unable to create an executable from two binaries, one compiled with embedded header" );
/* All done! */
error = clReleaseProgram( program );
test_error( error, "Unable to release program object" );
error = clReleaseProgram( header );
test_error( error, "Unable to release program object" );
error = clReleaseProgram( simple_program );
test_error( error, "Unable to release program object" );
error = clReleaseProgram( fully_linked_program );
test_error( error, "Unable to release program object" );
return 0;
}
const char* when_i_pondered_weak_and_weary = "When I pondered weak and weary!";
static void CL_CALLBACK simple_link_callback(cl_program program, void* user_data)
{
simple_user_data* simple_link_user_data = (simple_user_data*)user_data;
log_info("in the simple_link_callback: program %p just completed linking with '%s'\n", program, (const char*)simple_link_user_data->m_message);
if (strcmp(when_i_pondered_weak_and_weary, simple_link_user_data->m_message) != 0)
{
log_error("ERROR: in the simple_compile_callback: Expected '%s' and got %s! (in %s:%d)\n", when_i_pondered_weak_and_weary, simple_link_user_data->m_message, __FILE__, __LINE__);
}
int error;
log_info("in the simple_link_callback: program %p just completed linking with '%p'\n", program, simple_link_user_data->m_event);
error = clSetUserEventStatus(simple_link_user_data->m_event, CL_COMPLETE);
if (error != CL_SUCCESS)
{
log_error( "ERROR: simple_link_callback: Unable to set user event status to CL_COMPLETE! (%s in %s:%d)\n", IGetErrorString( error ), __FILE__, __LINE__ );
exit(-1);
}
log_info("in the simple_link_callback: Successfully signaled link_program_completion_event event!\n");
}
int test_simple_link_with_callback(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
{
int error;
cl_program program;
cl_event link_program_completion_event;
log_info("Testing a simple linking with callback...\n");
error = create_single_kernel_helper_create_program(context, &program, 1, &simple_kernel);
if( program == NULL || error != CL_SUCCESS )
{
log_error( "ERROR: Unable to create a simple test program! (%s in %s:%d)\n", IGetErrorString( error ), __FILE__, __LINE__ );
return -1;
}
error = clCompileProgram(program, 1, &deviceID, NULL, 0, NULL, NULL, NULL, NULL);
test_error( error, "Unable to compile a simple program" );
link_program_completion_event = clCreateUserEvent(context, &error);
test_error( error, "Unable to create a user event");
simple_user_data simple_link_user_data = {when_i_pondered_weak_and_weary, link_program_completion_event};
cl_program my_linked_library = clLinkProgram(context, 1, &deviceID, NULL, 1, &program, simple_link_callback, (void*)&simple_link_user_data, &error);
test_error( error, "Unable to link a simple program" );
error = clWaitForEvents(1, &link_program_completion_event);
test_error( error, "clWaitForEvents failed when waiting on link_program_completion_event");
/* All done! */
error = clReleaseEvent(link_program_completion_event);
test_error( error, "Unable to release event object" );
error = clReleaseProgram( program );
test_error( error, "Unable to release program object" );
error = clReleaseProgram( my_linked_library );
test_error( error, "Unable to release program object" );
return 0;
}
static void initBuffer(float* & srcBuffer, unsigned int cnDimension)
{
float num = 0.0f;
for( unsigned int i = 0; i < cnDimension; i++ )
{
if( ( i % 10 ) == 0 )
{
num = 0.0f;
}
srcBuffer[ i ] = num;
num = num + 1.0f;
}
}
static int verifyCopyBuffer(cl_context context, cl_command_queue queue, cl_kernel kernel)
{
int error, result = CL_SUCCESS;
const size_t cnDimension = 32;
// Allocate source buffer
float * srcBuffer = (float*)malloc(cnDimension * sizeof(float));
float * dstBuffer = (float*)malloc(cnDimension * sizeof(float));
if (srcBuffer == NULL) {
log_error( "ERROR: Unable to allocate srcBuffer float array with %lu floats! (in %s:%d)\n", cnDimension, __FILE__, __LINE__);
return -1;
}
if (dstBuffer == NULL) {
log_error( "ERROR: Unable to allocate dstBuffer float array with %lu floats! (in %s:%d)\n", cnDimension, __FILE__, __LINE__);
return -1;
}
if( srcBuffer && dstBuffer )
{
// initialize host memory
initBuffer(srcBuffer, cnDimension );
// Allocate device memory
cl_mem deviceMemSrc = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
cnDimension * sizeof( cl_float ), srcBuffer, &error);
test_error( error, "Unable to create a source memory buffer" );
cl_mem deviceMemDst = clCreateBuffer(context, CL_MEM_WRITE_ONLY,
cnDimension * sizeof( cl_float ), 0, &error);
test_error( error, "Unable to create a destination memory buffer" );
// Set kernel args
// Set parameter 0 to be the source buffer
error = clSetKernelArg(kernel, 0, sizeof( cl_mem ), ( void * )&deviceMemSrc );
test_error( error, "Unable to set the first kernel argument" );
// Set parameter 1 to be the destination buffer
error = clSetKernelArg(kernel, 1, sizeof( cl_mem ), ( void * )&deviceMemDst );
test_error( error, "Unable to set the second kernel argument" );
// Execute kernel
error = clEnqueueNDRangeKernel(queue, kernel, 1, NULL,
&cnDimension, 0, 0, NULL, NULL );
test_error( error, "Unable to enqueue kernel" );
error = clFlush( queue );
test_error( error, "Unable to flush the queue" );
// copy results from device back to host
error = clEnqueueReadBuffer(queue, deviceMemDst, CL_TRUE, 0, cnDimension * sizeof( cl_float ),
dstBuffer, 0, NULL, NULL );
test_error( error, "Unable to read the destination buffer" );
error = clFlush( queue );
test_error( error, "Unable to flush the queue" );
// Compare the source and destination buffers
const int* pSrc = (int*)srcBuffer;
const int* pDst = (int*)dstBuffer;
int mismatch = 0;
for( size_t i = 0; i < cnDimension; i++ )
{
if( pSrc[i] != pDst[i] )
{
if( mismatch < 4 )
{
log_info("Offset %08lX: Expected %08X, Got %08X\n", i * 4, pSrc[i], pDst[i] );
}
else
{
log_info(".");
}
mismatch++;
}
}
if( mismatch )
{
log_info("*** %d mismatches found, TEST FAILS! ***\n", mismatch );
result = -1;
}
else
{
log_info("Buffers match, test passes.\n");
}
free( srcBuffer );
srcBuffer = NULL;
free( dstBuffer );
dstBuffer = NULL;
if( deviceMemSrc )
{
error = clReleaseMemObject( deviceMemSrc );
test_error( error, "Unable to release memory object" );
}
if( deviceMemDst )
{
error = clReleaseMemObject( deviceMemDst );
test_error( error, "Unable to release memory object" );
}
}
return result;
}
int test_execute_after_simple_compile_and_link(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
{
int error;
cl_program program;
log_info("Testing execution after a simple compile and link...\n");
error = create_single_kernel_helper_create_program(context, &program, 1, &simple_kernel);
if( program == NULL || error != CL_SUCCESS )
{
log_error( "ERROR: Unable to create a simple test program! (%s in %s:%d)\n", IGetErrorString( error ), __FILE__, __LINE__ );
return -1;
}
error = clCompileProgram(program, 1, &deviceID, NULL, 0, NULL, NULL, NULL, NULL);
test_error( error, "Unable to compile a simple program" );
cl_program my_newly_linked_program = clLinkProgram(context, 1, &deviceID, NULL, 1, &program, NULL, NULL, &error);
test_error( error, "Unable to link a simple program" );
cl_kernel kernel = clCreateKernel(my_newly_linked_program, "CopyBuffer", &error);
test_error( error, "Unable to create a simple kernel" );
error = verifyCopyBuffer(context, queue, kernel);
if (error != CL_SUCCESS)
return error;
/* All done! */
error = clReleaseKernel( kernel );
test_error( error, "Unable to release kernel object" );
error = clReleaseProgram( program );
test_error( error, "Unable to release program object" );
error = clReleaseProgram( my_newly_linked_program );
test_error( error, "Unable to release program object" );
return 0;
}
int test_execute_after_simple_compile_and_link_no_device_info(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
{
int error;
cl_program program;
log_info("Testing execution after a simple compile and link with no device information provided...\n");
error = create_single_kernel_helper_create_program(context, &program, 1, &simple_kernel);
if( program == NULL || error != CL_SUCCESS )
{
log_error( "ERROR: Unable to create a simple test program! (%s in %s:%d)\n", IGetErrorString( error ), __FILE__, __LINE__ );
return -1;
}
error = clCompileProgram(program, 0, NULL, NULL, 0, NULL, NULL, NULL, NULL);
test_error( error, "Unable to compile a simple program" );
cl_program my_newly_linked_program = clLinkProgram(context, 0, NULL, NULL, 1, &program, NULL, NULL, &error);
test_error( error, "Unable to link a simple program" );
cl_kernel kernel = clCreateKernel(my_newly_linked_program, "CopyBuffer", &error);
test_error( error, "Unable to create a simple kernel" );
error = verifyCopyBuffer(context, queue, kernel);
if (error != CL_SUCCESS)
return error;
/* All done! */
error = clReleaseKernel( kernel );
test_error( error, "Unable to release kernel object" );
error = clReleaseProgram( program );
test_error( error, "Unable to release program object" );
error = clReleaseProgram( my_newly_linked_program );
test_error( error, "Unable to release program object" );
return 0;
}
int test_execute_after_simple_compile_and_link_with_defines(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
{
int error;
cl_program program;
log_info("Testing execution after a simple compile and link with defines...\n");
error = create_single_kernel_helper_create_program(context, &program, 1, &simple_kernel_with_defines, "-DFIRST=5 -DSECOND=37");
if( program == NULL || error != CL_SUCCESS )
{
log_error( "ERROR: Unable to create a simple test program! (%s in %s:%d)\n", IGetErrorString( error ), __FILE__, __LINE__ );
return -1;
}
error = clCompileProgram(program, 1, &deviceID, "-DFIRST=5 -DSECOND=37", 0, NULL, NULL, NULL, NULL);
test_error( error, "Unable to compile a simple program" );
cl_program my_newly_linked_program = clLinkProgram(context, 1, &deviceID, NULL, 1, &program, NULL, NULL, &error);
test_error( error, "Unable to link a simple program" );
cl_kernel kernel = clCreateKernel(my_newly_linked_program, "CopyBuffer", &error);
test_error( error, "Unable to create a simple kernel" );
error = verifyCopyBuffer(context, queue, kernel);
if (error != CL_SUCCESS)
return error;
/* All done! */
error = clReleaseKernel( kernel );
test_error( error, "Unable to release kernel object" );
error = clReleaseProgram( program );
test_error( error, "Unable to release program object" );
error = clReleaseProgram( my_newly_linked_program );
test_error( error, "Unable to release program object" );
return 0;
}
int test_execute_after_serialize_reload_object(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
{
int error;
cl_program program;
size_t binarySize;
unsigned char *binary;
log_info("Testing execution after serialization and reloading of the object...\n");
error = create_single_kernel_helper_create_program(context, &program, 1, &simple_kernel);
if( program == NULL || error != CL_SUCCESS )
{
log_error( "ERROR: Unable to create a simple test program! (%s in %s:%d)\n", IGetErrorString( error ), __FILE__, __LINE__ );
return -1;
}
error = clCompileProgram(program, 1, &deviceID, NULL, 0, NULL, NULL, NULL, NULL);
test_error( error, "Unable to compile a simple program" );
// Get the size of the resulting binary (only one device)
error = clGetProgramInfo( program, CL_PROGRAM_BINARY_SIZES, sizeof( binarySize ), &binarySize, NULL );
test_error( error, "Unable to get binary size" );
// Sanity check
if( binarySize == 0 )
{
log_error( "ERROR: Binary size of program is zero (in %s:%d)\n", __FILE__, __LINE__ );
return -1;
}
// Create a buffer and get the actual binary
binary = (unsigned char*)malloc(sizeof(unsigned char)*binarySize);
if (binary == NULL) {
log_error( "ERROR: Unable to allocate binary character array with %lu characters! (in %s:%d)\n", binarySize, __FILE__, __LINE__ );
return -1;
}
unsigned char *buffers[ 1 ] = { binary };
cl_int loadErrors[ 1 ];
// Do another sanity check here first
size_t size;
error = clGetProgramInfo( program, CL_PROGRAM_BINARIES, 0, NULL, &size );
test_error( error, "Unable to get expected size of binaries array" );
if( size != sizeof( buffers ) )
{
log_error( "ERROR: Expected size of binaries array in clGetProgramInfo is incorrect (should be %d, got %d) (in %s:%d)\n", (int)sizeof( buffers ), (int)size, __FILE__, __LINE__ );
free(binary);
return -1;
}
error = clGetProgramInfo( program, CL_PROGRAM_BINARIES, sizeof( buffers ), &buffers, NULL );
test_error( error, "Unable to get program binary" );
// use clCreateProgramWithBinary
cl_program program_with_binary = clCreateProgramWithBinary(context, 1, &deviceID, &binarySize, (const unsigned char**)buffers, loadErrors, &error);
test_error( error, "Unable to create program with binary" );
cl_program my_newly_linked_program = clLinkProgram(context, 1, &deviceID, NULL, 1, &program_with_binary, NULL, NULL, &error);
test_error( error, "Unable to link a simple program" );
cl_kernel kernel = clCreateKernel(my_newly_linked_program, "CopyBuffer", &error);
test_error( error, "Unable to create a simple kernel" );
error = verifyCopyBuffer(context, queue, kernel);
if (error != CL_SUCCESS)
return error;
/* All done! */
error = clReleaseKernel( kernel );
test_error( error, "Unable to release kernel object" );
error = clReleaseProgram( program );
test_error( error, "Unable to release program object" );
error = clReleaseProgram( my_newly_linked_program );
test_error( error, "Unable to release program object" );
error = clReleaseProgram( program_with_binary );
test_error( error, "Unable to release program object" );
free(binary);
return 0;
}
int test_execute_after_serialize_reload_library(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
{
int error;
cl_program program, another_program;
size_t binarySize;
unsigned char *binary;
log_info("Testing execution after linking a binary with a simple library...\n");
// we will test creation of a simple library from one file
error = create_single_kernel_helper_create_program(context, &program, 1, &simple_kernel);
if( program == NULL || error != CL_SUCCESS )
{
log_error( "ERROR: Unable to create a simple test program! (%s in %s:%d)\n", IGetErrorString( error ), __FILE__, __LINE__ );
return -1;
}
error = clCompileProgram(program, 1, &deviceID, NULL, 0, NULL, NULL, NULL, NULL);
test_error( error, "Unable to compile a simple program" );
cl_program my_newly_minted_library = clLinkProgram(context, 1, &deviceID, "-create-library", 1, &program, NULL, NULL, &error);
test_error( error, "Unable to create a simple library" );
// Get the size of the resulting library (only one device)
error = clGetProgramInfo( my_newly_minted_library, CL_PROGRAM_BINARY_SIZES, sizeof( binarySize ), &binarySize, NULL );
test_error( error, "Unable to get binary size" );
// Sanity check
if( binarySize == 0 )
{
log_error( "ERROR: Binary size of program is zero (in %s:%d)\n", __FILE__, __LINE__ );
return -1;
}
// Create a buffer and get the actual binary
binary = (unsigned char*)malloc(sizeof(unsigned char)*binarySize);
if (binary == NULL) {
log_error( "ERROR: Unable to allocate binary character array with %lu characters (in %s:%d)!", binarySize, __FILE__, __LINE__);
return -1;
}
unsigned char *buffers[ 1 ] = { binary };
cl_int loadErrors[ 1 ];
// Do another sanity check here first
size_t size;
error = clGetProgramInfo( my_newly_minted_library, CL_PROGRAM_BINARIES, 0, NULL, &size );
test_error( error, "Unable to get expected size of binaries array" );
if( size != sizeof( buffers ) )
{
log_error( "ERROR: Expected size of binaries array in clGetProgramInfo is incorrect (should be %d, got %d) (in %s:%d)\n", (int)sizeof( buffers ), (int)size, __FILE__, __LINE__ );
free(binary);
return -1;
}
error = clGetProgramInfo( my_newly_minted_library, CL_PROGRAM_BINARIES, sizeof( buffers ), &buffers, NULL );
test_error( error, "Unable to get program binary" );
// use clCreateProgramWithBinary
cl_program library_with_binary = clCreateProgramWithBinary(context, 1, &deviceID, &binarySize, (const unsigned char**)buffers, loadErrors, &error);
test_error( error, "Unable to create program with binary" );
error = create_single_kernel_helper_create_program(context, &another_program, 1, &another_simple_kernel);
if( another_program == NULL || error != CL_SUCCESS )
{
log_error( "ERROR: Unable to create a simple test program! (%s in %s:%d)\n", IGetErrorString( error ), __FILE__, __LINE__ );
return -1;
}
error = clCompileProgram(another_program, 1, &deviceID, NULL, 0, NULL, NULL, NULL, NULL);
test_error( error, "Unable to compile a simple program" );
cl_program program_and_archive[2] = { another_program, library_with_binary };
cl_program fully_linked_program = clLinkProgram(context, 1, &deviceID, "", 2, program_and_archive, NULL, NULL, &error);
test_error( error, "Unable to create an executable from a binary and a library" );
cl_kernel kernel = clCreateKernel(fully_linked_program, "CopyBuffer", &error);
test_error( error, "Unable to create a simple kernel" );
error = verifyCopyBuffer(context, queue, kernel);
if (error != CL_SUCCESS)
return error;
cl_kernel another_kernel = clCreateKernel(fully_linked_program, "AnotherCopyBuffer", &error);
test_error( error, "Unable to create another simple kernel" );
error = verifyCopyBuffer(context, queue, another_kernel);
if (error != CL_SUCCESS)
return error;
/* All done! */
error = clReleaseKernel( kernel );
test_error( error, "Unable to release kernel object" );
error = clReleaseKernel( another_kernel );
test_error( error, "Unable to release another kernel object" );
error = clReleaseProgram( program );
test_error( error, "Unable to release program object" );
error = clReleaseProgram( another_program );
test_error( error, "Unable to release program object" );
error = clReleaseProgram( my_newly_minted_library );
test_error( error, "Unable to release program object" );
error = clReleaseProgram( library_with_binary );
test_error( error, "Unable to release program object" );
error = clReleaseProgram( fully_linked_program );
test_error( error, "Unable to release program object" );
free(binary);
return 0;
}
static void CL_CALLBACK program_compile_completion_callback(cl_program program, void* user_data)
{
int error;
cl_event compile_program_completion_event = (cl_event)user_data;
log_info("in the program_compile_completion_callback: program %p just completed compiling with '%p'\n", program, compile_program_completion_event);
error = clSetUserEventStatus(compile_program_completion_event, CL_COMPLETE);
if (error != CL_SUCCESS)
{
log_error( "ERROR: in the program_compile_completion_callback: Unable to set user event status to CL_COMPLETE! (%s in %s:%d)\n", IGetErrorString( error ), __FILE__, __LINE__ );
exit(-1);
}
log_info("in the program_compile_completion_callback: Successfully signaled compile_program_completion_event event!\n");
}
static void CL_CALLBACK program_link_completion_callback(cl_program program, void* user_data)
{
int error;
cl_event link_program_completion_event = (cl_event)user_data;
log_info("in the program_link_completion_callback: program %p just completed linking with '%p'\n", program, link_program_completion_event);
error = clSetUserEventStatus(link_program_completion_event, CL_COMPLETE);
if (error != CL_SUCCESS)
{
log_error( "ERROR: in the program_link_completion_callback: Unable to set user event status to CL_COMPLETE! (%s in %s:%d)\n", IGetErrorString( error ), __FILE__, __LINE__ );
exit(-1);
}
log_info("in the program_link_completion_callback: Successfully signaled link_program_completion_event event!\n");
}
int test_execute_after_simple_compile_and_link_with_callbacks(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
{
int error;
cl_program program;
cl_event compile_program_completion_event, link_program_completion_event;
log_info("Testing execution after a simple compile and link with callbacks...\n");
error = create_single_kernel_helper_create_program(context, &program, 1, &simple_kernel);
if( program == NULL || error != CL_SUCCESS )
{
log_error( "ERROR: Unable to create a simple test program! (%s in %s:%d)\n", IGetErrorString( error ), __FILE__, __LINE__ );
return -1;
}
compile_program_completion_event = clCreateUserEvent(context, &error);
test_error( error, "Unable to create a user event");
error = clCompileProgram(program, 1, &deviceID, NULL, 0, NULL, NULL,
program_compile_completion_callback, (void*)compile_program_completion_event);
test_error( error, "Unable to compile a simple program" );
error = clWaitForEvents(1, &compile_program_completion_event);
test_error( error, "clWaitForEvents failed when waiting on compile_program_completion_event");
error = clReleaseEvent(compile_program_completion_event);
test_error( error, "Unable to release event object" );
link_program_completion_event = clCreateUserEvent(context, &error);
test_error( error, "Unable to create a user event");
cl_program my_newly_linked_program = clLinkProgram(context, 1, &deviceID, NULL, 1, &program,
program_link_completion_callback, (void*)link_program_completion_event, &error);
test_error( error, "Unable to link a simple program" );
error = clWaitForEvents(1, &link_program_completion_event);
test_error( error, "clWaitForEvents failed when waiting on link_program_completion_event");
error = clReleaseEvent(link_program_completion_event);
test_error( error, "Unable to release event object" );
cl_kernel kernel = clCreateKernel(my_newly_linked_program, "CopyBuffer", &error);
test_error( error, "Unable to create a simple kernel" );
error = verifyCopyBuffer(context, queue, kernel);
if (error != CL_SUCCESS)
return error;
/* All done! */
error = clReleaseKernel( kernel );
test_error( error, "Unable to release kernel object" );
error = clReleaseProgram( program );
test_error( error, "Unable to release program object" );
error = clReleaseProgram( my_newly_linked_program );
test_error( error, "Unable to release program object" );
return 0;
}
int test_simple_library_only(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
{
int error;
cl_program program;
log_info("Testing creation of a simple library...\n");
error = create_single_kernel_helper_create_program(context, &program, 1, &simple_kernel);
if( program == NULL || error != CL_SUCCESS )
{
log_error( "ERROR: Unable to create a simple test program! (%s in %s:%d)\n", IGetErrorString( error ), __FILE__, __LINE__ );
return -1;
}
error = clCompileProgram(program, 1, &deviceID, NULL, 0, NULL, NULL, NULL, NULL);
test_error( error, "Unable to compile a simple program" );
cl_program my_newly_minted_library = clLinkProgram(context, 1, &deviceID, "-create-library", 1, &program, NULL, NULL, &error);
test_error( error, "Unable to create a simple library" );
/* All done! */
error = clReleaseProgram( program );
test_error( error, "Unable to release program object" );
error = clReleaseProgram( my_newly_minted_library );
test_error( error, "Unable to release program object" );
return 0;
}
int test_simple_library_with_callback(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
{
int error;
cl_program program;
cl_event link_program_completion_event;
log_info("Testing creation of a simple library with a callback...\n");
error = create_single_kernel_helper_create_program(context, &program, 1, &simple_kernel);
if( program == NULL || error != CL_SUCCESS )
{
log_error( "ERROR: Unable to create a simple test program! (%s in %s:%d)\n", IGetErrorString( error ), __FILE__, __LINE__ );
return -1;
}
error = clCompileProgram(program, 1, &deviceID, NULL, 0, NULL, NULL, NULL, NULL);
test_error( error, "Unable to compile a simple program" );
link_program_completion_event = clCreateUserEvent(context, &error);
test_error( error, "Unable to create a user event");
simple_user_data simple_link_user_data = {when_i_pondered_weak_and_weary, link_program_completion_event};
cl_program my_newly_minted_library = clLinkProgram(context, 1, &deviceID, "-create-library", 1, &program,
simple_link_callback, (void*)&simple_link_user_data, &error);
test_error( error, "Unable to create a simple library" );
error = clWaitForEvents(1, &link_program_completion_event);
test_error( error, "clWaitForEvents failed when waiting on link_program_completion_event");
/* All done! */
error = clReleaseEvent(link_program_completion_event);
test_error( error, "Unable to release event object" );
error = clReleaseProgram( program );
test_error( error, "Unable to release program object" );
error = clReleaseProgram( my_newly_minted_library );
test_error( error, "Unable to release program object" );
return 0;
}
int test_simple_library_with_link(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
{
int error;
cl_program program, another_program;
log_info("Testing creation and linking with a simple library...\n");
error = create_single_kernel_helper_create_program(context, &program, 1, &simple_kernel);
if( program == NULL || error != CL_SUCCESS )
{
log_error( "ERROR: Unable to create a simple test program! (%s in %s:%d)\n", IGetErrorString( error ), __FILE__, __LINE__ );
return -1;
}
error = clCompileProgram(program, 1, &deviceID, NULL, 0, NULL, NULL, NULL, NULL);
test_error( error, "Unable to compile a simple program" );
cl_program my_newly_minted_library = clLinkProgram(context, 1, &deviceID, "-create-library", 1, &program, NULL, NULL, &error);
test_error( error, "Unable to create a simple library" );
error = create_single_kernel_helper_create_program(context, &another_program, 1, &another_simple_kernel);
if( another_program == NULL || error != CL_SUCCESS )
{
log_error( "ERROR: Unable to create a simple test program! (%s in %s:%d)\n", IGetErrorString( error ), __FILE__, __LINE__ );
return -1;
}
error = clCompileProgram(another_program, 1, &deviceID, NULL, 0, NULL, NULL, NULL, NULL);
test_error( error, "Unable to compile a simple program" );
cl_program program_and_archive[2] = { another_program, my_newly_minted_library };
cl_program fully_linked_program = clLinkProgram(context, 1, &deviceID, "", 2, program_and_archive, NULL, NULL, &error);
test_error( error, "Unable to create an executable from a binary and a library" );
/* All done! */
error = clReleaseProgram( program );
test_error( error, "Unable to release program object" );
error = clReleaseProgram( another_program );
test_error( error, "Unable to release program object" );
error = clReleaseProgram( my_newly_minted_library );
test_error( error, "Unable to release program object" );
error = clReleaseProgram( fully_linked_program );
test_error( error, "Unable to release program object" );
return 0;
}
int test_execute_after_simple_library_with_link(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
{
int error;
cl_program program, another_program;
log_info("Testing execution after linking a binary with a simple library...\n");
error = create_single_kernel_helper_create_program(context, &program, 1, &simple_kernel);
if( program == NULL || error != CL_SUCCESS )
{
log_error( "ERROR: Unable to create a simple test program! (%s in %s:%d)\n", IGetErrorString( error ), __FILE__, __LINE__ );
return -1;
}
error = clCompileProgram(program, 1, &deviceID, NULL, 0, NULL, NULL, NULL, NULL);
test_error( error, "Unable to compile a simple program" );
cl_program my_newly_minted_library = clLinkProgram(context, 1, &deviceID, "-create-library", 1, &program, NULL, NULL, &error);
test_error( error, "Unable to create a simple library" );
error = create_single_kernel_helper_create_program(context, &another_program, 1, &another_simple_kernel);
if( another_program == NULL || error != CL_SUCCESS )
{
log_error( "ERROR: Unable to create a simple test program! (%s in %s:%d)\n", IGetErrorString( error ), __FILE__, __LINE__ );
return -1;
}
error = clCompileProgram(another_program, 1, &deviceID, NULL, 0, NULL, NULL, NULL, NULL);
test_error( error, "Unable to compile a simple program" );
cl_program program_and_archive[2] = { another_program, my_newly_minted_library };
cl_program fully_linked_program = clLinkProgram(context, 1, &deviceID, "", 2, program_and_archive, NULL, NULL, &error);
test_error( error, "Unable to create an executable from a binary and a library" );
cl_kernel kernel = clCreateKernel(fully_linked_program, "CopyBuffer", &error);
test_error( error, "Unable to create a simple kernel" );
error = verifyCopyBuffer(context, queue, kernel);
if (error != CL_SUCCESS)
return error;
cl_kernel another_kernel = clCreateKernel(fully_linked_program, "AnotherCopyBuffer", &error);
test_error( error, "Unable to create another simple kernel" );
error = verifyCopyBuffer(context, queue, another_kernel);
if (error != CL_SUCCESS)
return error;
/* All done! */
error = clReleaseKernel( kernel );
test_error( error, "Unable to release kernel object" );
error = clReleaseKernel( another_kernel );
test_error( error, "Unable to release another kernel object" );
error = clReleaseProgram( program );
test_error( error, "Unable to release program object" );
error = clReleaseProgram( another_program );
test_error( error, "Unable to release program object" );
error = clReleaseProgram( my_newly_minted_library );
test_error( error, "Unable to release program object" );
error = clReleaseProgram( fully_linked_program );
test_error( error, "Unable to release program object" );
return 0;
}
int test_two_file_link(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
{
int error;
cl_program program, another_program;
log_info("Testing two file compiling and linking...\n");
error = create_single_kernel_helper_create_program(context, &program, 1, &simple_kernel);
if( program == NULL || error != CL_SUCCESS )
{
log_error( "ERROR: Unable to create a simple test program! (%s in %s:%d)\n", IGetErrorString( error ), __FILE__, __LINE__ );
return -1;
}
error = clCompileProgram(program, 1, &deviceID, NULL, 0, NULL, NULL, NULL, NULL);
test_error( error, "Unable to compile a simple program" );
error = create_single_kernel_helper_create_program(context, &another_program, 1, &another_simple_kernel);
if( another_program == NULL || error != CL_SUCCESS )
{
log_error( "ERROR: Unable to create a simple test program! (%s in %s:%d)\n", IGetErrorString( error ), __FILE__, __LINE__ );
return -1;
}
error = clCompileProgram(another_program, 1, &deviceID, NULL, 0, NULL, NULL, NULL, NULL);
test_error( error, "Unable to compile a simple program" );
cl_program two_programs[2] = { program, another_program };
cl_program fully_linked_program = clLinkProgram(context, 1, &deviceID, "", 2, two_programs, NULL, NULL, &error);
test_error( error, "Unable to create an executable from two binaries" );
/* All done! */
error = clReleaseProgram( program );
test_error( error, "Unable to release program object" );
error = clReleaseProgram( another_program );
test_error( error, "Unable to release program object" );
error = clReleaseProgram( fully_linked_program );
test_error( error, "Unable to release program object" );
return 0;
}
int test_execute_after_two_file_link(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
{
int error;
cl_program program, another_program;
log_info("Testing two file compiling and linking and execution of two kernels afterwards ...\n");
error = create_single_kernel_helper_create_program(context, &program, 1, &simple_kernel);
if( program == NULL || error != CL_SUCCESS )
{
log_error( "ERROR: Unable to create a simple test program! (%s in %s:%d)\n", IGetErrorString( error ), __FILE__, __LINE__ );
return -1;
}
error = clCompileProgram(program, 1, &deviceID, NULL, 0, NULL, NULL, NULL, NULL);
test_error( error, "Unable to compile a simple program" );
error = create_single_kernel_helper_create_program(context, &another_program, 1, &another_simple_kernel);
if( another_program == NULL || error != CL_SUCCESS )
{
log_error( "ERROR: Unable to create a simple test program! (%s in %s:%d)\n", IGetErrorString( error ), __FILE__, __LINE__ );
return -1;
}
error = clCompileProgram(another_program, 1, &deviceID, NULL, 0, NULL, NULL, NULL, NULL);
test_error( error, "Unable to compile a simple program" );
cl_program two_programs[2] = { program, another_program };
cl_program fully_linked_program = clLinkProgram(context, 1, &deviceID, "", 2, two_programs, NULL, NULL, &error);
test_error( error, "Unable to create an executable from two binaries" );
cl_kernel kernel = clCreateKernel(fully_linked_program, "CopyBuffer", &error);
test_error( error, "Unable to create a simple kernel" );
error = verifyCopyBuffer(context, queue, kernel);
if (error != CL_SUCCESS)
return error;
cl_kernel another_kernel = clCreateKernel(fully_linked_program, "AnotherCopyBuffer", &error);
test_error( error, "Unable to create another simple kernel" );
error = verifyCopyBuffer(context, queue, another_kernel);
if (error != CL_SUCCESS)
return error;
/* All done! */
error = clReleaseKernel( kernel );
test_error( error, "Unable to release kernel object" );
error = clReleaseKernel( another_kernel );
test_error( error, "Unable to release another kernel object"