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

Test all cluster sizes for cl_khr_subgroup_clustered_reduce (#1408)

Browse Source 
Signed-off-by: Stuart Brady <stuart.brady@arm.com>
This commit is contained in:
Stuart Brady
2022-04-12 17:42:55 +01:00
committed by GitHub
parent 2fcdde96d2
commit 7a0e7e767a
1 changed files with 25 additions and 14 deletions
Download Patch File Download Diff File Expand all files Collapse all files

39
test_conformance/subgroups/test_subgroup_clustered_reduce.cpp
View File

@@ -18,19 +18,29 @@
#include "subgroup_common_templates.h" #include "subgroup_common_templates.h"
#include "harness/typeWrappers.h" #include "harness/typeWrappers.h"
#define CLUSTER_SIZE 4
#define CLUSTER_SIZE_STR "4"
namespace { namespace {
std::string sub_group_clustered_reduce_source = R"( std::string sub_group_clustered_reduce_source = R"(
__kernel void test_%s(const __global Type *in, __global int4 *xy, __global Type *out) { __kernel void test_%s(const __global Type *in, __global int4 *xy, __global Type *out,
uint cluster_size) {
Type r;
int gid = get_global_id(0); int gid = get_global_id(0);
XY(xy,gid); XY(xy,gid);
xy[gid].w = 0; xy[gid].w = 0;
if (sizeof(in[gid]) == sizeof(%s(in[gid], )" CLUSTER_SIZE_STR R"())) { Type v = in[gid];
if (sizeof(in[gid]) == sizeof(%s(v, 1))) {
xy[gid].w = sizeof(in[gid]); xy[gid].w = sizeof(in[gid]);
} }
out[gid] = %s(in[gid], )" CLUSTER_SIZE_STR R"(); switch (cluster_size) {
case 1: r = %s(v, 1); break;
case 2: r = %s(v, 2); break;
case 4: r = %s(v, 4); break;
case 8: r = %s(v, 8); break;
case 16: r = %s(v, 16); break;
case 32: r = %s(v, 32); break;
case 64: r = %s(v, 64); break;
case 128: r = %s(v, 128); break;
}
out[gid] = r;
} }
)"; )";
@@ -94,32 +104,33 @@ template <typename Ty, ArithmeticOp operation> struct RED_CLU
int n = ii + ns > nw ? nw - ii : ns; int n = ii + ns > nw ? nw - ii : ns;
int midx = 4 * ii + 2; int midx = 4 * ii + 2;
std::vector<Ty> clusters_results; std::vector<Ty> clusters_results;
int clusters_counter = ns / CLUSTER_SIZE; int clusters_counter = ns / test_params.cluster_size;
clusters_results.resize(clusters_counter); clusters_results.resize(clusters_counter);
// Compute target // Compute target
Ty tr = mx[ii]; Ty tr = mx[ii];
for (int i = 0; i < n; ++i) for (int i = 0; i < n; ++i)
{ {
if (i % CLUSTER_SIZE == 0) if (i % test_params.cluster_size == 0)
tr = mx[ii + i]; tr = mx[ii + i];
else else
tr = calculate<Ty>(tr, mx[ii + i], operation); tr = calculate<Ty>(tr, mx[ii + i], operation);
clusters_results[i / CLUSTER_SIZE] = tr; clusters_results[i / test_params.cluster_size] = tr;
} }
// Check result // Check result
for (int i = 0; i < n; ++i) for (int i = 0; i < n; ++i)
{ {
Ty rr = my[ii + i]; Ty rr = my[ii + i];
tr = clusters_results[i / CLUSTER_SIZE]; tr = clusters_results[i / test_params.cluster_size];
if (!compare(rr, tr)) if (!compare(rr, tr))
{ {
log_error( log_error(
"ERROR: sub_group_clustered_reduce_%s(%s) mismatch " "ERROR: sub_group_clustered_reduce_%s(%s, %u) "
"for local id %d in sub group %d in group %d\n", "mismatch for local id %d in sub group %d in group "
"%d\n",
operation_names(operation), TypeManager<Ty>::name(), operation_names(operation), TypeManager<Ty>::name(),
i, j, k); test_params.cluster_size, i, j, k);
return TEST_FAIL; return TEST_FAIL;
} }
} }
@@ -184,7 +195,7 @@ int test_subgroup_functions_clustered_reduce(cl_device_id device,
constexpr size_t global_work_size = 2000; constexpr size_t global_work_size = 2000;
constexpr size_t local_work_size = 200; constexpr size_t local_work_size = 200;
WorkGroupParams test_params(global_work_size, local_work_size); WorkGroupParams test_params(global_work_size, local_work_size, -1, 3);
test_params.save_kernel_source(sub_group_clustered_reduce_source); test_params.save_kernel_source(sub_group_clustered_reduce_source);
RunTestForType rft(device, context, queue, num_elements, test_params); RunTestForType rft(device, context, queue, num_elements, test_params);