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

Make extended_async_copy tests type agnostic (#1619)

Browse Source 
The latest version of the cl_khr_extended_async_copies extension uses
element size rather the element type as its base. The means it can be
called with arbitrary and in particular non power of 2 sizes, such as 3
or 13.

Update the test_async_copy2D and test_async_copy3D tests to make them
element size based rather than type based.

As well as this run all tests that can fit into the memory of the
target rather than presumed large elements cannot fit.

Make some addtional good practice changes in terms of const usage,
declaring variables where they are use, and usage of iterators.

The test coverage increases from 1224 cases to 1332 cases for the
test_async_copy2D and test_async_copy3D cases.

Ticket: #1579

Signed-off-by: Chris Gearing <chris.gearing@mobileye.com>
Co-authored-by: Chris Gearing <chris.gearing@mobileye.com>
This commit is contained in:
Chris Gearing
2023-05-30 18:05:18 +02:00
committed by GitHub
parent 38abfc7d24
commit 19f4fc3f3d
2 changed files with 171 additions and 215 deletions
Download Patch File Download Diff File Expand all files Collapse all files

178
test_conformance/basic/test_async_copy2D.cpp
View File

@@ -27,17 +27,25 @@
static const char *async_global_to_local_kernel2D = R"OpenCLC(
#pragma OPENCL EXTENSION cl_khr_extended_async_copies : enable
%s // optional pragma string
__kernel void test_fn(const __global %s *src, __global %s *dst,
__local %s *localBuffer, int numElementsPerLine,
#define STRUCT_SIZE %d
typedef struct __attribute__((packed))
{
uchar byte[STRUCT_SIZE];
} VarSizeStruct __attribute__((aligned(1)));
__kernel void test_fn(const __global VarSizeStruct *src, __global VarSizeStruct *dst,
__local VarSizeStruct *localBuffer, int numElementsPerLine,
int lineCopiesPerWorkgroup, int lineCopiesPerWorkItem,
int srcStride, int dstStride) {
// Zero the local storage first
for (int i = 0; i < lineCopiesPerWorkItem; i++) {
for (int j = 0; j < numElementsPerLine; j++) {
const int index = (get_local_id(0) * lineCopiesPerWorkItem + i) * dstStride + j;
localBuffer[index] = (%s)(%s)0;
for (int k = 0; k < STRUCT_SIZE; k++) {
localBuffer[index].byte[k] = 0;
}
}
}
@@ -45,7 +53,7 @@ __kernel void test_fn(const __global %s *src, __global %s *dst,
// try the copy
barrier( CLK_LOCAL_MEM_FENCE );
event_t event = async_work_group_copy_2D2D(localBuffer, 0, src,
lineCopiesPerWorkgroup * get_group_id(0) * srcStride, sizeof(%s),
lineCopiesPerWorkgroup * get_group_id(0) * srcStride, sizeof(VarSizeStruct),
(size_t)numElementsPerLine, (size_t)lineCopiesPerWorkgroup, srcStride, dstStride, 0);
// Wait for the copy to complete, then verify by manually copying to the dest
@@ -63,16 +71,24 @@ __kernel void test_fn(const __global %s *src, __global %s *dst,
static const char *async_local_to_global_kernel2D = R"OpenCLC(
#pragma OPENCL EXTENSION cl_khr_extended_async_copies : enable
%s // optional pragma string
__kernel void test_fn(const __global %s *src, __global %s *dst, __local %s *localBuffer,
#define STRUCT_SIZE %d
typedef struct __attribute__((packed))
{
uchar byte[STRUCT_SIZE];
} VarSizeStruct __attribute__((aligned(1)));
__kernel void test_fn(const __global VarSizeStruct *src, __global VarSizeStruct *dst, __local VarSizeStruct *localBuffer,
int numElementsPerLine, int lineCopiesPerWorkgroup,
int lineCopiesPerWorkItem, int srcStride, int dstStride) {
// Zero the local storage first
for (int i = 0; i < lineCopiesPerWorkItem; i++) {
for (int j = 0; j < numElementsPerLine; j++) {
const int index = (get_local_id(0) * lineCopiesPerWorkItem + i) * srcStride + j;
localBuffer[index] = (%s)(%s)0;
for (int k = 0; k < STRUCT_SIZE; k++) {
localBuffer[index].byte[k] = 0;
}
}
}
@@ -90,36 +106,22 @@ __kernel void test_fn(const __global %s *src, __global %s *dst, __local %s *loca
// Do this to verify all kernels are done copying to the local buffer before we try the copy
barrier(CLK_LOCAL_MEM_FENCE);
event_t event = async_work_group_copy_2D2D(dst, lineCopiesPerWorkgroup * get_group_id(0) * dstStride,
localBuffer, 0, sizeof(%s), (size_t)numElementsPerLine, (size_t)lineCopiesPerWorkgroup, srcStride,
localBuffer, 0, sizeof(VarSizeStruct), (size_t)numElementsPerLine, (size_t)lineCopiesPerWorkgroup, srcStride,
dstStride, 0 );
wait_group_events(1, &event);
};
)OpenCLC";
int test_copy2D(cl_device_id deviceID, cl_context context,
cl_command_queue queue, const char *kernelCode,
ExplicitType vecType, int vecSize, int srcMargin, int dstMargin,
bool localIsDst)
int test_copy2D(const cl_device_id deviceID, const cl_context context,
const cl_command_queue queue, const char *const kernelCode,
const size_t elementSize, const int srcMargin,
const int dstMargin, const bool localIsDst)
{
int error;
clProgramWrapper program;
clKernelWrapper kernel;
clMemWrapper streams[2];
size_t threads[1], localThreads[1];
void *inBuffer, *outBuffer, *outBufferCopy;
MTdata d;
char vecNameString[64];
vecNameString[0] = 0;
if (vecSize == 1)
sprintf(vecNameString, "%s", get_explicit_type_name(vecType));
else
sprintf(vecNameString, "%s%d", get_explicit_type_name(vecType),
vecSize);
size_t elementSize = get_explicit_type_size(vecType) * vecSize;
log_info("Testing %s with srcMargin = %d, dstMargin = %d\n", vecNameString,
srcMargin, dstMargin);
log_info("Testing %d byte element with srcMargin = %d, dstMargin = %d\n",
elementSize, srcMargin, dstMargin);
cl_long max_local_mem_size;
error =
@@ -139,6 +141,13 @@ int test_copy2D(cl_device_id deviceID, cl_context context,
test_error(error,
"clGetDeviceInfo for CL_DEVICE_MAX_MEM_ALLOC_SIZE failed.");
cl_long max_work_group_size;
error = clGetDeviceInfo(deviceID, CL_DEVICE_MAX_WORK_GROUP_SIZE,
sizeof(max_work_group_size), &max_work_group_size,
NULL);
test_error(error,
"clGetDeviceInfo for CL_DEVICE_MAX_WORK_GROUP_SIZE failed.");
if (max_alloc_size > max_global_mem_size / 2)
max_alloc_size = max_global_mem_size / 2;
@@ -149,20 +158,17 @@ int test_copy2D(cl_device_id deviceID, cl_context context,
test_error(error,
"clGetDeviceInfo for CL_DEVICE_MAX_COMPUTE_UNITS failed.");
char programSource[4096];
programSource[0] = 0;
char *programPtr;
char programSource[4096] = { 0 };
const char *programPtr = programSource;
sprintf(programSource, kernelCode,
vecType == kDouble ? "#pragma OPENCL EXTENSION cl_khr_fp64 : enable"
: "",
vecNameString, vecNameString, vecNameString, vecNameString,
get_explicit_type_name(vecType), vecNameString);
sprintf(programSource, kernelCode, elementSize);
// log_info("program: %s\n", programSource);
programPtr = programSource;
clProgramWrapper program;
clKernelWrapper kernel;
error = create_single_kernel_helper(context, &program, &kernel, 1,
(const char **)&programPtr, "test_fn");
&programPtr, "test_fn");
test_error(error, "Unable to create testing kernel");
size_t max_workgroup_size;
@@ -188,9 +194,6 @@ int test_copy2D(cl_device_id deviceID, cl_context context,
const cl_int dstStride = numElementsPerLine + dstMargin;
const cl_int srcStride = numElementsPerLine + srcMargin;
elementSize =
get_explicit_type_size(vecType) * ((vecSize == 3) ? 4 : vecSize);
const size_t lineCopiesPerWorkItem = 13;
const size_t localStorageSpacePerWorkitem = lineCopiesPerWorkItem
* elementSize * (localIsDst ? dstStride : srcStride);
@@ -208,7 +211,6 @@ int test_copy2D(cl_device_id deviceID, cl_context context,
if (maxLocalWorkgroupSize > max_workgroup_size)
localWorkgroupSize = max_workgroup_size;
const size_t maxTotalLinesIn =
(max_alloc_size / elementSize + srcMargin) / srcStride;
const size_t maxTotalLinesOut =
@@ -231,9 +233,17 @@ int test_copy2D(cl_device_id deviceID, cl_context context,
const size_t globalWorkgroupSize =
numberOfLocalWorkgroups * localWorkgroupSize;
inBuffer = (void *)malloc(inBufferSize);
outBuffer = (void *)malloc(outBufferSize);
outBufferCopy = (void *)malloc(outBufferSize);
if ((localBufferSize / 4) > max_work_group_size)
{
log_info("Skipping due to resource requirements local:%db "
"max_work_group_size:%d\n",
localBufferSize, max_work_group_size);
return 0;
}
void *const inBuffer = (void *)malloc(inBufferSize);
void *const outBuffer = (void *)malloc(outBufferSize);
void *const outBufferCopy = (void *)malloc(outBufferSize);
const cl_int lineCopiesPerWorkItemInt =
static_cast<cl_int>(lineCopiesPerWorkItem);
@@ -250,18 +260,20 @@ int test_copy2D(cl_device_id deviceID, cl_context context,
(int)inBufferSize, (int)outBufferSize, lineCopiesPerWorkgroup,
lineCopiesPerWorkItemInt);
size_t threads[1], localThreads[1];
threads[0] = globalWorkgroupSize;
localThreads[0] = localWorkgroupSize;
d = init_genrand(gRandomSeed);
generate_random_data(
vecType, inBufferSize / get_explicit_type_size(vecType), d, inBuffer);
generate_random_data(
vecType, outBufferSize / get_explicit_type_size(vecType), d, outBuffer);
MTdata d = init_genrand(gRandomSeed);
generate_random_data(kChar, inBufferSize, d, inBuffer);
generate_random_data(kChar, outBufferSize, d, outBuffer);
free_mtdata(d);
d = NULL;
memcpy(outBufferCopy, outBuffer, outBufferSize);
clMemWrapper streams[2];
streams[0] = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR, inBufferSize,
inBuffer, &error);
test_error(error, "Unable to create input buffer");
@@ -301,8 +313,7 @@ int test_copy2D(cl_device_id deviceID, cl_context context,
// Verify
int failuresPrinted = 0;
// Verify
size_t typeSize = get_explicit_type_size(vecType) * vecSize;
for (int i = 0;
i < (int)globalWorkgroupSize * lineCopiesPerWorkItem * elementSize;
i += elementSize)
@@ -313,13 +324,12 @@ int test_copy2D(cl_device_id deviceID, cl_context context,
int inIdx = i * srcStride + j;
int outIdx = i * dstStride + j;
if (memcmp(((char *)inBuffer) + inIdx, ((char *)outBuffer) + outIdx,
typeSize)
elementSize)
!= 0)
{
unsigned char *inchar = (unsigned char *)inBuffer + inIdx;
unsigned char *outchar = (unsigned char *)outBuffer + outIdx;
char values[4096];
values[0] = 0;
char values[4096] = { 0 };
if (failuresPrinted == 0)
{
@@ -382,16 +392,14 @@ int test_copy2D_all_types(cl_device_id deviceID, cl_context context,
cl_command_queue queue, const char *kernelCode,
bool localIsDst)
{
ExplicitType vecType[] = {
kChar, kUChar, kShort, kUShort, kInt, kUInt, kLong,
kULong, kFloat, kDouble, kNumExplicitTypes
};
const unsigned int elemSizes[] = { 1, 2, 3, 4, 5, 6, 7,
8, 13, 16, 32, 47, 64 };
// The margins below represent the number of elements between the end of
// one line and the start of the next. The strides are equivalent to the
// length of the line plus the chosen margin.
unsigned int vecSizes[] = { 1, 2, 3, 4, 8, 16, 0 };
unsigned int smallTypesMarginSizes[] = { 0, 10, 100 };
unsigned int size, typeIndex, srcMargin, dstMargin;
// These have to be multipliers, because the margin must be a multiple of
// element size.
const unsigned int marginMultipliers[] = { 0, 10, 100 };
int errors = 0;
@@ -399,55 +407,27 @@ int test_copy2D_all_types(cl_device_id deviceID, cl_context context,
{
log_info(
"Device does not support extended async copies. Skipping test.\n");
return 0;
}
for (typeIndex = 0; vecType[typeIndex] != kNumExplicitTypes; typeIndex++)
else
{
if (vecType[typeIndex] == kDouble
&& !is_extension_available(deviceID, "cl_khr_fp64"))
continue;
if ((vecType[typeIndex] == kLong || vecType[typeIndex] == kULong)
&& !gHasLong)
continue;
for (size = 0; vecSizes[size] != 0; size++)
for (const unsigned int elemSize : elemSizes)
{
if (get_explicit_type_size(vecType[typeIndex]) * vecSizes[size]
<= 2) // small type
for (const unsigned int srcMarginMultiplier : marginMultipliers)
{
for (srcMargin = 0; srcMargin < sizeof(smallTypesMarginSizes)
/ sizeof(smallTypesMarginSizes[0]);
srcMargin++)
{
for (dstMargin = 0;
dstMargin < sizeof(smallTypesMarginSizes)
/ sizeof(smallTypesMarginSizes[0]);
dstMargin++)
for (const unsigned int dstMarginMultiplier : marginMultipliers)
{
if (test_copy2D(deviceID, context, queue, kernelCode,
vecType[typeIndex], vecSizes[size],
smallTypesMarginSizes[srcMargin],
smallTypesMarginSizes[dstMargin],
localIsDst))
elemSize, srcMarginMultiplier * elemSize,
dstMarginMultiplier * elemSize, localIsDst))
{
errors++;
}
}
}
}
// not a small type, check only zero stride
else if (test_copy2D(deviceID, context, queue, kernelCode,
vecType[typeIndex], vecSizes[size], 0, 0,
localIsDst))
{
errors++;
}
}
}
if (errors) return -1;
return 0;
return errors ? -1 : 0;
}
int test_async_copy_global_to_local2D(cl_device_id deviceID, cl_context context,

196
test_conformance/basic/test_async_copy3D.cpp
View File

@@ -27,9 +27,14 @@
static const char *async_global_to_local_kernel3D = R"OpenCLC(
#pragma OPENCL EXTENSION cl_khr_extended_async_copies : enable
%s // optional pragma string
__kernel void test_fn(const __global %s *src, __global %s *dst, __local %s *localBuffer,
#define STRUCT_SIZE %d
typedef struct __attribute__((packed))
{
uchar byte[STRUCT_SIZE];
} VarSizeStruct __attribute__((aligned(1)));
__kernel void test_fn(const __global VarSizeStruct *src, __global VarSizeStruct *dst, __local VarSizeStruct *localBuffer,
int numElementsPerLine, int numLines, int planesCopiesPerWorkgroup,
int planesCopiesPerWorkItem, int srcLineStride,
int dstLineStride, int srcPlaneStride, int dstPlaneStride ) {
@@ -38,7 +43,9 @@ __kernel void test_fn(const __global %s *src, __global %s *dst, __local %s *loca
for (int j = 0; j < numLines; j++) {
for (int k = 0; k < numElementsPerLine; k++) {
const int index = (get_local_id(0) * planesCopiesPerWorkItem + i) * dstPlaneStride + j * dstLineStride + k;
localBuffer[index] = (%s)(%s)0;
for (int k = 0; k < STRUCT_SIZE; k++) {
localBuffer[index].byte[k] = 0;
}
}
}
}
@@ -48,7 +55,7 @@ __kernel void test_fn(const __global %s *src, __global %s *dst, __local %s *loca
event_t event = async_work_group_copy_3D3D(localBuffer, 0, src,
planesCopiesPerWorkgroup * get_group_id(0) * srcPlaneStride,
sizeof(%s), (size_t)numElementsPerLine, (size_t)numLines,
sizeof(VarSizeStruct), (size_t)numElementsPerLine, (size_t)numLines,
planesCopiesPerWorkgroup, srcLineStride, srcPlaneStride, dstLineStride,
dstPlaneStride, 0);
@@ -69,9 +76,14 @@ __kernel void test_fn(const __global %s *src, __global %s *dst, __local %s *loca
static const char *async_local_to_global_kernel3D = R"OpenCLC(
#pragma OPENCL EXTENSION cl_khr_extended_async_copies : enable
%s // optional pragma string
__kernel void test_fn(const __global %s *src, __global %s *dst, __local %s *localBuffer,
#define STRUCT_SIZE %d
typedef struct __attribute__((packed))
{
uchar byte[STRUCT_SIZE];
} VarSizeStruct __attribute__((aligned(1)));
__kernel void test_fn(const __global VarSizeStruct *src, __global VarSizeStruct *dst, __local VarSizeStruct *localBuffer,
int numElementsPerLine, int numLines, int planesCopiesPerWorkgroup,
int planesCopiesPerWorkItem, int srcLineStride,
int dstLineStride, int srcPlaneStride, int dstPlaneStride) {
@@ -80,7 +92,9 @@ __kernel void test_fn(const __global %s *src, __global %s *dst, __local %s *loca
for (int j = 0; j < numLines; j++) {
for (int k = 0; k < numElementsPerLine; k++) {
const int index = (get_local_id(0) * planesCopiesPerWorkItem + i) * srcPlaneStride + j * srcLineStride + k;
localBuffer[index] = (%s)(%s)0;
for (int k = 0; k < STRUCT_SIZE; k++) {
localBuffer[index].byte[k] = 0;
}
}
}
}
@@ -103,38 +117,25 @@ __kernel void test_fn(const __global %s *src, __global %s *dst, __local %s *loca
event_t event = async_work_group_copy_3D3D(dst,
planesCopiesPerWorkgroup * get_group_id(0) * dstPlaneStride, localBuffer, 0,
sizeof(%s), (size_t)numElementsPerLine, (size_t)numLines, planesCopiesPerWorkgroup,
sizeof(VarSizeStruct), (size_t)numElementsPerLine, (size_t)numLines, planesCopiesPerWorkgroup,
srcLineStride, srcPlaneStride, dstLineStride, dstPlaneStride, 0);
wait_group_events(1, &event);
}
)OpenCLC";
int test_copy3D(cl_device_id deviceID, cl_context context,
cl_command_queue queue, const char *kernelCode,
ExplicitType vecType, int vecSize, int srcLineMargin,
int dstLineMargin, int srcPlaneMargin, int dstPlaneMargin,
bool localIsDst)
int test_copy3D(const cl_device_id deviceID, const cl_context context,
const cl_command_queue queue, const char *const kernelCode,
const size_t elementSize, const int srcLineMargin,
const int dstLineMargin, const int srcPlaneMargin,
const int dstPlaneMargin, const bool localIsDst)
{
int error;
clProgramWrapper program;
clKernelWrapper kernel;
clMemWrapper streams[2];
size_t threads[1], localThreads[1];
void *inBuffer, *outBuffer, *outBufferCopy;
MTdata d;
char vecNameString[64];
vecNameString[0] = 0;
if (vecSize == 1)
sprintf(vecNameString, "%s", get_explicit_type_name(vecType));
else
sprintf(vecNameString, "%s%d", get_explicit_type_name(vecType),
vecSize);
size_t elementSize = get_explicit_type_size(vecType) * vecSize;
log_info("Testing %s with srcLineMargin = %d, dstLineMargin = %d, "
log_info(
"Testing %d byte element with srcLineMargin = %d, dstLineMargin = %d, "
"srcPlaneMargin = %d, dstPlaneMargin = %d\n",
vecNameString, srcLineMargin, dstLineMargin, srcPlaneMargin,
elementSize, srcLineMargin, dstLineMargin, srcPlaneMargin,
dstPlaneMargin);
cl_long max_local_mem_size;
@@ -165,20 +166,16 @@ int test_copy3D(cl_device_id deviceID, cl_context context,
test_error(error,
"clGetDeviceInfo for CL_DEVICE_MAX_COMPUTE_UNITS failed.");
char programSource[4096];
programSource[0] = 0;
char *programPtr;
char programSource[4096] = { 0 };
const char *programPtr = programSource;
sprintf(programSource, kernelCode,
vecType == kDouble ? "#pragma OPENCL EXTENSION cl_khr_fp64 : enable"
: "",
vecNameString, vecNameString, vecNameString, vecNameString,
get_explicit_type_name(vecType), vecNameString, vecNameString);
sprintf(programSource, kernelCode, elementSize);
// log_info("program: %s\n", programSource);
programPtr = programSource;
clProgramWrapper program;
clKernelWrapper kernel;
error = create_single_kernel_helper(context, &program, &kernel, 1,
(const char **)&programPtr, "test_fn");
&programPtr, "test_fn");
test_error(error, "Unable to create testing kernel");
size_t max_workgroup_size;
@@ -196,6 +193,13 @@ int test_copy3D(cl_device_id deviceID, cl_context context,
test_error(error,
"clGetDeviceInfo failed for CL_DEVICE_MAX_WORK_ITEM_SIZES");
cl_long max_work_group_size;
error = clGetDeviceInfo(deviceID, CL_DEVICE_MAX_WORK_GROUP_SIZE,
sizeof(max_work_group_size), &max_work_group_size,
NULL);
test_error(error,
"clGetDeviceInfo for CL_DEVICE_MAX_WORK_GROUP_SIZE failed.");
// Pick the minimum of the device and the kernel
if (max_workgroup_size > max_local_workgroup_size[0])
max_workgroup_size = max_local_workgroup_size[0];
@@ -208,8 +212,6 @@ int test_copy3D(cl_device_id deviceID, cl_context context,
const cl_int dstPlaneStride = (numLines * dstLineStride) + dstPlaneMargin;
const cl_int srcPlaneStride = (numLines * srcLineStride) + srcPlaneMargin;
elementSize =
get_explicit_type_size(vecType) * ((vecSize == 3) ? 4 : vecSize);
const size_t planesCopiesPerWorkItem = 2;
const size_t localStorageSpacePerWorkitem = elementSize
* planesCopiesPerWorkItem
@@ -251,9 +253,17 @@ int test_copy3D(cl_device_id deviceID, cl_context context,
const size_t globalWorkgroupSize =
numberOfLocalWorkgroups * localWorkgroupSize;
inBuffer = (void *)malloc(inBufferSize);
outBuffer = (void *)malloc(outBufferSize);
outBufferCopy = (void *)malloc(outBufferSize);
if ((localBufferSize / 4) > max_work_group_size)
{
log_info("Skipping due to resource requirements local:%db "
"max_work_group_size:%d\n",
localBufferSize, max_work_group_size);
return 0;
}
void *const inBuffer = (void *)malloc(inBufferSize);
void *const outBuffer = (void *)malloc(outBufferSize);
void *const outBufferCopy = (void *)malloc(outBufferSize);
const cl_int planesCopiesPerWorkItemInt =
static_cast<cl_int>(planesCopiesPerWorkItem);
@@ -270,18 +280,20 @@ int test_copy3D(cl_device_id deviceID, cl_context context,
(int)localBufferSize, (int)inBufferSize, (int)outBufferSize,
planesCopiesPerWorkgroup, planesCopiesPerWorkItemInt);
size_t threads[1], localThreads[1];
threads[0] = globalWorkgroupSize;
localThreads[0] = localWorkgroupSize;
d = init_genrand(gRandomSeed);
generate_random_data(
vecType, inBufferSize / get_explicit_type_size(vecType), d, inBuffer);
generate_random_data(
vecType, outBufferSize / get_explicit_type_size(vecType), d, outBuffer);
MTdata d = init_genrand(gRandomSeed);
generate_random_data(kChar, inBufferSize, d, inBuffer);
generate_random_data(kChar, outBufferSize, d, outBuffer);
free_mtdata(d);
d = NULL;
memcpy(outBufferCopy, outBuffer, outBufferSize);
clMemWrapper streams[2];
streams[0] = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR, inBufferSize,
inBuffer, &error);
test_error(error, "Unable to create input buffer");
@@ -327,8 +339,7 @@ int test_copy3D(cl_device_id deviceID, cl_context context,
// Verify
int failuresPrinted = 0;
// Verify
size_t typeSize = get_explicit_type_size(vecType) * vecSize;
for (int i = 0;
i < (int)globalWorkgroupSize * planesCopiesPerWorkItem * elementSize;
i += elementSize)
@@ -341,14 +352,13 @@ int test_copy3D(cl_device_id deviceID, cl_context context,
int inIdx = i * srcPlaneStride + j * srcLineStride + k;
int outIdx = i * dstPlaneStride + j * dstLineStride + k;
if (memcmp(((char *)inBuffer) + inIdx,
((char *)outBuffer) + outIdx, typeSize)
((char *)outBuffer) + outIdx, elementSize)
!= 0)
{
unsigned char *inchar = (unsigned char *)inBuffer + inIdx;
unsigned char *outchar =
(unsigned char *)outBuffer + outIdx;
char values[4096];
values[0] = 0;
char values[4096] = { 0 };
if (failuresPrinted == 0)
{
@@ -439,17 +449,14 @@ int test_copy3D_all_types(cl_device_id deviceID, cl_context context,
cl_command_queue queue, const char *kernelCode,
bool localIsDst)
{
ExplicitType vecType[] = {
kChar, kUChar, kShort, kUShort, kInt, kUInt, kLong,
kULong, kFloat, kDouble, kNumExplicitTypes
};
const unsigned int elemSizes[] = { 1, 2, 3, 4, 5, 6, 7,
8, 13, 16, 32, 47, 64 };
// The margins below represent the number of elements between the end of
// one line or plane and the start of the next. The strides are equivalent
// to the size of the line or plane plus the chosen margin.
unsigned int vecSizes[] = { 1, 2, 3, 4, 8, 16, 0 };
unsigned int smallTypesMarginSizes[] = { 0, 10, 100 };
unsigned int size, typeIndex, srcLineMargin, dstLineMargin, srcPlaneMargin,
dstPlaneMargin;
// one line and the start of the next. The strides are equivalent to the
// size of the line or plane plus the chosen margin.
// These have to be multipliers, because the margin must be a multiple of
// element size.
const unsigned int marginMultipliers[] = { 0, 10, 100 };
int errors = 0;
@@ -457,51 +464,28 @@ int test_copy3D_all_types(cl_device_id deviceID, cl_context context,
{
log_info(
"Device does not support extended async copies. Skipping test.\n");
return 0;
}
for (typeIndex = 0; vecType[typeIndex] != kNumExplicitTypes; typeIndex++)
else
{
if (vecType[typeIndex] == kDouble
&& !is_extension_available(deviceID, "cl_khr_fp64"))
continue;
if ((vecType[typeIndex] == kLong || vecType[typeIndex] == kULong)
&& !gHasLong)
continue;
for (size = 0; vecSizes[size] != 0; size++)
for (const unsigned int elemSize : elemSizes)
{
if (get_explicit_type_size(vecType[typeIndex]) * vecSizes[size]
<= 2) // small type
for (const unsigned int srcLineMarginMultiplier : marginMultipliers)
{
for (srcLineMargin = 0;
srcLineMargin < sizeof(smallTypesMarginSizes)
/ sizeof(smallTypesMarginSizes[0]);
srcLineMargin++)
for (const unsigned int dstLineMarginMultiplier :
marginMultipliers)
{
for (dstLineMargin = 0;
dstLineMargin < sizeof(smallTypesMarginSizes)
/ sizeof(smallTypesMarginSizes[0]);
dstLineMargin++)
for (const unsigned int srcPlaneMarginMultiplier :
marginMultipliers)
{
for (srcPlaneMargin = 0;
srcPlaneMargin < sizeof(smallTypesMarginSizes)
/ sizeof(smallTypesMarginSizes[0]);
srcPlaneMargin++)
for (const unsigned int dstPlaneMarginMultiplier :
marginMultipliers)
{
for (dstPlaneMargin = 0;
dstPlaneMargin < sizeof(smallTypesMarginSizes)
/ sizeof(smallTypesMarginSizes[0]);
dstPlaneMargin++)
{
if (test_copy3D(
deviceID, context, queue, kernelCode,
vecType[typeIndex], vecSizes[size],
smallTypesMarginSizes[srcLineMargin],
smallTypesMarginSizes[dstLineMargin],
smallTypesMarginSizes[srcPlaneMargin],
smallTypesMarginSizes[dstPlaneMargin],
if (test_copy3D(deviceID, context, queue,
kernelCode, elemSize,
srcLineMarginMultiplier * elemSize,
dstLineMarginMultiplier * elemSize,
srcPlaneMarginMultiplier * elemSize,
dstPlaneMarginMultiplier * elemSize,
localIsDst))
{
errors++;
@@ -511,14 +495,6 @@ int test_copy3D_all_types(cl_device_id deviceID, cl_context context,
}
}
}
// not a small type, check only zero stride
else if (test_copy3D(deviceID, context, queue, kernelCode,
vecType[typeIndex], vecSizes[size], 0, 0, 0, 0,
localIsDst))
{
errors++;
}
}
}
if (errors) return -1;
return 0;