blob: 6e99d462ced4071b05bf62ee1a1780d2199071bf [file] [log] [blame]
//
// Copyright (c) 2017 The Khronos Group Inc.
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
//
#include "harness/compat.h"
#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#include <sys/types.h>
#include <sys/stat.h>
#include "procs.h"
const char *kernel_code = R"(
__kernel void test_kernel(
char%s c, uchar%s uc, short%s s, ushort%s us, int%s i, uint%s ui, float%s f,
__global float%s *result)
{
result[0] = %s(c);
result[1] = %s(uc);
result[2] = %s(s);
result[3] = %s(us);
result[4] = %s(i);
result[5] = %s(ui);
result[6] = f;
})";
const char *kernel_code_long = R"(
__kernel void test_kernel_long(
long%s l, ulong%s ul,
__global float%s *result)
{
result[0] = %s(l);
result[1] = %s(ul);
})";
int test_parameter_types_long(cl_device_id device, cl_context context,
cl_command_queue queue, int num_elements)
{
clMemWrapper results;
int error;
size_t global[3] = { 1, 1, 1 };
float results_back[2 * 16];
int count, index;
const char *types[] = { "long", "ulong" };
char kernel_string[8192];
int sizes[] = { 1, 2, 4, 8, 16 };
const char *size_strings[] = { "", "2", "4", "8", "16" };
float expected;
int total_errors = 0;
int size_to_test;
char *ptr;
char convert_string[1024];
size_t max_parameter_size;
// We don't really care about the contents since we're just testing that the
// types work.
cl_long l[16] = { -21, -1, 2, -3, 4, -5, 6, -7,
8, -9, 10, -11, 12, -13, 14, -15 };
cl_ulong ul[16] = { 22, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15 };
// Calculate how large our paramter size is to the kernel
size_t parameter_size = sizeof(cl_long) + sizeof(cl_ulong);
// Init our strings.
kernel_string[0] = '\0';
convert_string[0] = '\0';
// Get the maximum parameter size allowed
error =
clGetDeviceInfo(device, CL_DEVICE_MAX_PARAMETER_SIZE,
sizeof(max_parameter_size), &max_parameter_size, NULL);
test_error(error, "Unable to get max parameter size from device");
// Create the results buffer
results = clCreateBuffer(context, CL_MEM_READ_WRITE,
sizeof(cl_float) * 2 * 16, NULL, &error);
test_error(error, "clCreateBuffer failed");
// Go over all the vector sizes
for (size_to_test = 0; size_to_test < 5; size_to_test++)
{
clProgramWrapper program;
clKernelWrapper kernel;
size_t total_parameter_size =
parameter_size * sizes[size_to_test] + sizeof(cl_mem);
if (total_parameter_size > max_parameter_size)
{
log_info(
"Can not test with vector size %d because it would exceed the "
"maximum allowed parameter size to the kernel. (%d > %d)\n",
(int)sizes[size_to_test], (int)total_parameter_size,
(int)max_parameter_size);
continue;
}
log_info("Testing vector size %d\n", sizes[size_to_test]);
// If size is > 1, then we need a explicit convert call.
if (sizes[size_to_test] > 1)
{
sprintf(convert_string, "convert_float%s",
size_strings[size_to_test]);
}
else
{
sprintf(convert_string, " ");
}
// Build the kernel
sprintf(kernel_string, kernel_code_long, size_strings[size_to_test],
size_strings[size_to_test], size_strings[size_to_test],
convert_string, convert_string);
ptr = kernel_string;
error = create_single_kernel_helper(context, &program, &kernel, 1,
(const char **)&ptr,
"test_kernel_long");
test_error(error, "create single kernel failed");
// Set the arguments
for (count = 0; count < 2; count++)
{
switch (count)
{
case 0:
error = clSetKernelArg(
kernel, count, sizeof(cl_long) * sizes[size_to_test],
&l);
break;
case 1:
error = clSetKernelArg(
kernel, count, sizeof(cl_ulong) * sizes[size_to_test],
&ul);
break;
default: log_error("Test error"); break;
}
if (error)
log_error("Setting kernel arg %d %s%s: ", count, types[count],
size_strings[size_to_test]);
test_error(error, "clSetKernelArgs failed");
}
error = clSetKernelArg(kernel, 2, sizeof(cl_mem), &results);
test_error(error, "clSetKernelArgs failed");
// Execute
error = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, global, NULL, 0,
NULL, NULL);
test_error(error, "clEnqueueNDRangeKernel failed");
// Read back the results
error = clEnqueueReadBuffer(queue, results, CL_TRUE, 0,
sizeof(cl_float) * 2 * 16, results_back, 0,
NULL, NULL);
test_error(error, "clEnqueueReadBuffer failed");
// Verify the results
for (count = 0; count < 2; count++)
{
for (index = 0; index < sizes[size_to_test]; index++)
{
switch (count)
{
case 0: expected = (float)l[index]; break;
case 1: expected = (float)ul[index]; break;
default: log_error("Test error"); break;
}
if (results_back[count * sizes[size_to_test] + index]
!= expected)
{
total_errors++;
log_error("Conversion from %s%s failed: index %d got %g, "
"expected %g.\n",
types[count], size_strings[size_to_test], index,
results_back[count * sizes[size_to_test] + index],
expected);
}
}
}
}
return total_errors;
}
int test_parameter_types(cl_device_id device, cl_context context,
cl_command_queue queue, int num_elements)
{
clMemWrapper results;
int error;
size_t global[3] = { 1, 1, 1 };
float results_back[7 * 16];
int count, index;
const char *types[] = { "char", "uchar", "short", "ushort",
"int", "uint", "float" };
char kernel_string[8192];
int sizes[] = { 1, 2, 4, 8, 16 };
const char *size_strings[] = { "", "2", "4", "8", "16" };
float expected;
int total_errors = 0;
int size_to_test;
char *ptr;
char convert_string[1024];
size_t max_parameter_size;
// We don't really care about the contents since we're just testing that the
// types work.
cl_char c[16] = { 0, -1, 2, -3, 4, -5, 6, -7,
8, -9, 10, -11, 12, -13, 14, -15 };
cl_uchar uc[16] = { 16, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15 };
cl_short s[16] = { -17, -1, 2, -3, 4, -5, 6, -7,
8, -9, 10, -11, 12, -13, 14, -15 };
cl_ushort us[16] = {
18, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15
};
cl_int i[16] = { -19, -1, 2, -3, 4, -5, 6, -7,
8, -9, 10, -11, 12, -13, 14, -15 };
cl_uint ui[16] = { 20, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15 };
cl_float f[16] = { -23, -1, 2, -3, 4, -5, 6, -7,
8, -9, 10, -11, 12, -13, 14, -15 };
// Calculate how large our paramter size is to the kernel
size_t parameter_size = sizeof(cl_char) + sizeof(cl_uchar)
+ sizeof(cl_short) + sizeof(cl_ushort) + sizeof(cl_int)
+ sizeof(cl_uint) + sizeof(cl_float);
// Init our strings.
kernel_string[0] = '\0';
convert_string[0] = '\0';
// Get the maximum parameter size allowed
error =
clGetDeviceInfo(device, CL_DEVICE_MAX_PARAMETER_SIZE,
sizeof(max_parameter_size), &max_parameter_size, NULL);
test_error(error, "Unable to get max parameter size from device");
// Create the results buffer
results = clCreateBuffer(context, CL_MEM_READ_WRITE,
sizeof(cl_float) * 7 * 16, NULL, &error);
test_error(error, "clCreateBuffer failed");
// Go over all the vector sizes
for (size_to_test = 0; size_to_test < 5; size_to_test++)
{
clProgramWrapper program;
clKernelWrapper kernel;
size_t total_parameter_size =
parameter_size * sizes[size_to_test] + sizeof(cl_mem);
if (total_parameter_size > max_parameter_size)
{
log_info(
"Can not test with vector size %d because it would exceed the "
"maximum allowed parameter size to the kernel. (%d > %d)\n",
(int)sizes[size_to_test], (int)total_parameter_size,
(int)max_parameter_size);
continue;
}
log_info("Testing vector size %d\n", sizes[size_to_test]);
// If size is > 1, then we need a explicit convert call.
if (sizes[size_to_test] > 1)
{
sprintf(convert_string, "convert_float%s",
size_strings[size_to_test]);
}
else
{
sprintf(convert_string, " ");
}
// Build the kernel
sprintf(kernel_string, kernel_code, size_strings[size_to_test],
size_strings[size_to_test], size_strings[size_to_test],
size_strings[size_to_test], size_strings[size_to_test],
size_strings[size_to_test], size_strings[size_to_test],
size_strings[size_to_test], convert_string, convert_string,
convert_string, convert_string, convert_string, convert_string);
ptr = kernel_string;
error = create_single_kernel_helper(context, &program, &kernel, 1,
(const char **)&ptr, "test_kernel");
test_error(error, "create single kernel failed");
// Set the arguments
for (count = 0; count < 7; count++)
{
switch (count)
{
case 0:
error = clSetKernelArg(
kernel, count, sizeof(cl_char) * sizes[size_to_test],
&c);
break;
case 1:
error = clSetKernelArg(
kernel, count, sizeof(cl_uchar) * sizes[size_to_test],
&uc);
break;
case 2:
error = clSetKernelArg(
kernel, count, sizeof(cl_short) * sizes[size_to_test],
&s);
break;
case 3:
error = clSetKernelArg(
kernel, count, sizeof(cl_ushort) * sizes[size_to_test],
&us);
break;
case 4:
error = clSetKernelArg(kernel, count,
sizeof(cl_int) * sizes[size_to_test],
&i);
break;
case 5:
error = clSetKernelArg(
kernel, count, sizeof(cl_uint) * sizes[size_to_test],
&ui);
break;
case 6:
error = clSetKernelArg(
kernel, count, sizeof(cl_float) * sizes[size_to_test],
&f);
break;
default: log_error("Test error"); break;
}
if (error)
log_error("Setting kernel arg %d %s%s: ", count, types[count],
size_strings[size_to_test]);
test_error(error, "clSetKernelArgs failed");
}
error = clSetKernelArg(kernel, 7, sizeof(cl_mem), &results);
test_error(error, "clSetKernelArgs failed");
// Execute
error = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, global, NULL, 0,
NULL, NULL);
test_error(error, "clEnqueueNDRangeKernel failed");
// Read back the results
error = clEnqueueReadBuffer(queue, results, CL_TRUE, 0,
sizeof(cl_float) * 7 * 16, results_back, 0,
NULL, NULL);
test_error(error, "clEnqueueReadBuffer failed");
// Verify the results
for (count = 0; count < 7; count++)
{
for (index = 0; index < sizes[size_to_test]; index++)
{
switch (count)
{
case 0: expected = (float)c[index]; break;
case 1: expected = (float)uc[index]; break;
case 2: expected = (float)s[index]; break;
case 3: expected = (float)us[index]; break;
case 4: expected = (float)i[index]; break;
case 5: expected = (float)ui[index]; break;
case 6: expected = (float)f[index]; break;
default: log_error("Test error"); break;
}
if (results_back[count * sizes[size_to_test] + index]
!= expected)
{
total_errors++;
log_error("Conversion from %s%s failed: index %d got %g, "
"expected %g.\n",
types[count], size_strings[size_to_test], index,
results_back[count * sizes[size_to_test] + index],
expected);
}
}
}
}
if (gHasLong)
{
log_info("Testing long types...\n");
total_errors +=
test_parameter_types_long(device, context, queue, num_elements);
}
else
{
log_info("Longs unsupported, skipping.");
}
return total_errors;
}