You can not select more than 25 topics
Topics must start with a letter or number, can include dashes ('-') and can be up to 35 characters long.
405 lines
14 KiB
405 lines
14 KiB
//
|
|
// 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 "harness/conversions.h"
|
|
#include "harness/typeWrappers.h"
|
|
|
|
#if !defined(__APPLE__)
|
|
#include <CL/cl_gl.h>
|
|
#endif
|
|
|
|
static const char *bufferKernelPattern =
|
|
"__kernel void sample_test( __global %s%s *source, __global %s%s *clDest, "
|
|
"__global %s%s *glDest )\n"
|
|
"{\n"
|
|
" int tid = get_global_id(0);\n"
|
|
" clDest[ tid ] = source[ tid ] + (%s%s)(1);\n"
|
|
" glDest[ tid ] = source[ tid ] + (%s%s)(2);\n"
|
|
"}\n";
|
|
|
|
#define TYPE_CASE(enum, type, range, offset) \
|
|
case enum: { \
|
|
cl_##type *ptr = (cl_##type *)outData; \
|
|
for (i = 0; i < count; i++) \
|
|
ptr[i] = (cl_##type)((genrand_int32(d) & range) - offset); \
|
|
break; \
|
|
}
|
|
|
|
void gen_input_data(ExplicitType type, size_t count, MTdata d, void *outData)
|
|
{
|
|
size_t i;
|
|
|
|
switch (type)
|
|
{
|
|
case kBool: {
|
|
bool *boolPtr = (bool *)outData;
|
|
for (i = 0; i < count; i++)
|
|
{
|
|
boolPtr[i] = (genrand_int32(d) & 1) ? true : false;
|
|
}
|
|
break;
|
|
}
|
|
|
|
TYPE_CASE(kChar, char, 250, 127)
|
|
TYPE_CASE(kUChar, uchar, 250, 0)
|
|
TYPE_CASE(kShort, short, 65530, 32767)
|
|
TYPE_CASE(kUShort, ushort, 65530, 0)
|
|
TYPE_CASE(kInt, int, 0x0fffffff, 0x70000000)
|
|
TYPE_CASE(kUInt, uint, 0x0fffffff, 0)
|
|
|
|
case kLong: {
|
|
cl_long *longPtr = (cl_long *)outData;
|
|
for (i = 0; i < count; i++)
|
|
{
|
|
longPtr[i] = (cl_long)genrand_int32(d)
|
|
| ((cl_ulong)genrand_int32(d) << 32);
|
|
}
|
|
break;
|
|
}
|
|
|
|
case kULong: {
|
|
cl_ulong *ulongPtr = (cl_ulong *)outData;
|
|
for (i = 0; i < count; i++)
|
|
{
|
|
ulongPtr[i] = (cl_ulong)genrand_int32(d)
|
|
| ((cl_ulong)genrand_int32(d) << 32);
|
|
}
|
|
break;
|
|
}
|
|
|
|
case kFloat: {
|
|
cl_float *floatPtr = (float *)outData;
|
|
for (i = 0; i < count; i++)
|
|
floatPtr[i] = get_random_float(-100000.f, 100000.f, d);
|
|
break;
|
|
}
|
|
|
|
default:
|
|
log_error(
|
|
"ERROR: Invalid type passed in to generate_random_data!\n");
|
|
break;
|
|
}
|
|
}
|
|
|
|
#define INC_CASE(enum, type) \
|
|
case enum: { \
|
|
cl_##type *src = (cl_##type *)inData; \
|
|
cl_##type *dst = (cl_##type *)outData; \
|
|
*dst = *src + 1; \
|
|
break; \
|
|
}
|
|
|
|
void get_incremented_value(void *inData, void *outData, ExplicitType type)
|
|
{
|
|
switch (type)
|
|
{
|
|
INC_CASE(kChar, char)
|
|
INC_CASE(kUChar, uchar)
|
|
INC_CASE(kShort, short)
|
|
INC_CASE(kUShort, ushort)
|
|
INC_CASE(kInt, int)
|
|
INC_CASE(kUInt, uint)
|
|
INC_CASE(kLong, long)
|
|
INC_CASE(kULong, ulong)
|
|
INC_CASE(kFloat, float)
|
|
default: break;
|
|
}
|
|
}
|
|
|
|
int test_buffer_kernel(cl_context context, cl_command_queue queue,
|
|
ExplicitType vecType, size_t vecSize, int numElements,
|
|
int validate_only, MTdata d)
|
|
{
|
|
clProgramWrapper program;
|
|
clKernelWrapper kernel;
|
|
clMemWrapper streams[3];
|
|
size_t dataSize = numElements * 16 * sizeof(cl_long);
|
|
#if !(defined(_WIN32) && defined(_MSC_VER))
|
|
cl_long inData[numElements * 16], outDataCL[numElements * 16],
|
|
outDataGL[numElements * 16];
|
|
#else
|
|
cl_long *inData = (cl_long *)_malloca(dataSize);
|
|
cl_long *outDataCL = (cl_long *)_malloca(dataSize);
|
|
cl_long *outDataGL = (cl_long *)_malloca(dataSize);
|
|
#endif
|
|
glBufferWrapper inGLBuffer, outGLBuffer;
|
|
int i;
|
|
size_t bufferSize;
|
|
|
|
int error;
|
|
size_t threads[1], localThreads[1];
|
|
char kernelSource[10240];
|
|
char *programPtr;
|
|
char sizeName[4];
|
|
|
|
/* Create the source */
|
|
if (vecSize == 1)
|
|
sizeName[0] = 0;
|
|
else
|
|
sprintf(sizeName, "%d", (int)vecSize);
|
|
|
|
sprintf(kernelSource, bufferKernelPattern, get_explicit_type_name(vecType),
|
|
sizeName, get_explicit_type_name(vecType), sizeName,
|
|
get_explicit_type_name(vecType), sizeName,
|
|
get_explicit_type_name(vecType), sizeName,
|
|
get_explicit_type_name(vecType), sizeName);
|
|
|
|
/* Create kernels */
|
|
programPtr = kernelSource;
|
|
if (create_single_kernel_helper(context, &program, &kernel, 1,
|
|
(const char **)&programPtr, "sample_test"))
|
|
{
|
|
return -1;
|
|
}
|
|
|
|
bufferSize = numElements * vecSize * get_explicit_type_size(vecType);
|
|
|
|
/* Generate some almost-random input data */
|
|
gen_input_data(vecType, vecSize * numElements, d, inData);
|
|
memset(outDataCL, 0, dataSize);
|
|
memset(outDataGL, 0, dataSize);
|
|
|
|
/* Generate some GL buffers to go against */
|
|
glGenBuffers(1, &inGLBuffer);
|
|
glGenBuffers(1, &outGLBuffer);
|
|
|
|
glBindBuffer(GL_ARRAY_BUFFER, inGLBuffer);
|
|
glBufferData(GL_ARRAY_BUFFER, bufferSize, inData, GL_STATIC_DRAW);
|
|
|
|
// Note: we need to bind the output buffer, even though we don't care about
|
|
// its values yet, because CL needs it to get the buffer size
|
|
glBindBuffer(GL_ARRAY_BUFFER, outGLBuffer);
|
|
glBufferData(GL_ARRAY_BUFFER, bufferSize, outDataGL, GL_STATIC_DRAW);
|
|
|
|
glBindBuffer(GL_ARRAY_BUFFER, 0);
|
|
glFinish();
|
|
|
|
|
|
/* Generate some streams. The first and last ones are GL, middle one just
|
|
* vanilla CL */
|
|
streams[0] = (*clCreateFromGLBuffer_ptr)(context, CL_MEM_READ_ONLY,
|
|
inGLBuffer, &error);
|
|
test_error(error, "Unable to create input GL buffer");
|
|
|
|
streams[1] =
|
|
clCreateBuffer(context, CL_MEM_READ_WRITE, bufferSize, NULL, &error);
|
|
test_error(error, "Unable to create output CL buffer");
|
|
|
|
streams[2] = (*clCreateFromGLBuffer_ptr)(context, CL_MEM_WRITE_ONLY,
|
|
outGLBuffer, &error);
|
|
test_error(error, "Unable to create output GL buffer");
|
|
|
|
|
|
/* Validate the info */
|
|
if (validate_only)
|
|
{
|
|
int result = (CheckGLObjectInfo(streams[0], CL_GL_OBJECT_BUFFER,
|
|
(GLuint)inGLBuffer, (GLenum)0, 0)
|
|
| CheckGLObjectInfo(streams[2], CL_GL_OBJECT_BUFFER,
|
|
(GLuint)outGLBuffer, (GLenum)0, 0));
|
|
for (i = 0; i < 3; i++)
|
|
{
|
|
streams[i].reset();
|
|
}
|
|
|
|
glDeleteBuffers(1, &inGLBuffer);
|
|
inGLBuffer = 0;
|
|
glDeleteBuffers(1, &outGLBuffer);
|
|
outGLBuffer = 0;
|
|
|
|
return result;
|
|
}
|
|
|
|
/* Assign streams and execute */
|
|
for (int i = 0; i < 3; i++)
|
|
{
|
|
error = clSetKernelArg(kernel, i, sizeof(streams[i]), &streams[i]);
|
|
test_error(error, "Unable to set kernel arguments");
|
|
}
|
|
error =
|
|
(*clEnqueueAcquireGLObjects_ptr)(queue, 1, &streams[0], 0, NULL, NULL);
|
|
test_error(error, "Unable to acquire GL obejcts");
|
|
error =
|
|
(*clEnqueueAcquireGLObjects_ptr)(queue, 1, &streams[2], 0, NULL, NULL);
|
|
test_error(error, "Unable to acquire GL obejcts");
|
|
|
|
/* Run the kernel */
|
|
threads[0] = numElements;
|
|
|
|
error = get_max_common_work_group_size(context, kernel, threads[0],
|
|
&localThreads[0]);
|
|
test_error(error, "Unable to get work group size to use");
|
|
|
|
error = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, threads,
|
|
localThreads, 0, NULL, NULL);
|
|
test_error(error, "Unable to execute test kernel");
|
|
|
|
error =
|
|
(*clEnqueueReleaseGLObjects_ptr)(queue, 1, &streams[0], 0, NULL, NULL);
|
|
test_error(error, "clEnqueueReleaseGLObjects failed");
|
|
error =
|
|
(*clEnqueueReleaseGLObjects_ptr)(queue, 1, &streams[2], 0, NULL, NULL);
|
|
test_error(error, "clEnqueueReleaseGLObjects failed");
|
|
|
|
// Get the results from both CL and GL and make sure everything looks
|
|
// correct
|
|
error = clEnqueueReadBuffer(queue, streams[1], CL_TRUE, 0, bufferSize,
|
|
outDataCL, 0, NULL, NULL);
|
|
test_error(error, "Unable to read output CL array!");
|
|
|
|
glBindBuffer(GL_ARRAY_BUFFER, outGLBuffer);
|
|
void *glMem = glMapBuffer(GL_ARRAY_BUFFER, GL_READ_ONLY);
|
|
memcpy(outDataGL, glMem, bufferSize);
|
|
glUnmapBuffer(GL_ARRAY_BUFFER);
|
|
|
|
char *inP = (char *)inData, *glP = (char *)outDataGL,
|
|
*clP = (char *)outDataCL;
|
|
error = 0;
|
|
for (size_t i = 0; i < numElements * vecSize; i++)
|
|
{
|
|
cl_long expectedCLValue, expectedGLValue;
|
|
get_incremented_value(inP, &expectedCLValue, vecType);
|
|
get_incremented_value(&expectedCLValue, &expectedGLValue, vecType);
|
|
|
|
if (memcmp(clP, &expectedCLValue, get_explicit_type_size(vecType)) != 0)
|
|
{
|
|
char scratch[64];
|
|
log_error(
|
|
"ERROR: Data sample %d from the CL output did not validate!\n",
|
|
(int)i);
|
|
log_error("\t Input: %s\n",
|
|
GetDataVectorString(inP, get_explicit_type_size(vecType),
|
|
1, scratch));
|
|
log_error("\tExpected: %s\n",
|
|
GetDataVectorString(&expectedCLValue,
|
|
get_explicit_type_size(vecType), 1,
|
|
scratch));
|
|
log_error("\t Actual: %s\n",
|
|
GetDataVectorString(clP, get_explicit_type_size(vecType),
|
|
1, scratch));
|
|
error = -1;
|
|
}
|
|
|
|
if (memcmp(glP, &expectedGLValue, get_explicit_type_size(vecType)) != 0)
|
|
{
|
|
char scratch[64];
|
|
log_error(
|
|
"ERROR: Data sample %d from the GL output did not validate!\n",
|
|
(int)i);
|
|
log_error("\t Input: %s\n",
|
|
GetDataVectorString(inP, get_explicit_type_size(vecType),
|
|
1, scratch));
|
|
log_error("\tExpected: %s\n",
|
|
GetDataVectorString(&expectedGLValue,
|
|
get_explicit_type_size(vecType), 1,
|
|
scratch));
|
|
log_error("\t Actual: %s\n",
|
|
GetDataVectorString(glP, get_explicit_type_size(vecType),
|
|
1, scratch));
|
|
error = -1;
|
|
}
|
|
|
|
if (error) return error;
|
|
|
|
inP += get_explicit_type_size(vecType);
|
|
glP += get_explicit_type_size(vecType);
|
|
clP += get_explicit_type_size(vecType);
|
|
}
|
|
|
|
for (i = 0; i < 3; i++)
|
|
{
|
|
streams[i].reset();
|
|
}
|
|
|
|
glDeleteBuffers(1, &inGLBuffer);
|
|
inGLBuffer = 0;
|
|
glDeleteBuffers(1, &outGLBuffer);
|
|
outGLBuffer = 0;
|
|
|
|
return 0;
|
|
}
|
|
|
|
int test_buffers(cl_device_id device, cl_context context,
|
|
cl_command_queue queue, int numElements)
|
|
{
|
|
ExplicitType vecType[] = {
|
|
kChar, kUChar, kShort, kUShort, kInt,
|
|
kUInt, kLong, kULong, kFloat, kNumExplicitTypes
|
|
};
|
|
unsigned int vecSizes[] = { 1, 2, 4, 8, 16, 0 };
|
|
unsigned int index, typeIndex;
|
|
int retVal = 0;
|
|
RandomSeed seed(gRandomSeed);
|
|
|
|
|
|
for (typeIndex = 0; vecType[typeIndex] != kNumExplicitTypes; typeIndex++)
|
|
{
|
|
for (index = 0; vecSizes[index] != 0; index++)
|
|
{
|
|
// Test!
|
|
if (test_buffer_kernel(context, queue, vecType[typeIndex],
|
|
vecSizes[index], numElements, 0, seed)
|
|
!= 0)
|
|
{
|
|
char sizeNames[][4] = { "", "", "2", "", "4", "", "", "", "8",
|
|
"", "", "", "", "", "", "", "16" };
|
|
log_error(" Buffer test %s%s FAILED\n",
|
|
get_explicit_type_name(vecType[typeIndex]),
|
|
sizeNames[vecSizes[index]]);
|
|
retVal++;
|
|
}
|
|
}
|
|
}
|
|
|
|
return retVal;
|
|
}
|
|
|
|
|
|
int test_buffers_getinfo(cl_device_id device, cl_context context,
|
|
cl_command_queue queue, int numElements)
|
|
{
|
|
ExplicitType vecType[] = {
|
|
kChar, kUChar, kShort, kUShort, kInt,
|
|
kUInt, kLong, kULong, kFloat, kNumExplicitTypes
|
|
};
|
|
unsigned int vecSizes[] = { 1, 2, 4, 8, 16, 0 };
|
|
unsigned int index, typeIndex;
|
|
int retVal = 0;
|
|
RandomSeed seed(gRandomSeed);
|
|
|
|
|
|
for (typeIndex = 0; vecType[typeIndex] != kNumExplicitTypes; typeIndex++)
|
|
{
|
|
for (index = 0; vecSizes[index] != 0; index++)
|
|
{
|
|
// Test!
|
|
if (test_buffer_kernel(context, queue, vecType[typeIndex],
|
|
vecSizes[index], numElements, 1, seed)
|
|
!= 0)
|
|
{
|
|
char sizeNames[][4] = { "", "", "2", "", "4", "", "", "", "8",
|
|
"", "", "", "", "", "", "", "16" };
|
|
log_error(" Buffer test %s%s FAILED\n",
|
|
get_explicit_type_name(vecType[typeIndex]),
|
|
sizeNames[vecSizes[index]]);
|
|
retVal++;
|
|
}
|
|
}
|
|
}
|
|
|
|
return retVal;
|
|
}
|