mirror of
https://github.com/KhronosGroup/OpenCL-CTS.git
synced 2026-03-19 06:09:01 +00:00
e92140f82d
In C++17 the return type of std::filesystem::path::u8string() is std::string, but in C++20 the return type changed to std::u8string. Add a helper to copy a std::u8string to a std::string, to be used when a std::string is required. This fixes the build for C++20.
3936 lines
140 KiB
C++
3936 lines
140 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"
|
|
|
|
#include <filesystem>
|
|
|
|
#if defined(_WIN32)
|
|
#include <time.h>
|
|
#elif defined(__linux__) || defined(__APPLE__)
|
|
#include <sys/time.h>
|
|
#include <unistd.h>
|
|
#endif
|
|
#include "harness/conversions.h"
|
|
#include "harness/stringHelpers.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 = 0.0f;\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";
|
|
|
|
static 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;
|
|
}
|
|
|
|
REGISTER_TEST(large_compile)
|
|
{
|
|
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, device, 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
|
|
|
|
static 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;
|
|
}
|
|
|
|
REGISTER_TEST(multi_file_libraries)
|
|
{
|
|
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, device, 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;
|
|
}
|
|
|
|
static 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 = 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;
|
|
}
|
|
|
|
REGISTER_TEST(multiple_embedded_headers)
|
|
{
|
|
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, device, 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); }
|
|
|
|
static 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;
|
|
}
|
|
|
|
REGISTER_TEST(multiple_libraries)
|
|
{
|
|
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, device, 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;
|
|
}
|
|
|
|
static 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;
|
|
}
|
|
|
|
REGISTER_TEST(multiple_files_multiple_libraries)
|
|
{
|
|
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, device, 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;
|
|
}
|
|
|
|
static 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;
|
|
}
|
|
|
|
REGISTER_TEST(multiple_files)
|
|
{
|
|
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, device, 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;
|
|
}
|
|
|
|
REGISTER_TEST(simple_compile_only)
|
|
{
|
|
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, &device, 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;
|
|
}
|
|
|
|
REGISTER_TEST(simple_static_compile_only)
|
|
{
|
|
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, &device, 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, &device, 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, &device, 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;
|
|
}
|
|
|
|
REGISTER_TEST(simple_extern_compile_only)
|
|
{
|
|
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, &device, 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, &device, 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, &device, 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, &device, 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");
|
|
}
|
|
|
|
REGISTER_TEST(simple_compile_with_callback)
|
|
{
|
|
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, &device, 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;
|
|
}
|
|
|
|
REGISTER_TEST(simple_embedded_header_compile)
|
|
{
|
|
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, &device, 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;
|
|
}
|
|
|
|
REGISTER_TEST(simple_link_only)
|
|
{
|
|
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, &device, NULL, 0, NULL, NULL, NULL, NULL);
|
|
test_error(error, "Unable to compile a simple program");
|
|
|
|
cl_program my_newly_linked_program = clLinkProgram(
|
|
context, 1, &device, 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;
|
|
}
|
|
|
|
REGISTER_TEST(two_file_regular_variable_access)
|
|
{
|
|
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, &device, 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, &device, 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, &device, 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;
|
|
}
|
|
|
|
REGISTER_TEST(two_file_regular_struct_access)
|
|
{
|
|
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, &device, 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, &device, 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, &device, 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;
|
|
}
|
|
|
|
|
|
REGISTER_TEST(two_file_regular_function_access)
|
|
{
|
|
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, &device, 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, &device, 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, &device, 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;
|
|
}
|
|
|
|
REGISTER_TEST(simple_embedded_header_link)
|
|
{
|
|
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, &device, 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, &device, 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, &device, "", 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");
|
|
}
|
|
|
|
REGISTER_TEST(simple_link_with_callback)
|
|
{
|
|
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, &device, 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, &device, 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 %zu "
|
|
"floats! (in %s:%d)\n",
|
|
cnDimension, __FILE__, __LINE__);
|
|
return -1;
|
|
}
|
|
if (dstBuffer == NULL)
|
|
{
|
|
log_error("ERROR: Unable to allocate dstBuffer float array with %zu "
|
|
"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 %08zX: 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;
|
|
}
|
|
|
|
REGISTER_TEST(execute_after_simple_compile_and_link)
|
|
{
|
|
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, &device, NULL, 0, NULL, NULL, NULL, NULL);
|
|
test_error(error, "Unable to compile a simple program");
|
|
|
|
cl_program my_newly_linked_program = clLinkProgram(
|
|
context, 1, &device, 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;
|
|
}
|
|
|
|
REGISTER_TEST(execute_after_simple_compile_and_link_no_device_info)
|
|
{
|
|
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;
|
|
}
|
|
|
|
REGISTER_TEST(execute_after_simple_compile_and_link_with_defines)
|
|
{
|
|
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, &device, "-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, &device, 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;
|
|
}
|
|
|
|
REGISTER_TEST(execute_after_serialize_reload_object)
|
|
{
|
|
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, &device, 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 %zu "
|
|
"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, &device, &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, &device, 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;
|
|
}
|
|
|
|
REGISTER_TEST(execute_after_serialize_reload_library)
|
|
{
|
|
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, &device, NULL, 0, NULL, NULL, NULL, NULL);
|
|
test_error(error, "Unable to compile a simple program");
|
|
|
|
cl_program my_newly_minted_library =
|
|
clLinkProgram(context, 1, &device, "-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 %zu "
|
|
"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, &device, &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, &device, 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, &device, "", 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");
|
|
}
|
|
|
|
REGISTER_TEST(execute_after_simple_compile_and_link_with_callbacks)
|
|
{
|
|
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, &device, 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, &device, 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;
|
|
}
|
|
|
|
REGISTER_TEST(simple_library_only)
|
|
{
|
|
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, &device, NULL, 0, NULL, NULL, NULL, NULL);
|
|
test_error(error, "Unable to compile a simple program");
|
|
|
|
cl_program my_newly_minted_library =
|
|
clLinkProgram(context, 1, &device, "-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;
|
|
}
|
|
|
|
REGISTER_TEST(simple_library_with_callback)
|
|
{
|
|
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, &device, 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, &device, "-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;
|
|
}
|
|
|
|
REGISTER_TEST(simple_library_with_link)
|
|
{
|
|
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, &device, NULL, 0, NULL, NULL, NULL, NULL);
|
|
test_error(error, "Unable to compile a simple program");
|
|
|
|
cl_program my_newly_minted_library =
|
|
clLinkProgram(context, 1, &device, "-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, &device, 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, &device, "", 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;
|
|
}
|
|
|
|
REGISTER_TEST(execute_after_simple_library_with_link)
|
|
{
|
|
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, &device, NULL, 0, NULL, NULL, NULL, NULL);
|
|
test_error(error, "Unable to compile a simple program");
|
|
|
|
cl_program my_newly_minted_library =
|
|
clLinkProgram(context, 1, &device, "-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, &device, 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, &device, "", 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;
|
|
}
|
|
|
|
REGISTER_TEST(two_file_link)
|
|
{
|
|
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, &device, 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, &device, 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, &device, "", 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;
|
|
}
|
|
|
|
REGISTER_TEST(execute_after_two_file_link)
|
|
{
|
|
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, &device, 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, &device, 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, &device, "", 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");
|
|
|
|
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;
|
|
}
|
|
|
|
REGISTER_TEST(execute_after_embedded_header_link)
|
|
{
|
|
int error;
|
|
cl_program program, header, simple_program;
|
|
|
|
log_info("Testing execution after embedded header link...\n");
|
|
// we will test execution after compiling and linking with embedded headers
|
|
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, &device, NULL, 1, &header,
|
|
&simple_header_name, NULL, NULL);
|
|
test_error(error,
|
|
"Unable to compile a simple program with embedded header");
|
|
|
|
simple_program =
|
|
clCreateProgramWithSource(context, 1, &simple_kernel, NULL, &error);
|
|
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, &device, 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, &device, "", 2, two_programs, NULL, NULL, &error);
|
|
test_error(error,
|
|
"Unable to create an executable from two binaries, one compiled "
|
|
"with embedded header");
|
|
|
|
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(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;
|
|
}
|
|
|
|
REGISTER_TEST(execute_after_included_header_link)
|
|
{
|
|
int error;
|
|
cl_program program, simple_program;
|
|
|
|
log_info("Testing execution after included header link...\n");
|
|
// we will test execution after compiling and linking with included headers
|
|
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;
|
|
}
|
|
|
|
/* setup */
|
|
std::error_code ec;
|
|
auto temp_dir_path = std::filesystem::temp_directory_path(ec);
|
|
if (ec)
|
|
{
|
|
log_error("ERROR: Unable to get the temporary directory path\n");
|
|
return -1;
|
|
}
|
|
temp_dir_path = temp_dir_path / "foo" / "bar";
|
|
std::filesystem::create_directories(temp_dir_path, ec);
|
|
if (ec)
|
|
{
|
|
log_error("ERROR: Unable to create directory: %s, error: %d (%s)\n",
|
|
temp_dir_path.u8string().c_str(), ec.value(),
|
|
ec.message().c_str());
|
|
return -1;
|
|
}
|
|
|
|
const auto simple_header_path = temp_dir_path / simple_header_name;
|
|
const std::string simple_header_path_str =
|
|
to_string(simple_header_path.u8string());
|
|
FILE *simple_header_file = fopen(simple_header_path_str.c_str(), "w");
|
|
if (simple_header_file == NULL)
|
|
{
|
|
log_error("ERROR: Unable to create simple header file %s! (in %s:%d)\n",
|
|
simple_header_path_str.c_str(), __FILE__, __LINE__);
|
|
return -1;
|
|
}
|
|
if (fprintf(simple_header_file, "%s", simple_header) < 0)
|
|
{
|
|
log_error(
|
|
"ERROR: Unable to write to simple header file %s! (in %s:%d)\n",
|
|
simple_header_path.u8string().c_str(), __FILE__, __LINE__);
|
|
return -1;
|
|
}
|
|
if (fclose(simple_header_file) != 0)
|
|
{
|
|
log_error("ERROR: Unable to close simple header file %s! (in %s:%d)\n",
|
|
simple_header_path.u8string().c_str(), __FILE__, __LINE__);
|
|
return -1;
|
|
}
|
|
|
|
const std::string include_path =
|
|
std::string("-I") + to_string(temp_dir_path.generic_u8string());
|
|
error = clCompileProgram(program, 1, &device, include_path.c_str(), 0, NULL,
|
|
NULL, NULL, NULL);
|
|
test_error(error,
|
|
"Unable to compile a simple program with included header");
|
|
|
|
/* cleanup */
|
|
std::filesystem::remove_all(temp_dir_path, ec);
|
|
if (ec)
|
|
{
|
|
log_error("ERROR: Unable to delete directory: %s, error: %d (%s)",
|
|
temp_dir_path.u8string().c_str(), ec.value(),
|
|
ec.message().c_str());
|
|
return -1;
|
|
}
|
|
|
|
simple_program =
|
|
clCreateProgramWithSource(context, 1, &simple_kernel, NULL, &error);
|
|
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, &device, 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, &device, "", 2, two_programs, NULL, NULL, &error);
|
|
test_error(error,
|
|
"Unable to create an executable from two binaries, one compiled "
|
|
"with embedded header");
|
|
|
|
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(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;
|
|
}
|
|
|
|
REGISTER_TEST(program_binary_type)
|
|
{
|
|
int error;
|
|
cl_program program, another_program, program_with_binary,
|
|
fully_linked_program_with_binary;
|
|
cl_program_binary_type program_type = -1;
|
|
size_t size;
|
|
size_t binarySize;
|
|
unsigned char *binary;
|
|
|
|
log_info("Testing querying of program binary type...\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, &device, NULL, 0, NULL, NULL, NULL, NULL);
|
|
test_error(error, "Unable to compile a simple program");
|
|
|
|
error = clGetProgramBuildInfo(program, device, CL_PROGRAM_BINARY_TYPE,
|
|
sizeof(cl_program_binary_type), &program_type,
|
|
NULL);
|
|
test_error(error, "Unable to get program binary type");
|
|
if (program_type != CL_PROGRAM_BINARY_TYPE_COMPILED_OBJECT)
|
|
{
|
|
log_error("ERROR: Expected program type of a just compiled program to "
|
|
"be CL_PROGRAM_BINARY_TYPE_COMPILED_OBJECT (in %s:%d)\n",
|
|
__FILE__, __LINE__);
|
|
return -1;
|
|
}
|
|
program_type = -1;
|
|
|
|
// 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 "
|
|
"%zu 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
|
|
program_with_binary = clCreateProgramWithBinary(
|
|
context, 1, &device, &binarySize, (const unsigned char **)buffers,
|
|
loadErrors, &error);
|
|
test_error(error, "Unable to create program with binary");
|
|
|
|
error = clGetProgramBuildInfo(
|
|
program_with_binary, device, CL_PROGRAM_BINARY_TYPE,
|
|
sizeof(cl_program_binary_type), &program_type, NULL);
|
|
test_error(error, "Unable to get program binary type");
|
|
if (program_type != CL_PROGRAM_BINARY_TYPE_COMPILED_OBJECT)
|
|
{
|
|
log_error("ERROR: Expected program type of a program created from "
|
|
"compiled object to be "
|
|
"CL_PROGRAM_BINARY_TYPE_COMPILED_OBJECT (in %s:%d)\n",
|
|
__FILE__, __LINE__);
|
|
return -1;
|
|
}
|
|
program_type = -1;
|
|
free(binary);
|
|
}
|
|
|
|
cl_program my_newly_minted_library =
|
|
clLinkProgram(context, 1, &device, "-create-library", 1,
|
|
&program_with_binary, NULL, NULL, &error);
|
|
test_error(error, "Unable to create a simple library");
|
|
error = clGetProgramBuildInfo(
|
|
my_newly_minted_library, device, CL_PROGRAM_BINARY_TYPE,
|
|
sizeof(cl_program_binary_type), &program_type, NULL);
|
|
test_error(error, "Unable to get program binary type");
|
|
if (program_type != CL_PROGRAM_BINARY_TYPE_LIBRARY)
|
|
{
|
|
log_error("ERROR: Expected program type of a just linked library to be "
|
|
"CL_PROGRAM_BINARY_TYPE_LIBRARY (in %s:%d)\n",
|
|
__FILE__, __LINE__);
|
|
return -1;
|
|
}
|
|
program_type = -1;
|
|
|
|
// 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 %zu "
|
|
"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
|
|
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, &device, &binarySize, (const unsigned char **)buffers,
|
|
loadErrors, &error);
|
|
test_error(error, "Unable to create program with binary");
|
|
error = clGetProgramBuildInfo(
|
|
library_with_binary, device, CL_PROGRAM_BINARY_TYPE,
|
|
sizeof(cl_program_binary_type), &program_type, NULL);
|
|
test_error(error, "Unable to get program binary type");
|
|
if (program_type != CL_PROGRAM_BINARY_TYPE_LIBRARY)
|
|
{
|
|
log_error("ERROR: Expected program type of a library loaded with "
|
|
"binary to be CL_PROGRAM_BINARY_TYPE_LIBRARY (in %s:%d)\n",
|
|
__FILE__, __LINE__);
|
|
return -1;
|
|
}
|
|
program_type = -1;
|
|
free(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, &device, 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, &device, "", 2, program_and_archive, NULL, NULL, &error);
|
|
test_error(error,
|
|
"Unable to create an executable from a binary and a library");
|
|
|
|
error = clGetProgramBuildInfo(
|
|
fully_linked_program, device, CL_PROGRAM_BINARY_TYPE,
|
|
sizeof(cl_program_binary_type), &program_type, NULL);
|
|
test_error(error, "Unable to get program binary type");
|
|
if (program_type != CL_PROGRAM_BINARY_TYPE_EXECUTABLE)
|
|
{
|
|
log_error("ERROR: Expected program type of a newly build executable to "
|
|
"be CL_PROGRAM_BINARY_TYPE_EXECUTABLE (in %s:%d)\n",
|
|
__FILE__, __LINE__);
|
|
return -1;
|
|
}
|
|
program_type = -1;
|
|
|
|
// Get the size of the resulting binary (only one device)
|
|
error = clGetProgramInfo(fully_linked_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 "
|
|
"%zu 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(fully_linked_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(fully_linked_program, CL_PROGRAM_BINARIES,
|
|
sizeof(buffers), &buffers, NULL);
|
|
test_error(error, "Unable to get program binary");
|
|
|
|
// use clCreateProgramWithBinary
|
|
fully_linked_program_with_binary = clCreateProgramWithBinary(
|
|
context, 1, &device, &binarySize, (const unsigned char **)buffers,
|
|
loadErrors, &error);
|
|
test_error(error, "Unable to create program with binary");
|
|
|
|
error = clGetProgramBuildInfo(
|
|
fully_linked_program_with_binary, device, CL_PROGRAM_BINARY_TYPE,
|
|
sizeof(cl_program_binary_type), &program_type, NULL);
|
|
test_error(error, "Unable to get program binary type");
|
|
if (program_type != CL_PROGRAM_BINARY_TYPE_EXECUTABLE)
|
|
{
|
|
log_error("ERROR: Expected program type of a program created from "
|
|
"a fully linked executable binary to be "
|
|
"CL_PROGRAM_BINARY_TYPE_EXECUTABLE (in %s:%d)\n",
|
|
__FILE__, __LINE__);
|
|
return -1;
|
|
}
|
|
program_type = -1;
|
|
free(binary);
|
|
}
|
|
|
|
error = clBuildProgram(fully_linked_program_with_binary, 1, &device, NULL,
|
|
NULL, NULL);
|
|
test_error(error, "Unable to build a simple program");
|
|
|
|
cl_kernel kernel =
|
|
clCreateKernel(fully_linked_program_with_binary, "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_with_binary,
|
|
"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");
|
|
|
|
/* Oh, one more thing. Steve Jobs and apparently Herb Sutter. The question
|
|
* is "Who is copying whom?" */
|
|
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 = clBuildProgram(program, 1, &device, NULL, NULL, NULL);
|
|
test_error(error, "Unable to build a simple program");
|
|
error = clGetProgramBuildInfo(program, device, CL_PROGRAM_BINARY_TYPE,
|
|
sizeof(cl_program_binary_type), &program_type,
|
|
NULL);
|
|
test_error(error, "Unable to get program binary type");
|
|
if (program_type != CL_PROGRAM_BINARY_TYPE_EXECUTABLE)
|
|
{
|
|
log_error(
|
|
"ERROR: Expected program type of a program created from compiled "
|
|
"object to be CL_PROGRAM_BINARY_TYPE_EXECUTABLE (in %s:%d)\n",
|
|
__FILE__, __LINE__);
|
|
return -1;
|
|
}
|
|
program_type = -1;
|
|
|
|
/* All's well that ends well. William Shakespeare */
|
|
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");
|
|
|
|
error = clReleaseProgram(fully_linked_program_with_binary);
|
|
test_error(error, "Unable to release program object");
|
|
|
|
error = clReleaseProgram(program_with_binary);
|
|
test_error(error, "Unable to release program object");
|
|
|
|
return 0;
|
|
}
|
|
|
|
volatile int compileNotificationSent;
|
|
|
|
void CL_CALLBACK test_notify_compile_complete(cl_program program,
|
|
void *userData)
|
|
{
|
|
if (userData == NULL || strcmp((char *)userData, "compilation") != 0)
|
|
{
|
|
log_error("ERROR: User data passed in to compile notify function was "
|
|
"not correct! (in %s:%d)\n",
|
|
__FILE__, __LINE__);
|
|
compileNotificationSent = -1;
|
|
}
|
|
else
|
|
compileNotificationSent = 1;
|
|
log_info("\n <-- program successfully compiled\n");
|
|
}
|
|
|
|
volatile int libraryCreationNotificationSent;
|
|
|
|
void CL_CALLBACK test_notify_create_library_complete(cl_program program,
|
|
void *userData)
|
|
{
|
|
if (userData == NULL || strcmp((char *)userData, "create library") != 0)
|
|
{
|
|
log_error("ERROR: User data passed in to library creation notify "
|
|
"function was not correct! (in %s:%d)\n",
|
|
__FILE__, __LINE__);
|
|
libraryCreationNotificationSent = -1;
|
|
}
|
|
else
|
|
libraryCreationNotificationSent = 1;
|
|
log_info("\n <-- library successfully created\n");
|
|
}
|
|
|
|
volatile int linkNotificationSent;
|
|
|
|
void CL_CALLBACK test_notify_link_complete(cl_program program, void *userData)
|
|
{
|
|
if (userData == NULL || strcmp((char *)userData, "linking") != 0)
|
|
{
|
|
log_error("ERROR: User data passed in to link notify function was not "
|
|
"correct! (in %s:%d)\n",
|
|
__FILE__, __LINE__);
|
|
linkNotificationSent = -1;
|
|
}
|
|
else
|
|
linkNotificationSent = 1;
|
|
log_info("\n <-- program successfully linked\n");
|
|
}
|
|
|
|
static int test_large_compile_and_link_status_options_log(
|
|
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];
|
|
char *compile_log;
|
|
char *compile_options;
|
|
char *library_log;
|
|
char *library_options;
|
|
char *linking_log;
|
|
char *linking_options;
|
|
cl_build_status status;
|
|
size_t size_ret;
|
|
|
|
compileNotificationSent = libraryCreationNotificationSent =
|
|
linkNotificationSent = 0;
|
|
|
|
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, &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;
|
|
}
|
|
|
|
/* Lets check that the compilation status is CL_BUILD_NONE */
|
|
error = clGetProgramBuildInfo(program, deviceID, CL_PROGRAM_BUILD_STATUS,
|
|
sizeof(status), &status, NULL);
|
|
test_error(error, "Unable to get program compile status");
|
|
if (status != CL_BUILD_NONE)
|
|
{
|
|
log_error("ERROR: Expected compile status to be CL_BUILD_NONE prior to "
|
|
"the beginning of the compilation! (status: %d in %s:%d)\n",
|
|
(int)status, __FILE__, __LINE__);
|
|
return -1;
|
|
}
|
|
|
|
/* Compile it */
|
|
error =
|
|
clCompileProgram(program, 1, &deviceID, NULL, 0, NULL, NULL,
|
|
test_notify_compile_complete, (void *)"compilation");
|
|
test_error(error, "Unable to compile a simple program");
|
|
|
|
/* Wait for compile to complete (just keep polling, since we're just a test
|
|
*/
|
|
error = clGetProgramBuildInfo(program, deviceID, CL_PROGRAM_BUILD_STATUS,
|
|
sizeof(status), &status, NULL);
|
|
test_error(error, "Unable to get program compile status");
|
|
|
|
while ((int)status == CL_BUILD_IN_PROGRESS)
|
|
{
|
|
log_info("\n -- still waiting for compile... (status is %d)", status);
|
|
sleep(1);
|
|
error =
|
|
clGetProgramBuildInfo(program, deviceID, CL_PROGRAM_BUILD_STATUS,
|
|
sizeof(status), &status, NULL);
|
|
test_error(error, "Unable to get program compile status");
|
|
}
|
|
if (status != CL_BUILD_SUCCESS)
|
|
{
|
|
log_error("ERROR: compile failed! (status: %d in %s:%d)\n", (int)status,
|
|
__FILE__, __LINE__);
|
|
return -1;
|
|
}
|
|
|
|
error = clGetProgramBuildInfo(program, deviceID, CL_PROGRAM_BUILD_LOG, 0,
|
|
NULL, &size_ret);
|
|
test_error(error, "Device failed to return compile log size");
|
|
compile_log = (char *)malloc(size_ret);
|
|
error = clGetProgramBuildInfo(program, deviceID, CL_PROGRAM_BUILD_LOG,
|
|
size_ret, compile_log, NULL);
|
|
if (error != CL_SUCCESS)
|
|
{
|
|
log_error("Device failed to return a compile log (in %s:%d)\n",
|
|
__FILE__, __LINE__);
|
|
test_error(error, "clGetProgramBuildInfo CL_PROGRAM_BUILD_LOG failed");
|
|
}
|
|
log_info("BUILD LOG: %s\n", compile_log);
|
|
free(compile_log);
|
|
|
|
error = clGetProgramBuildInfo(program, deviceID, CL_PROGRAM_BUILD_OPTIONS,
|
|
0, NULL, &size_ret);
|
|
test_error(error, "Device failed to return compile options size");
|
|
compile_options = (char *)malloc(size_ret);
|
|
error = clGetProgramBuildInfo(program, deviceID, CL_PROGRAM_BUILD_OPTIONS,
|
|
size_ret, compile_options, NULL);
|
|
test_error(
|
|
error,
|
|
"Device failed to return compile options.\nclGetProgramBuildInfo "
|
|
"CL_PROGRAM_BUILD_OPTIONS failed");
|
|
|
|
log_info("BUILD OPTIONS: %s\n", compile_options);
|
|
free(compile_options);
|
|
|
|
/* 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)",
|
|
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,
|
|
test_notify_create_library_complete, (void *)"create library", &error);
|
|
test_error(error, "Unable to create a multi-line library");
|
|
|
|
/* Wait for library creation to complete (just keep polling, since we're
|
|
* just a test */
|
|
error = clGetProgramBuildInfo(my_newly_minted_library, deviceID,
|
|
CL_PROGRAM_BUILD_STATUS, sizeof(status),
|
|
&status, NULL);
|
|
test_error(error, "Unable to get library creation link status");
|
|
|
|
while ((int)status == CL_BUILD_IN_PROGRESS)
|
|
{
|
|
log_info("\n -- still waiting for library creation... (status is %d)",
|
|
status);
|
|
sleep(1);
|
|
error = clGetProgramBuildInfo(my_newly_minted_library, deviceID,
|
|
CL_PROGRAM_BUILD_STATUS, sizeof(status),
|
|
&status, NULL);
|
|
test_error(error, "Unable to get library creation link status");
|
|
}
|
|
if (status != CL_BUILD_SUCCESS)
|
|
{
|
|
log_error("ERROR: library creation failed! (status: %d in %s:%d)\n",
|
|
(int)status, __FILE__, __LINE__);
|
|
return -1;
|
|
}
|
|
error = clGetProgramBuildInfo(my_newly_minted_library, deviceID,
|
|
CL_PROGRAM_BUILD_LOG, 0, NULL, &size_ret);
|
|
test_error(error, "Device failed to return a library creation log size");
|
|
library_log = (char *)malloc(size_ret);
|
|
error = clGetProgramBuildInfo(my_newly_minted_library, deviceID,
|
|
CL_PROGRAM_BUILD_LOG, size_ret, library_log,
|
|
NULL);
|
|
if (error != CL_SUCCESS)
|
|
{
|
|
log_error("Device failed to return a library creation log (in %s:%d)\n",
|
|
__FILE__, __LINE__);
|
|
test_error(error, "clGetProgramBuildInfo CL_PROGRAM_BUILD_LOG failed");
|
|
}
|
|
log_info("CREATE LIBRARY LOG: %s\n", library_log);
|
|
free(library_log);
|
|
|
|
error = clGetProgramBuildInfo(my_newly_minted_library, deviceID,
|
|
CL_PROGRAM_BUILD_OPTIONS, 0, NULL, &size_ret);
|
|
test_error(error, "Device failed to return library creation options size");
|
|
library_options = (char *)malloc(size_ret);
|
|
error = clGetProgramBuildInfo(my_newly_minted_library, deviceID,
|
|
CL_PROGRAM_BUILD_OPTIONS, size_ret,
|
|
library_options, NULL);
|
|
test_error(
|
|
error,
|
|
"Device failed to return library creation "
|
|
"options.\nclGetProgramBuildInfo CL_PROGRAM_BUILD_OPTIONS failed");
|
|
|
|
log_info("CREATE LIBRARY OPTIONS: %s\n", library_options);
|
|
free(library_options);
|
|
|
|
/* 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,
|
|
test_notify_link_complete, (void *)"linking", &error);
|
|
test_error(error, "Unable to link a program with a library");
|
|
|
|
/* Wait for linking to complete (just keep polling, since we're just a test
|
|
*/
|
|
error = clGetProgramBuildInfo(my_newly_linked_program, deviceID,
|
|
CL_PROGRAM_BUILD_STATUS, sizeof(status),
|
|
&status, NULL);
|
|
test_error(error, "Unable to get program link status");
|
|
|
|
while ((int)status == CL_BUILD_IN_PROGRESS)
|
|
{
|
|
log_info("\n -- still waiting for program linking... (status is %d)",
|
|
status);
|
|
sleep(1);
|
|
error = clGetProgramBuildInfo(my_newly_linked_program, deviceID,
|
|
CL_PROGRAM_BUILD_STATUS, sizeof(status),
|
|
&status, NULL);
|
|
test_error(error, "Unable to get program link status");
|
|
}
|
|
if (status != CL_BUILD_SUCCESS)
|
|
{
|
|
log_error("ERROR: program linking failed! (status: %d in %s:%d)\n",
|
|
(int)status, __FILE__, __LINE__);
|
|
return -1;
|
|
}
|
|
error = clGetProgramBuildInfo(my_newly_linked_program, deviceID,
|
|
CL_PROGRAM_BUILD_LOG, 0, NULL, &size_ret);
|
|
test_error(error, "Device failed to return a linking log size");
|
|
linking_log = (char *)malloc(size_ret);
|
|
error = clGetProgramBuildInfo(my_newly_linked_program, deviceID,
|
|
CL_PROGRAM_BUILD_LOG, size_ret, linking_log,
|
|
NULL);
|
|
if (error != CL_SUCCESS)
|
|
{
|
|
log_error("Device failed to return a linking log (in %s:%d).\n",
|
|
__FILE__, __LINE__);
|
|
test_error(error, "clGetProgramBuildInfo CL_PROGRAM_BUILD_LOG failed");
|
|
}
|
|
log_info("BUILDING LOG: %s\n", linking_log);
|
|
free(linking_log);
|
|
|
|
error = clGetProgramBuildInfo(my_newly_linked_program, deviceID,
|
|
CL_PROGRAM_BUILD_OPTIONS, 0, NULL, &size_ret);
|
|
test_error(error, "Device failed to return linking options size");
|
|
linking_options = (char *)malloc(size_ret);
|
|
error = clGetProgramBuildInfo(my_newly_linked_program, deviceID,
|
|
CL_PROGRAM_BUILD_OPTIONS, size_ret,
|
|
linking_options, NULL);
|
|
test_error(
|
|
error,
|
|
"Device failed to return linking options.\nclGetProgramBuildInfo "
|
|
"CL_PROGRAM_BUILD_OPTIONS failed");
|
|
|
|
log_info("BUILDING OPTIONS: %s\n", linking_options);
|
|
free(linking_options);
|
|
|
|
// 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 = clReleaseKernel(kernel);
|
|
test_error(error, "Unable to release kernel object");
|
|
|
|
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 = 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;
|
|
}
|
|
|
|
REGISTER_TEST(compile_and_link_status_options_log)
|
|
{
|
|
unsigned int toTest[] = { 256, 0 }; // 512, 1024, 8192, 16384, 32768, 0 };
|
|
unsigned int i;
|
|
|
|
log_info("Testing Compile and Link Status, Options and Logging ...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_compile_and_link_status_options_log(context, device,
|
|
queue, toTest[i])
|
|
!= 0)
|
|
{
|
|
log_error(
|
|
"ERROR: large program compilation, linking, status, options "
|
|
"and logging 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;
|
|
}
|