Cold-Air-Networks/mirror-opencl-cts
2
0
Fork 0
Code Issues Pull requests Projects Releases Packages Wiki Activity Actions

vectors: Parallelize alignment tests (#2715)

Browse source 
Modernizes the vector alignment conformance tests by replacing slow
nested procedural loops with thread pool parallelization, significantly
improving test execution times on multi-core systems.

Merged as agreed on 14Jul26 call
This commit is contained in:
Romaric Jodin 2026-07-14 17:46:22 +02:00 committed by GitHub
parent f6d1e8711f
commit 790f45d4cd
No known key found for this signature in database
GPG key ID: B5690EEEBB952194
4 changed files with 157 additions and 174 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

13
ci/pocl/golden.json
View file

@ -140,5 +140,18 @@
"--num-worker-threads 2 CL_FILTER_NEAREST CL_R": {
"1D": "fail"
}
},
"test_vectors": {
"--wimpy": {
"step_type": "pass",
"step_var": "pass",
"step_typedef_type": "pass",
"step_typedef_var": "pass",
"vec_align_array": "pass",
"vec_align_struct": "pass",
"vec_align_packed_struct": "pass",
"vec_align_struct_arr": "pass",
"vec_align_packed_struct_arr": "pass"
}
}
}

12
test_conformance/vectors/defines.h
View file

@ -24,17 +24,17 @@
// 1,2,4,8,16,3
#define NUM_VECTOR_SIZES 6
extern int g_arrVecSizes[NUM_VECTOR_SIZES];
extern int g_arrVecSteps[NUM_VECTOR_SIZES];
extern const int g_arrVecSizes[NUM_VECTOR_SIZES];
extern const int g_arrVecSteps[NUM_VECTOR_SIZES];
extern const char *g_arrVecSizeNames[NUM_VECTOR_SIZES];
extern size_t g_arrVecAlignMasks[NUM_VECTOR_SIZES];
extern const char* const g_arrVecSizeNames[NUM_VECTOR_SIZES];
extern const size_t g_arrVecAlignMasks[NUM_VECTOR_SIZES];
// Define the buffer size that we want to block our test with
#define BUFFER_SIZE (1024 * 1024)
#define KPAGESIZE 4096
extern ExplicitType types[];
extern const ExplicitType types[];
extern const char *g_arrTypeNames[];
extern const char* const g_arrTypeNames[];
extern const size_t g_arrTypeSizes[];

18
test_conformance/vectors/globals.cpp
View file

@ -18,12 +18,11 @@
// 1,2,3,4,8,16 or
// 1,2,4,8,16,3
int g_arrVecSizes[NUM_VECTOR_SIZES] = { 1, 2, 3, 4, 8, 16 };
int g_arrVecSteps[NUM_VECTOR_SIZES] = { 1, 2, 4, 4, 8, 16 };
const char *g_arrVecSizeNames[NUM_VECTOR_SIZES] = {
"", "2", "3", "4", "8", "16"
};
size_t g_arrVecAlignMasks[NUM_VECTOR_SIZES] = {
const int g_arrVecSizes[NUM_VECTOR_SIZES] = { 1, 2, 3, 4, 8, 16 };
const int g_arrVecSteps[NUM_VECTOR_SIZES] = { 1, 2, 4, 4, 8, 16 };
const char* const g_arrVecSizeNames[NUM_VECTOR_SIZES] = { "", "2", "3",
"4", "8", "16" };
const size_t g_arrVecAlignMasks[NUM_VECTOR_SIZES] = {
(size_t)0,
(size_t)0x1, // 2
(size_t)0x3, // 3
@ -32,13 +31,14 @@ size_t g_arrVecAlignMasks[NUM_VECTOR_SIZES] = {
(size_t)0xf // 16
};
ExplicitType types[] = {
const ExplicitType types[] = {
kChar, kUChar, kShort, kUShort, kInt, kUInt, kLong,
kULong, kFloat, kDouble, kNumExplicitTypes
};
const char *g_arrTypeNames[] = { "char", "uchar", "short", "ushort", "int",
"uint", "long", "ulong", "float", "double" };
const char* const g_arrTypeNames[] = { "char", "uchar", "short", "ushort",
"int", "uint", "long", "ulong",
"float", "double" };
extern const size_t g_arrTypeSizes[] = { 1, 1, 2, 2, 4, 4, 8, 8, 4, 8 };

288
test_conformance/vectors/test_vec_align.cpp
View file

@ -19,6 +19,7 @@
#include "harness/conversions.h"
#include "harness/typeWrappers.h"
#include "harness/testHarness.h"
#include "harness/ThreadPool.h"
#include "structs.h"
@ -26,6 +27,7 @@
#include "type_replacer.h"
#include <array>
size_t get_align(size_t vecSize)
{
@ -77,7 +79,7 @@ int test_vec_internal(cl_device_id deviceID, cl_context context,
cl_command_queue queue, const char* pattern,
const char* testName, size_t bufSize, size_t preSize,
size_t typeMultiplePreSize, size_t postSize,
size_t typeMultiplePostSize)
size_t typeMultiplePostSize, bool supports_fp64)
{
int err;
int typeIdx, vecSizeIdx;
@ -104,7 +106,7 @@ int test_vec_internal(cl_device_id deviceID, cl_context context,
// Skip doubles if it is not supported otherwise enable pragma
if (types[typeIdx] == kDouble)
{
if (!is_extension_available(deviceID, "cl_khr_fp64"))
if (!supports_fp64)
{
continue;
}
@ -243,7 +245,7 @@ int test_vec_internal(cl_device_id deviceID, cl_context context,
}
static const char* patterns[] = {
static const std::vector<const char*> patterns = {
".PRAGMA..STATE.\n"
"__kernel void test_vec_align_array(.SRC_SCOPE. .TYPE..NUM. *source, "
".DST_SCOPE. uint *dest)\n"
@ -306,50 +308,108 @@ static const char* patterns[] = {
// __attribute__ ((packed))
};
constexpr unsigned ARR_SIZE = 6;
const char* pre_substitution_arr[] = { "",
"char c;\n",
"short3 s;",
".TYPE.3 tPre;\n",
".TYPE. arrPre[5];\n",
".TYPE. arrPre[12];\n",
NULL };
static const std::array<const char*, ARR_SIZE> pre_substitution_arr = {
"",
"char c;\n",
"short3 s;",
".TYPE.3 tPre;\n",
".TYPE. arrPre[5];\n",
".TYPE. arrPre[12];\n"
};
// alignments of everything in pre_substitution_arr as raw alignments
// 0 if such a thing is meaningless
size_t pre_align_arr[] = { 0,
sizeof(cl_char),
4 * sizeof(cl_short),
0, // taken care of in type_multiple_pre_align_arr
0,
0 };
static const std::array<size_t, ARR_SIZE> pre_align_arr = {
0,
sizeof(cl_char),
4 * sizeof(cl_short),
0, // taken care of in type_multiple_pre_align_arr
0,
0
};
// alignments of everything in pre_substitution_arr as multiples of
// sizeof(.TYPE.)
// 0 if such a thing is meaningless
size_t type_multiple_pre_align_arr[] = { 0, 0, 0, 4, 5, 12 };
static const std::array<size_t, ARR_SIZE> type_multiple_pre_align_arr = {
0, 0, 0, 4, 5, 12
};
const char* post_substitution_arr[] = { "",
"char cPost;\n",
".TYPE. arrPost[3];\n",
".TYPE. arrPost[5];\n",
".TYPE.3 arrPost;\n",
".TYPE. arrPost[12];\n",
NULL };
static const std::array<const char*, ARR_SIZE> post_substitution_arr = {
"",
"char cPost;\n",
".TYPE. arrPost[3];\n",
".TYPE. arrPost[5];\n",
".TYPE.3 arrPost;\n",
".TYPE. arrPost[12];\n"
};
// alignments of everything in post_substitution_arr as raw alignments
// 0 if such a thing is meaningless
size_t post_align_arr[] = { 0, sizeof(cl_char),
0, // taken care of in type_multiple_post_align_arr
0, 0,
0 };
static const std::array<size_t, ARR_SIZE> post_align_arr = {
0, sizeof(cl_char),
0, // taken care of in type_multiple_post_align_arr
0, 0,
0
};
// alignments of everything in post_substitution_arr as multiples of
// sizeof(.TYPE.)
// 0 if such a thing is meaningless
size_t type_multiple_post_align_arr[] = { 0, 0, 3, 5, 4, 12 };
static const std::array<size_t, ARR_SIZE> type_multiple_post_align_arr = {
0, 0, 3, 5, 4, 12
};
struct test_vec_thread_info
{
cl_device_id device;
cl_context context;
cl_command_queue queue;
const char* testName;
size_t bufSize;
bool packed;
const char* source;
bool supports_fp64;
};
cl_int test_vec_thread(cl_uint job_id, cl_uint thread_id, void* userInfo)
{
test_vec_thread_info* info = (test_vec_thread_info*)userInfo;
char tmp[2048];
int preIdx = job_id / ARR_SIZE;
int postIdx = job_id % ARR_SIZE;
size_t preSize = 0;
size_t typeMultiplePreSize = 0;
size_t postSize = 0;
size_t typeMultiplePostSize = 0;
if (info->packed)
{
preSize = pre_align_arr[preIdx];
typeMultiplePreSize = type_multiple_pre_align_arr[preIdx];
postSize = post_align_arr[postIdx];
typeMultiplePostSize = type_multiple_post_align_arr[postIdx];
}
const char* replaceWith1 = pre_substitution_arr[preIdx];
const char* replaceWith2 = post_substitution_arr[postIdx];
doReplace(tmp, (size_t)2048, info->source, ".PRE.", replaceWith1, ".POST.",
replaceWith2);
return test_vec_internal(info->device, info->context, info->queue, tmp,
info->testName, info->bufSize, preSize,
typeMultiplePreSize, postSize,
typeMultiplePostSize, info->supports_fp64);
}
bool supports_fp64(cl_device_id device)
{
return is_extension_available(device, "cl_khr_fp64");
}
// there hsould be a packed version of this?
REGISTER_TEST(vec_align_array)
@ -360,176 +420,86 @@ REGISTER_TEST(vec_align_array)
log_info("Testing global\n");
doReplace(tmp, (size_t)2048, patterns[0], ".SRC_SCOPE.", "__global",
".DST_SCOPE.", "__global"); //
result = test_vec_internal(device, context, queue, tmp,
"test_vec_align_array", BUFFER_SIZE, 0, 0, 0, 0);
result =
test_vec_internal(device, context, queue, tmp, "test_vec_align_array",
BUFFER_SIZE, 0, 0, 0, 0, supports_fp64(device));
return result;
}
REGISTER_TEST(vec_align_struct)
{
char tmp1[2048], tmp2[2048];
int result = 0;
int preIdx, postIdx;
char tmp[2048];
log_info("testing __private\n");
doReplace(tmp2, (size_t)2048, patterns[1], ".SRC_SCOPE.", "__private",
doReplace(tmp, (size_t)2048, patterns[1], ".SRC_SCOPE.", "__private",
".DST_SCOPE.", "__global"); //
for (preIdx = 0; pre_substitution_arr[preIdx] != NULL; ++preIdx)
test_vec_thread_info info{ device, context, queue, "test_vec_align_struct",
512, false, tmp, supports_fp64(device) };
cl_int result = ThreadPool_Do(test_vec_thread, ARR_SIZE * ARR_SIZE, &info);
if (result != CL_SUCCESS)
{
for (postIdx = 0; post_substitution_arr[postIdx] != NULL; ++postIdx)
{
doReplace(tmp1, (size_t)2048, tmp2, ".PRE.",
pre_substitution_arr[preIdx], ".POST.",
post_substitution_arr[postIdx]);
result =
test_vec_internal(device, context, queue, tmp1,
"test_vec_align_struct", 512, 0, 0, 0, 0);
if (result != 0)
{
return result;
}
}
return result;
}
log_info("testing __local\n");
doReplace(tmp2, (size_t)2048, patterns[1], ".SRC_SCOPE.", "__local",
doReplace(tmp, (size_t)2048, patterns[1], ".SRC_SCOPE.", "__local",
".DST_SCOPE.", "__global"); //
for (preIdx = 0; pre_substitution_arr[preIdx] != NULL; ++preIdx)
{
for (postIdx = 0; post_substitution_arr[postIdx] != NULL; ++postIdx)
{
doReplace(tmp1, (size_t)2048, tmp2, ".PRE.",
pre_substitution_arr[preIdx], ".POST.",
post_substitution_arr[postIdx]);
result =
test_vec_internal(device, context, queue, tmp1,
"test_vec_align_struct", 512, 0, 0, 0, 0);
if (result != 0)
{
return result;
}
}
}
return 0;
info.testName = "test_vec_align_struct";
return ThreadPool_Do(test_vec_thread, ARR_SIZE * ARR_SIZE, &info);
}
REGISTER_TEST(vec_align_packed_struct)
{
char tmp1[2048], tmp2[2048];
int result = 0;
int preIdx, postIdx;
char tmp[2048];
log_info("Testing __private\n");
doReplace(tmp2, (size_t)2048, patterns[2], ".SRC_SCOPE.", "__private",
doReplace(tmp, (size_t)2048, patterns[2], ".SRC_SCOPE.", "__private",
".DST_SCOPE.", "__global"); //
for (preIdx = 0; pre_substitution_arr[preIdx] != NULL; ++preIdx)
test_vec_thread_info info{ device, context,
queue, "test_vec_align_packed_struct",
512, true,
tmp, supports_fp64(device) };
cl_int result = ThreadPool_Do(test_vec_thread, ARR_SIZE * ARR_SIZE, &info);
if (result != CL_SUCCESS)
{
for (postIdx = 0; post_substitution_arr[postIdx] != NULL; ++postIdx)
{
doReplace(tmp1, (size_t)2048, tmp2, ".PRE.",
pre_substitution_arr[preIdx], ".POST.",
post_substitution_arr[postIdx]);
result = test_vec_internal(
device, context, queue, tmp1, "test_vec_align_packed_struct",
512, pre_align_arr[preIdx], type_multiple_pre_align_arr[preIdx],
post_align_arr[postIdx], type_multiple_post_align_arr[postIdx]);
if (result != 0)
{
return result;
}
}
return result;
}
log_info("testing __local\n");
doReplace(tmp2, (size_t)2048, patterns[2], ".SRC_SCOPE.", "__local",
doReplace(tmp, (size_t)2048, patterns[2], ".SRC_SCOPE.", "__local",
".DST_SCOPE.", "__global"); //
for (preIdx = 0; pre_substitution_arr[preIdx] != NULL; ++preIdx)
{
for (postIdx = 0; post_substitution_arr[postIdx] != NULL; ++postIdx)
{
doReplace(tmp1, (size_t)2048, tmp2, ".PRE.",
pre_substitution_arr[preIdx], ".POST.",
post_substitution_arr[postIdx]);
result = test_vec_internal(
device, context, queue, tmp1, "test_vec_align_packed_struct",
512, pre_align_arr[preIdx], type_multiple_pre_align_arr[preIdx],
post_align_arr[postIdx], type_multiple_post_align_arr[postIdx]);
if (result != 0)
{
return result;
}
}
}
return 0;
info.testName = "test_vec_align_packed_struct";
return ThreadPool_Do(test_vec_thread, ARR_SIZE * ARR_SIZE, &info);
}
REGISTER_TEST(vec_align_struct_arr)
{
char tmp1[2048], tmp2[2048];
int result = 0;
int preIdx, postIdx;
char tmp[2048];
log_info("testing __global\n");
doReplace(tmp2, (size_t)2048, patterns[3], ".SRC_SCOPE.", "__global",
doReplace(tmp, (size_t)2048, patterns[3], ".SRC_SCOPE.", "__global",
".DST_SCOPE.", "__global"); //
for (preIdx = 0; pre_substitution_arr[preIdx] != NULL; ++preIdx)
{
for (postIdx = 0; post_substitution_arr[postIdx] != NULL; ++postIdx)
{
doReplace(tmp1, (size_t)2048, tmp2, ".PRE.",
pre_substitution_arr[preIdx], ".POST.",
post_substitution_arr[postIdx]);
result = test_vec_internal(device, context, queue, tmp1,
"test_vec_align_struct_arr", BUFFER_SIZE,
0, 0, 0, 0);
if (result != 0)
{
return result;
}
}
}
return 0;
test_vec_thread_info info{ device, context,
queue, "test_vec_align_struct_arr",
BUFFER_SIZE, false,
tmp, supports_fp64(device) };
return ThreadPool_Do(test_vec_thread, ARR_SIZE * ARR_SIZE, &info);
}
REGISTER_TEST(vec_align_packed_struct_arr)
{
char tmp1[2048], tmp2[2048];
int result = 0;
int preIdx, postIdx;
char tmp[2048];
log_info("Testing __global\n");
doReplace(tmp2, (size_t)2048, patterns[4], ".SRC_SCOPE.", "__global",
doReplace(tmp, (size_t)2048, patterns[4], ".SRC_SCOPE.", "__global",
".DST_SCOPE.", "__global"); //
for (preIdx = 0; pre_substitution_arr[preIdx] != NULL; ++preIdx)
{
for (postIdx = 0; post_substitution_arr[postIdx] != NULL; ++postIdx)
{
doReplace(tmp1, (size_t)2048, tmp2, ".PRE.",
pre_substitution_arr[preIdx], ".POST.",
post_substitution_arr[postIdx]);
result = test_vec_internal(
device, context, queue, tmp1,
"test_vec_align_packed_struct_arr", BUFFER_SIZE,
pre_align_arr[preIdx], type_multiple_pre_align_arr[preIdx],
post_align_arr[postIdx], type_multiple_post_align_arr[postIdx]);
if (result != 0) return result;
}
}
return 0;
test_vec_thread_info info{ device, context,
queue, "test_vec_align_packed_struct_arr",
BUFFER_SIZE, true,
tmp, supports_fp64(device) };
return ThreadPool_Do(test_vec_thread, ARR_SIZE * ARR_SIZE, &info);
}