ahmed/OpenCL-CTS
1
0
Fork 0
mirror of https://github.com/KhronosGroup/OpenCL-CTS.git synced 2026-03-19 06:09:01 +00:00
Code Issues Packages Projects Releases Wiki Activity
Files
e1cf741f1c800db5e27d1be4cf7647429d8d74ea
OpenCL-CTS/test_conformance/generic_address_space/advanced_tests.cpp
Kévin Petit 5e2f5b857e Accept OpenCL 3.0 in version parsing code and use where appropriate (#752) 
* Accept OpenCL 3.0 in version parsing code and use where appropriate

There were a number of tests against 2.2 that are clearer against 3.0.

Fixes #751

Signed-off-by: Kévin Petit <kpet@free.fr>

* Remove CL_EXPERIMENTAL guards

Signed-off-by: Kévin Petit <kpet@free.fr>

* formatting

* Configure the headers for OpenCL 3.0

* more format fixes
2020-04-30 11:51:06 +01:00

999 lines
37 KiB
C++
Raw Blame History

//
// 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/testHarness.h"
#include "harness/typeWrappers.h"
#include "base.h"
#include <string>
#include <vector>
#include <algorithm>
#include <sstream>
typedef enum {
ARG_TYPE_NONE,
ARG_TYPE_HOST_PTR,
ARG_TYPE_HOST_LOCAL,
ARG_TYPE_COARSE_GRAINED_SVM,
ARG_TYPE_FINE_GRAINED_BUFFER_SVM,
ARG_TYPE_FINE_GRAINED_SYSTEM_SVM,
ARG_TYPE_ATOMICS_SVM
} ExtraKernelArgMemType;
class CSVMWrapper {
public:
CSVMWrapper() : ptr_(NULL), context_(NULL) { }
void Attach(cl_context context, void *ptr) {
context_ = context;
ptr_ = ptr;
}
~CSVMWrapper() {
if (ptr_)
clSVMFree(context_, ptr_);
}
operator void *() {
return ptr_;
}
private:
void *ptr_;
cl_context context_;
};
class CAdvancedTest : public CTest {
public:
CAdvancedTest(const std::vector<std::string>& kernel) : CTest(), _kernels(kernel), _extraKernelArgMemType(ARG_TYPE_NONE) {
}
CAdvancedTest(const std::string& library, const std::vector<std::string>& kernel) : CTest(), _libraryCode(library), _kernels(kernel), _extraKernelArgMemType(ARG_TYPE_NONE) {
}
CAdvancedTest(const std::string& kernel, ExtraKernelArgMemType argType = ARG_TYPE_NONE) : CTest(), _kernels(1, kernel), _extraKernelArgMemType(argType) {
}
CAdvancedTest(const std::string& library, const std::string& kernel) : CTest(), _libraryCode(library), _kernels(1, kernel), _extraKernelArgMemType(ARG_TYPE_NONE) {
}
int PrintCompilationLog(cl_program program, cl_device_id device) {
cl_int error;
size_t buildLogSize = 0;
error = clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, 0, NULL, &buildLogSize);
test_error(error, "clGetProgramBuildInfo failed");
std::string log;
log.resize(buildLogSize);
error = clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, buildLogSize, &log[0], NULL);
test_error(error, "clGetProgramBuildInfo failed");
log_error("Build log for device is:\n------------\n");
log_error("%s\n", log.c_str() );
log_error( "\n----------\n" );
return CL_SUCCESS;
}
int ExecuteSubcase(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements, const std::string& src) {
cl_int error;
clProgramWrapper program, preCompiledLibrary, library, finalProgram;
clKernelWrapper kernel;
const char *srcPtr = src.c_str();
if (!_libraryCode.empty()) {
program = clCreateProgramWithSource(context, 1, &srcPtr, NULL, &error);
test_error(error, "clCreateProgramWithSource failed");
error = clCompileProgram(program, 1, &deviceID, "-cl-std=CL2.0", 0, NULL, NULL, NULL, NULL);
if (error != CL_SUCCESS)
PrintCompilationLog(program, deviceID);
test_error(error, "clCompileProgram failed");
const char *srcPtrLibrary = _libraryCode.c_str();
preCompiledLibrary = clCreateProgramWithSource(context, 1, &srcPtrLibrary, NULL, &error);
test_error(error, "clCreateProgramWithSource failed");
error = clCompileProgram(preCompiledLibrary, 1, &deviceID, "-cl-std=CL2.0", 0, NULL, NULL, NULL, NULL);
if (error != CL_SUCCESS)
PrintCompilationLog(preCompiledLibrary, deviceID);
test_error(error, "clCompileProgram failed");
library = clLinkProgram(context, 1, &deviceID, "-create-library", 1, &preCompiledLibrary, NULL, NULL, &error);
test_error(error, "clLinkProgram failed");
cl_program objects[] = { program, library };
finalProgram = clLinkProgram(context, 1, &deviceID, "", 2, objects, NULL, NULL, &error);
test_error(error, "clLinkProgram failed");
kernel = clCreateKernel(finalProgram, "testKernel", &error);
test_error(error, "clCreateKernel failed");
}
else {
if (create_single_kernel_helper_with_build_options(context, &program, &kernel, 1, &srcPtr, "testKernel", "-cl-std=CL2.0")) {
log_error("create_single_kernel_helper failed\n");
return -1;
}
}
size_t bufferSize = num_elements * sizeof(cl_uint);
clMemWrapper buffer = clCreateBuffer(context, CL_MEM_WRITE_ONLY, bufferSize, NULL, &error);
test_error(error, "clCreateBuffer failed");
error = clSetKernelArg(kernel, 0, sizeof(buffer), &buffer);
test_error(error, "clSetKernelArg(0) failed");
// Warning: the order below is very important as SVM buffer cannot be free'd before corresponding mem_object
CSVMWrapper svmWrapper;
clMemWrapper extraArg;
std::vector<cl_uint> extraArgData(num_elements);
for (cl_uint i = 0; i < (cl_uint)num_elements; i++)
extraArgData[i] = i;
if (_extraKernelArgMemType != ARG_TYPE_NONE) {
if (_extraKernelArgMemType == ARG_TYPE_HOST_PTR) {
extraArg = clCreateBuffer(context, CL_MEM_USE_HOST_PTR | CL_MEM_READ_WRITE, bufferSize, &extraArgData[0], &error);
test_error(error, "clCreateBuffer failed");
}
else {
void *ptr = NULL;
switch (_extraKernelArgMemType) {
case ARG_TYPE_COARSE_GRAINED_SVM:
ptr = clSVMAlloc(context, CL_MEM_READ_WRITE, bufferSize, 0);
break;
case ARG_TYPE_FINE_GRAINED_BUFFER_SVM:
ptr = clSVMAlloc(context, CL_MEM_SVM_FINE_GRAIN_BUFFER | CL_MEM_READ_WRITE, bufferSize, 0);
break;
case ARG_TYPE_FINE_GRAINED_SYSTEM_SVM:
ptr = &extraArgData[0];
break;
case ARG_TYPE_ATOMICS_SVM:
ptr = clSVMAlloc(context, CL_MEM_SVM_FINE_GRAIN_BUFFER | CL_MEM_SVM_ATOMICS | CL_MEM_READ_WRITE, bufferSize, 0);
break;
default:
break;
}
if(_extraKernelArgMemType != ARG_TYPE_HOST_LOCAL) {
if (!ptr) {
log_error("Allocation failed\n");
return -1;
}
if (_extraKernelArgMemType != ARG_TYPE_FINE_GRAINED_SYSTEM_SVM) {
svmWrapper.Attach(context, ptr);
}
if (_extraKernelArgMemType == ARG_TYPE_COARSE_GRAINED_SVM) {
error = clEnqueueSVMMap(queue, CL_TRUE, CL_MAP_WRITE, ptr, bufferSize, 0, NULL, NULL);
test_error(error, "clEnqueueSVMMap failed");
}
memcpy(ptr, &extraArgData[0], bufferSize);
if (_extraKernelArgMemType == ARG_TYPE_COARSE_GRAINED_SVM) {
error = clEnqueueSVMUnmap(queue, ptr, 0, NULL, NULL);
test_error(error, "clEnqueueSVMUnmap failed");
clFinish(queue);
}
extraArg = clCreateBuffer(context, CL_MEM_USE_HOST_PTR | CL_MEM_READ_WRITE, bufferSize, ptr, &error);
test_error(error, "clCreateBuffer from SVM buffer failed");
}
}
if(_extraKernelArgMemType == ARG_TYPE_HOST_LOCAL)
error = clSetKernelArg(kernel, 1, bufferSize, NULL);
else
error = clSetKernelArg(kernel, 1, sizeof(extraArg), &extraArg);
test_error(error, "clSetKernelArg(1) failed");
}
size_t globalWorkGroupSize = num_elements;
size_t localWorkGroupSize = 0;
error = get_max_common_work_group_size(context, kernel, globalWorkGroupSize, &localWorkGroupSize);
test_error(error, "Unable to get common work group size");
error = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &globalWorkGroupSize, &localWorkGroupSize, 0, NULL, NULL);
test_error(error, "clEnqueueNDRangeKernel failed");
// verify results
std::vector<cl_uint> results(num_elements);
error = clEnqueueReadBuffer(queue, buffer, CL_TRUE, 0, bufferSize, &results[0], 0, NULL, NULL);
test_error(error, "clEnqueueReadBuffer failed");
size_t passCount = std::count(results.begin(), results.end(), 1);
if (passCount != results.size()) {
std::vector<cl_uint>::iterator iter = std::find(results.begin(), results.end(), 0);
log_error("Verification on device failed at index %ld\n", std::distance(results.begin(), iter));
log_error("%ld out of %ld failed\n", (results.size()-passCount), results.size());
return -1;
}
return CL_SUCCESS;
}
int Execute(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
cl_int result = CL_SUCCESS;
for (std::vector<std::string>::const_iterator it = _kernels.begin(); it != _kernels.end(); ++it) {
log_info("Executing subcase #%ld out of %ld\n", (it - _kernels.begin() + 1), _kernels.size());
result |= ExecuteSubcase(deviceID, context, queue, num_elements, *it);
}
return result;
}
private:
const std::string _libraryCode;
const std::vector<std::string> _kernels;
const ExtraKernelArgMemType _extraKernelArgMemType;
};
int test_library_function(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
const std::string LIBRARY_FUNCTION = common::CONFORMANCE_VERIFY_FENCE +
NL
NL "bool helperFunction(float *floatp, float val) {"
NL " if (!isFenceValid(get_fence(floatp)))"
NL " return false;"
NL
NL " if (*floatp != val)"
NL " return false;"
NL
NL " return true;"
NL "}"
NL;
const std::string KERNEL_FUNCTION =
NL
NL "extern bool helperFunction(float *floatp, float val);"
NL
NL "__global float gfloat = 1.0f;"
NL
NL "__kernel void testKernel(__global uint *results) {"
NL " uint tid = get_global_id(0);"
NL
NL " __global float *gfloatp = &gfloat;"
NL " __local float lfloat;"
NL " lfloat = 2.0f;"
NL " __local float *lfloatp = &lfloat;"
NL " float pfloat = 3.0f;"
NL " __private float *pfloatp = &pfloat;"
NL
NL " uint failures = 0;"
NL
NL " failures += helperFunction(gfloatp, gfloat) ? 0 : 1;"
NL " failures += helperFunction(lfloatp, lfloat) ? 0 : 1;"
NL " failures += helperFunction(pfloatp, pfloat) ? 0 : 1;"
NL
NL " results[tid] = failures == 0;"
NL "}"
NL;
CAdvancedTest test(LIBRARY_FUNCTION, KERNEL_FUNCTION);
return test.Execute(deviceID, context, queue, num_elements);
}
int test_generic_variable_volatile(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
std::vector<std::string> KERNEL_FUNCTIONS;
KERNEL_FUNCTIONS.push_back(common::CONFORMANCE_VERIFY_FENCE +
NL
NL "bool helperFunction(float *floatp, float val) {"
NL " if (!isFenceValid(get_fence(floatp)))"
NL " return false;"
NL
NL " if (*floatp != val)"
NL " return false;"
NL
NL " return true;"
NL "}"
NL
NL "__kernel void testKernel(__global uint *results) {"
NL " uint tid = get_global_id(0);"
NL
NL " static __global float val;"
NL " val = 0.1f;"
NL " float * volatile ptr = &val;"
NL
NL " results[tid] = helperFunction(ptr, val);"
NL "}"
NL
);
KERNEL_FUNCTIONS.push_back(common::CONFORMANCE_VERIFY_FENCE +
NL
NL "bool helperFunction(float *floatp, float val) {"
NL " if (!isFenceValid(get_fence(floatp)))"
NL " return false;"
NL
NL " if (*floatp != val)"
NL " return false;"
NL
NL " return true;"
NL "}"
NL
NL "__kernel void testKernel(__global uint *results) {"
NL " uint tid = get_global_id(0);"
NL
NL " __local float val;"
NL " val = 0.1f;"
NL " float * ptr = &val;"
NL
NL " results[tid] = helperFunction(ptr, val);"
NL "}"
NL
);
KERNEL_FUNCTIONS.push_back(common::CONFORMANCE_VERIFY_FENCE +
NL
NL "bool helperFunction(float *floatp, float val) {"
NL " if (!isFenceValid(get_fence(floatp)))"
NL " return false;"
NL
NL " if (*floatp != val)"
NL " return false;"
NL
NL " return true;"
NL "}"
NL
NL "__kernel void testKernel(__global uint *results) {"
NL " uint tid = get_global_id(0);"
NL
NL " __private float val;"
NL " val = 0.1f;"
NL " float * volatile ptr = &val;"
NL
NL " results[tid] = helperFunction(ptr, val);"
NL "}"
NL
);
CAdvancedTest test(KERNEL_FUNCTIONS);
return test.Execute(deviceID, context, queue, num_elements);
}
int test_generic_variable_const(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
std::vector<std::string> KERNEL_FUNCTIONS;
KERNEL_FUNCTIONS.push_back(common::CONFORMANCE_VERIFY_FENCE +
NL
NL "bool helperFunction(const float *floatp, float val) {"
NL " if (!isFenceValid(get_fence(floatp)))"
NL " return false;"
NL
NL " if (*floatp != val)"
NL " return false;"
NL
NL " return true;"
NL "}"
NL
NL "__kernel void testKernel(__global uint *results) {"
NL " uint tid = get_global_id(0);"
NL
NL " const __private float val = 0.1f;"
NL " const float * ptr = &val;"
NL
NL " results[tid] = helperFunction(ptr, val);"
NL "}"
NL
);
KERNEL_FUNCTIONS.push_back(common::CONFORMANCE_VERIFY_FENCE +
NL
NL "bool helperFunction(const float *floatp, float val) {"
NL " if (!isFenceValid(get_fence(floatp)))"
NL " return false;"
NL
NL " if (*floatp != val)"
NL " return false;"
NL
NL " return true;"
NL "}"
NL
NL "__kernel void testKernel(__global uint *results) {"
NL " uint tid = get_global_id(0);"
NL
NL " const static __global float val = 0.1f;"
NL " const float * ptr = &val;"
NL
NL " results[tid] = helperFunction(ptr, val);"
NL "}"
NL
);
CAdvancedTest test(KERNEL_FUNCTIONS);
return test.Execute(deviceID, context, queue, num_elements);
}
int test_generic_variable_gentype(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
const std::string KERNEL_FUNCTION_TEMPLATE = common::CONFORMANCE_VERIFY_FENCE +
NL
NL "%s"
NL
NL "bool helperFunction(const %s *%sp, %s val) {"
NL " if (!isFenceValid(get_fence(%sp)))"
NL " return false;"
NL
NL " return %s(*%sp == val);"
NL "}"
NL
NL "__kernel void testKernel(__global uint *results) {"
NL " uint tid = get_global_id(0);"
NL
NL " %s %s val = (%s)1;"
NL " %s * ptr = &val;"
NL
NL " results[tid] = helperFunction(ptr, val);"
NL "}"
NL;
/* Qualcomm fix: 12502 Gen Addr Space - Fix kernel for generic variable gentype (half) test
const std::string KERNEL_FUNCTION_TEMPLATE_HALF = common::CONFORMANCE_VERIFY_FENCE */
const std::string vector_sizes[] = { "", "2", "3", "4", "8", "16" };
const std::string gentype_base[] = { "float", "char", "uchar", "short", "ushort", "int", "uint", "long", "ulong" };
const std::string gentype_others[] = { "bool", "size_t", "ptrdiff_t", "intptr_t", "uintptr_t" };
const std::string address_spaces[] = { "static __global", "__private" };
const std::string vector_cmp = "all";
std::vector<std::string> KERNEL_FUNCTIONS;
// Add base types plus theirs vector variants
for (size_t i = 0; i < sizeof(gentype_base) / sizeof(gentype_base[0]); i++) {
for (size_t j = 0; j < sizeof(vector_sizes) / sizeof(vector_sizes[0]); j++) {
for (size_t k = 0; k < sizeof(address_spaces) / sizeof(address_spaces[0]); k++) {
char temp_kernel[1024];
const std::string fulltype = gentype_base[i] + vector_sizes[j];
sprintf(temp_kernel, KERNEL_FUNCTION_TEMPLATE.c_str(),
"",
fulltype.c_str(), fulltype.c_str(), fulltype.c_str(), fulltype.c_str(),
(j > 0 ? vector_cmp.c_str() : ""),
fulltype.c_str(), address_spaces[k].c_str(), fulltype.c_str(), fulltype.c_str(),
fulltype.c_str());
KERNEL_FUNCTIONS.push_back(temp_kernel);
}
}
}
const std::string cl_khr_fp64_pragma = "#pragma OPENCL EXTENSION cl_khr_fp64 : enable";
// Add double floating types if they are supported
if (is_extension_available(deviceID, "cl_khr_fp64")) {
for (size_t j = 0; j < sizeof(vector_sizes) / sizeof(vector_sizes[0]); j++) {
for (size_t k = 0; k < sizeof(address_spaces) / sizeof(address_spaces[0]); k++) {
char temp_kernel[1024];
const std::string fulltype = std::string("double") + vector_sizes[j];
sprintf(temp_kernel, KERNEL_FUNCTION_TEMPLATE.c_str(),
cl_khr_fp64_pragma.c_str(),
fulltype.c_str(), fulltype.c_str(), fulltype.c_str(), fulltype.c_str(),
(j > 0 ? vector_cmp.c_str() : ""),
fulltype.c_str(), address_spaces[k].c_str(), fulltype.c_str(), fulltype.c_str(),
fulltype.c_str());
KERNEL_FUNCTIONS.push_back(temp_kernel);
}
}
}
/* Qualcomm fix: 12502 Gen Addr Space - Fix kernel for generic variable gentype (half) test */
const std::string cl_khr_fp16_pragma = "#pragma OPENCL EXTENSION cl_khr_fp16 : enable";
// Add half floating types if they are supported
if (is_extension_available(deviceID, "cl_khr_fp16")) {
for (size_t j = 0; j < sizeof(vector_sizes) / sizeof(vector_sizes[0]); j++) {
for (size_t k = 0; k < sizeof(address_spaces) / sizeof(address_spaces[0]); k++) {
char temp_kernel[1024];
const std::string fulltype = std::string("half") + vector_sizes[j];
sprintf(temp_kernel, KERNEL_FUNCTION_TEMPLATE.c_str(),
cl_khr_fp16_pragma.c_str(),
fulltype.c_str(), fulltype.c_str(), fulltype.c_str(), fulltype.c_str(),
(j > 0 ? vector_cmp.c_str() : ""),
fulltype.c_str(), address_spaces[k].c_str(), fulltype.c_str(), fulltype.c_str(),
fulltype.c_str());
/* Qualcomm fix: end */
KERNEL_FUNCTIONS.push_back(temp_kernel);
}
}
}
// Add other types that do not have vector variants
for (size_t i = 0; i < sizeof(gentype_others) / sizeof(gentype_others[0]); i++) {
for (size_t k = 0; k < sizeof(address_spaces) / sizeof(address_spaces[0]); k++) {
char temp_kernel[1024];
const std::string fulltype = gentype_others[i];
sprintf(temp_kernel, KERNEL_FUNCTION_TEMPLATE.c_str(),
"",
fulltype.c_str(), fulltype.c_str(), fulltype.c_str(), fulltype.c_str(),
"",
fulltype.c_str(), address_spaces[k].c_str(), fulltype.c_str(), fulltype.c_str(),
fulltype.c_str());
KERNEL_FUNCTIONS.push_back(temp_kernel);
}
}
CAdvancedTest test(KERNEL_FUNCTIONS);
return test.Execute(deviceID, context, queue, num_elements);
}
void create_math_kernels(std::vector<std::string>& KERNEL_FUNCTIONS) {
const std::string KERNEL_FUNCTION_TEMPLATE =
NL
NL "__kernel void testKernel(__global uint *results) {"
NL " uint tid = get_global_id(0);"
NL
NL " const %s param1 = %s;"
NL " %s param2_generic;"
NL " %s param2_reference;"
NL " %s * ptr = &param2_generic;"
NL " %s return_value_generic;"
NL " %s return_value_reference;"
NL
NL " return_value_generic = %s(param1, ptr);"
NL " return_value_reference = %s(param1, &param2_reference);"
NL
NL " results[tid] = (%s(*ptr == param2_reference) && %s(return_value_generic == return_value_reference));"
NL "}"
NL;
typedef struct {
std::string bulitin_name;
std::string base_gentype;
std::string pointer_gentype;
std::string first_param_value;
std::string compare_fn;
} BuiltinDescriptor;
BuiltinDescriptor builtins[] = {
{ "fract", "float", "float", "133.55f", "" },
{ "frexp", "float2", "int2", "(float2)(24.12f, 99999.7f)", "all" },
{ "frexp", "float", "int", "1234.5f", "" },
{ "lgamma_r", "float2", "int2", "(float2)(1000.0f, 9999.5f)", "all" },
{ "lgamma_r", "float", "int", "1000.0f", "" },
{ "modf", "float", "float", "1234.56789f", "" },
{ "sincos", "float", "float", "3.141592f", "" }
};
for (size_t i = 0; i < sizeof(builtins) / sizeof(builtins[0]); i++) {
char temp_kernel[1024];
sprintf(temp_kernel, KERNEL_FUNCTION_TEMPLATE.c_str(), builtins[i].base_gentype.c_str(), builtins[i].first_param_value.c_str(),
builtins[i].pointer_gentype.c_str(), builtins[i].pointer_gentype.c_str(), builtins[i].pointer_gentype.c_str(), builtins[i].base_gentype.c_str(),
builtins[i].base_gentype.c_str(), builtins[i].bulitin_name.c_str(), builtins[i].bulitin_name.c_str(),
builtins[i].compare_fn.c_str(), builtins[i].compare_fn.c_str());
KERNEL_FUNCTIONS.push_back(temp_kernel);
}
// add special case for remquo (3 params)
KERNEL_FUNCTIONS.push_back(
NL
NL "__kernel void testKernel(__global uint *results) {"
NL " uint tid = get_global_id(0);"
NL
NL " const float param1 = 1234.56789f;"
NL " const float param2 = 123.456789f;"
NL " int param3_generic;"
NL " int param3_reference;"
NL " int * ptr = &param3_generic;"
NL " float return_value_generic;"
NL " float return_value_reference;"
NL
NL " return_value_generic = remquo(param1, param2, ptr);"
NL " return_value_reference = remquo(param1, param2, &param3_reference);"
NL
NL " results[tid] = (*ptr == param3_reference && return_value_generic == return_value_reference);"
NL "}"
NL
);
}
std::string get_default_data_for_type(const std::string& type) {
std::string result;
if (type == "float") {
for (int i = 0; i < 10; i++) {
for (int j = 0; j < 10; j++) {
char temp[10];
sprintf(temp, "%d.%df, ", i, j);
result += std::string(temp);
}
}
}
else if (type == "double") {
for (int i = 0; i < 10; i++) {
for (int j = 0; j < 10; j++) {
char temp[10];
sprintf(temp, "%d.%d, ", i, j);
result += std::string(temp);
}
}
}
else {
for (int i = 0; i < 100; i++) {
char temp[10];
sprintf(temp, "%d, ", i);
result += std::string(temp);
}
}
return result;
}
void create_vload_kernels(std::vector<std::string>& KERNEL_FUNCTIONS, cl_device_id deviceID) {
const std::string KERNEL_FUNCTION_TEMPLATE_GLOBAL =
NL
NL "%s"
NL "__global %s data[] = { %s };"
NL
NL "__kernel void testKernel(__global uint *results) {"
NL " uint tid = get_global_id(0);"
NL
NL " // Testing: %s"
NL " const %s * ptr = data;"
NL " %s%s result_generic = vload%s(2, ptr);"
NL " %s%s result_reference = vload%s(2, data);"
NL
NL " results[tid] = all(result_generic == result_reference);"
NL "}"
NL;
const std::string KERNEL_FUNCTION_TEMPLATE_LOCAL =
NL
NL "%s"
NL "__constant %s to_copy_from[] = { %s };"
NL
NL "__kernel void testKernel(__global uint *results) {"
NL " uint tid = get_global_id(0);"
NL
NL " __local %s data[100];"
NL " for (int i = 0; i < sizeof(to_copy_from) / sizeof(to_copy_from[0]); i++)"
NL " data[i] = to_copy_from[i];"
NL
NL " const %s * ptr = data;"
NL " %s%s result_generic = vload%s(2, ptr);"
NL " %s%s result_reference = vload%s(2, data);"
NL
NL " results[tid] = all(result_generic == result_reference);"
NL "}"
NL;
const std::string KERNEL_FUNCTION_TEMPLATE_PRIVATE =
NL
NL "%s"
NL "__kernel void testKernel(__global uint *results) {"
NL " uint tid = get_global_id(0);"
NL
NL " %s data[] = { %s };"
NL " // Testing: %s"
NL " const %s * ptr = data;"
NL " %s%s result_generic = vload%s(2, ptr);"
NL " %s%s result_reference = vload%s(2, data);"
NL
NL " results[tid] = all(result_generic == result_reference);"
NL "}"
NL;
const std::string vector_sizes[] = { "2", "3", "4", "8", "16" };
const std::string gentype_base[] = { "double", "float", "char", "uchar", "short", "ushort", "int", "uint", "long", "ulong" };
const std::string kernel_variants[] = { KERNEL_FUNCTION_TEMPLATE_GLOBAL, KERNEL_FUNCTION_TEMPLATE_LOCAL, KERNEL_FUNCTION_TEMPLATE_PRIVATE };
const std::string cl_khr_fp64_pragma = "#pragma OPENCL EXTENSION cl_khr_fp64 : enable";
for (size_t i = 0; i < sizeof(gentype_base) / sizeof(gentype_base[0]); i++) {
const char *pragma_str = "";
if (i == 0) {
if (!is_extension_available(deviceID, "cl_khr_fp64"))
continue;
else
pragma_str = cl_khr_fp64_pragma.c_str();
}
for (size_t j = 0; j < sizeof(vector_sizes) / sizeof(vector_sizes[0]); j++) {
for (size_t k = 0; k < sizeof(kernel_variants) / sizeof(kernel_variants[0]); k++) {
char temp_kernel[4098];
sprintf(temp_kernel, kernel_variants[k].c_str(),
pragma_str,
gentype_base[i].c_str(),
get_default_data_for_type(gentype_base[i]).c_str(),
gentype_base[i].c_str(),
gentype_base[i].c_str(),
gentype_base[i].c_str(), vector_sizes[j].c_str(), vector_sizes[j].c_str(),
gentype_base[i].c_str(), vector_sizes[j].c_str(), vector_sizes[j].c_str()
);
KERNEL_FUNCTIONS.push_back(temp_kernel);
}
}
}
}
void create_vstore_kernels(std::vector<std::string>& KERNEL_FUNCTIONS, cl_device_id deviceID) {
const std::string KERNEL_FUNCTION_TEMPLATE_GLOBAL =
NL
NL "%s"
NL "__global %s data_generic[] = { %s };"
NL "__global %s data_reference[] = { %s };"
NL
NL "__kernel void testKernel(__global uint *results) {"
NL " uint tid = get_global_id(0);"
NL
NL " %s%s input = (%s%s)(1);"
NL " %s * ptr = data_generic;"
NL
NL " vstore%s(input, 2, ptr);"
NL " vstore%s(input, 2, data_reference);"
NL
NL " bool result = true;"
NL " for (int i = 0; i < sizeof(data_generic) / sizeof(data_generic[0]); i++)"
NL " if (data_generic[i] != data_reference[i])"
NL " result = false;"
NL
NL " results[tid] = result;"
NL "}"
NL;
const std::string KERNEL_FUNCTION_TEMPLATE_LOCAL =
NL
NL "%s"
NL "__constant %s to_copy_from[] = { %s };"
NL
NL "__kernel void testKernel(__global uint *results) {"
NL " uint tid = get_global_id(0);"
NL
NL " __local %s data_generic[100];"
NL " for (int i = 0; i < sizeof(to_copy_from) / sizeof(to_copy_from[0]); i++)"
NL " data_generic[i] = to_copy_from[i];"
NL
NL " __local %s data_reference[100];"
NL " for (int i = 0; i < sizeof(to_copy_from) / sizeof(to_copy_from[0]); i++)"
NL " data_reference[i] = to_copy_from[i];"
NL
NL " %s%s input = (%s%s)(1);"
NL " %s * ptr = data_generic;"
NL
NL " vstore%s(input, 2, ptr);"
NL " vstore%s(input, 2, data_reference);"
NL
NL " work_group_barrier(CLK_LOCAL_MEM_FENCE);"
NL
NL " bool result = true;"
NL " for (int i = 0; i < sizeof(data_generic) / sizeof(data_generic[0]); i++)"
NL " if (data_generic[i] != data_reference[i])"
NL " result = false;"
NL
NL " results[tid] = result;"
NL "}"
NL;
const std::string KERNEL_FUNCTION_TEMPLATE_PRIVATE =
NL
NL "%s"
NL "__kernel void testKernel(__global uint *results) {"
NL " uint tid = get_global_id(0);"
NL
NL " __private %s data_generic[] = { %s };"
NL " __private %s data_reference[] = { %s };"
NL
NL " %s%s input = (%s%s)(1);"
NL " %s * ptr = data_generic;"
NL
NL " vstore%s(input, 2, ptr);"
NL " vstore%s(input, 2, data_reference);"
NL
NL " bool result = true;"
NL " for (int i = 0; i < sizeof(data_generic) / sizeof(data_generic[0]); i++)"
NL " if (data_generic[i] != data_reference[i])"
NL " result = false;"
NL
NL " results[tid] = result;"
NL "}"
NL;
const std::string vector_sizes[] = { "2", "3", "4", "8", "16" };
const std::string gentype_base[] = { "double", "float", "char", "uchar", "short", "ushort", "int", "uint", "long", "ulong" };
const std::string kernel_variants[] = { KERNEL_FUNCTION_TEMPLATE_GLOBAL, KERNEL_FUNCTION_TEMPLATE_LOCAL, KERNEL_FUNCTION_TEMPLATE_PRIVATE };
const std::string cl_khr_fp64_pragma = "#pragma OPENCL EXTENSION cl_khr_fp64 : enable";
for (size_t i = 0; i < sizeof(gentype_base) / sizeof(gentype_base[0]); i++) {
const char *pragma_str = "";
if (i == 0) {
if (!is_extension_available(deviceID, "cl_khr_fp64"))
continue;
else
pragma_str = cl_khr_fp64_pragma.c_str();
}
for (size_t j = 0; j < sizeof(vector_sizes) / sizeof(vector_sizes[0]); j++) {
for (size_t k = 0; k < sizeof(kernel_variants) / sizeof(kernel_variants[0]); k++) {
char temp_kernel[4098];
switch (k) {
case 0: // global template
case 2: // private template
sprintf(temp_kernel, kernel_variants[k].c_str(),
pragma_str,
gentype_base[i].c_str(), get_default_data_for_type(gentype_base[i]).c_str(),
gentype_base[i].c_str(), get_default_data_for_type(gentype_base[i]).c_str(),
gentype_base[i].c_str(), vector_sizes[j].c_str(), gentype_base[i].c_str(), vector_sizes[j].c_str(),
gentype_base[i].c_str(),
vector_sizes[j].c_str(),
vector_sizes[j].c_str()
);
break;
case 1: // local template
sprintf(temp_kernel, kernel_variants[k].c_str(),
pragma_str,
gentype_base[i].c_str(), get_default_data_for_type(gentype_base[i]).c_str(),
gentype_base[i].c_str(),
gentype_base[i].c_str(),
gentype_base[i].c_str(), vector_sizes[j].c_str(), gentype_base[i].c_str(), vector_sizes[j].c_str(),
gentype_base[i].c_str(),
vector_sizes[j].c_str(),
vector_sizes[j].c_str()
);
break;
}
KERNEL_FUNCTIONS.push_back(temp_kernel);
}
}
}
}
int test_builtin_functions(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
std::vector<std::string> KERNEL_FUNCTIONS;
create_math_kernels(KERNEL_FUNCTIONS);
create_vload_kernels(KERNEL_FUNCTIONS, deviceID);
create_vstore_kernels(KERNEL_FUNCTIONS, deviceID);
CAdvancedTest test(KERNEL_FUNCTIONS);
return test.Execute(deviceID, context, queue, num_elements);
}
int test_generic_advanced_casting(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
std::vector<std::string> KERNEL_FUNCTIONS;
KERNEL_FUNCTIONS.push_back(
NL
NL "__global char arr[16] = { 0, 0, 0, 0, 1, 1, 1, 1, 2, 2, 2, 2, 3, 3, 3, 3 };"
NL
NL "__kernel void testKernel(__global uint *results) {"
NL " uint tid = get_global_id(0);"
NL
NL " const int * volatile ptr = (const int *)arr;"
NL
NL " results[tid] = (ptr[0] == 0x00000000) && (ptr[1] == 0x01010101) && (ptr[2] == 0x02020202) && (ptr[3] == 0x03030303);"
NL "}"
NL
);
KERNEL_FUNCTIONS.push_back(
NL
NL "__kernel void testKernel(__global uint *results) {"
NL " uint tid = get_global_id(0);"
NL
NL " __local int i;"
NL " i = 0x11112222;"
NL " short *ptr = (short *)&i;"
NL " local int *lptr = (local int *)ptr;"
NL
NL " results[tid] = (lptr == &i) && (*lptr == i);"
NL "}"
NL
);
KERNEL_FUNCTIONS.push_back(
NL
NL "__kernel void testKernel(__global uint *results) {"
NL " uint tid = get_global_id(0);"
NL
NL " int i = 0x11112222;"
NL
NL " void *ptr = &i;"
NL " int copy = *((int *)ptr);"
NL
NL " results[tid] = (copy == i);"
NL "}"
NL
);
CAdvancedTest test(KERNEL_FUNCTIONS);
return test.Execute(deviceID, context, queue, num_elements);
}
int test_generic_ptr_to_host_mem_svm(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
cl_int result = CL_SUCCESS;
/* Test SVM capabilities and select matching tests */
cl_device_svm_capabilities caps;
auto version = get_device_cl_version(deviceID);
auto expected_min_version = Version(2, 0);
cl_int error = clGetDeviceInfo(deviceID, CL_DEVICE_SVM_CAPABILITIES, sizeof(caps), &caps, NULL);
test_error(error, "clGetDeviceInfo(CL_DEVICE_SVM_CAPABILITIES) failed");
if ((version < expected_min_version)
|| (version >= Version(3, 0) && caps == 0))
return TEST_SKIPPED_ITSELF;
if (caps & CL_DEVICE_SVM_COARSE_GRAIN_BUFFER) {
CAdvancedTest test_global_svm_ptr(common::GLOBAL_KERNEL_FUNCTION, ARG_TYPE_COARSE_GRAINED_SVM);
result |= test_global_svm_ptr.Execute(deviceID, context, queue, num_elements);
}
if (caps & CL_DEVICE_SVM_FINE_GRAIN_BUFFER) {
CAdvancedTest test_global_svm_ptr(common::GLOBAL_KERNEL_FUNCTION, ARG_TYPE_FINE_GRAINED_BUFFER_SVM);
result |= test_global_svm_ptr.Execute(deviceID, context, queue, num_elements);
}
if (caps & CL_DEVICE_SVM_FINE_GRAIN_SYSTEM) {
CAdvancedTest test_global_svm_ptr(common::GLOBAL_KERNEL_FUNCTION, ARG_TYPE_FINE_GRAINED_SYSTEM_SVM);
result |= test_global_svm_ptr.Execute(deviceID, context, queue, num_elements);
}
if (caps & CL_DEVICE_SVM_ATOMICS) {
CAdvancedTest test_global_svm_ptr(common::GLOBAL_KERNEL_FUNCTION, ARG_TYPE_ATOMICS_SVM);
result |= test_global_svm_ptr.Execute(deviceID, context, queue, num_elements);
}
return result;
}
int test_generic_ptr_to_host_mem(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
cl_int result = CL_SUCCESS;
CAdvancedTest test_global_ptr(common::GLOBAL_KERNEL_FUNCTION, ARG_TYPE_HOST_PTR);
result |= test_global_ptr.Execute(deviceID, context, queue, num_elements);
CAdvancedTest test_local_ptr(common::LOCAL_KERNEL_FUNCTION, ARG_TYPE_HOST_LOCAL);
result |= test_local_ptr.Execute(deviceID, context, queue, num_elements / 64);
return result;
}
Reference in New Issue View Git Blame Copy Permalink