ahmed/OpenCL-CTS
1
0
Fork 0
mirror of https://github.com/KhronosGroup/OpenCL-CTS.git synced 2026-03-25 08:19:02 +00:00
Code Packages Projects Releases Wiki Activity

Improve cl_khr_subgroup_shuffle* test coverage (#1402)

Browse Source 
Test cases where the index/mask/delta is greater than or equal to the
maximum subgroup size.  These are cases that return undefined results
but are not undefined behavior.

The index/mask/delta values now include values less than twice the
subgroup size, and 0xffffffff.

Testing for sub_group_shuffle_xor() already allowed inputs that were
greater or equal to the subgroup size for the last subgroup in a
workgroup, but did not properly account for this in the verification
function, potentially resulting in out of bounds accesses.

Signed-off-by: Stuart Brady <stuart.brady@arm.com>
This commit is contained in:
Stuart Brady
2022-03-29 19:39:06 +01:00
committed by GitHub
parent f6dbc5b9b5
commit c42cf518da
1 changed files with 34 additions and 46 deletions
Download Patch File Download Diff File Expand all files Collapse all files

80
test_conformance/subgroups/subgroup_common_templates.h
View File

@@ -481,12 +481,12 @@ template <typename Ty, ShuffleOp operation> struct SHF
static void gen(Ty *x, Ty *t, cl_int *m, const WorkGroupParams &test_params) static void gen(Ty *x, Ty *t, cl_int *m, const WorkGroupParams &test_params)
{ {
int i, ii, j, k, l, n, delta; int i, ii, j, k, n, delta;
cl_uint l;
int nw = test_params.local_workgroup_size; int nw = test_params.local_workgroup_size;
int ns = test_params.subgroup_size; int ns = test_params.subgroup_size;
int ng = test_params.global_workgroup_size; int ng = test_params.global_workgroup_size;
int nj = (nw + ns - 1) / ns; int nj = (nw + ns - 1) / ns;
int d = ns > 100 ? 100 : ns;
ii = 0; ii = 0;
ng = ng / nw; ng = ng / nw;
for (k = 0; k < ng; ++k) for (k = 0; k < ng; ++k)
@@ -498,33 +498,10 @@ template <typename Ty, ShuffleOp operation> struct SHF
for (i = 0; i < n; ++i) for (i = 0; i < n; ++i)
{ {
int midx = 4 * ii + 4 * i + 2; int midx = 4 * ii + 4 * i + 2;
l = (int)(genrand_int32(gMTdata) & 0x7fffffff) l = (((cl_uint)(genrand_int32(gMTdata) & 0x7fffffff) + 1)
% (d > n ? n : d); % (ns * 2 + 1))
switch (operation) - 1;
{ m[midx] = l;
case ShuffleOp::shuffle:
case ShuffleOp::shuffle_xor:
// storing information about shuffle index
m[midx] = (cl_int)l;
break;
case ShuffleOp::shuffle_up:
delta = l; // calculate delta for shuffle up
if (i - delta < 0)
{
delta = i;
}
m[midx] = (cl_int)delta;
break;
case ShuffleOp::shuffle_down:
delta = l; // calculate delta for shuffle down
if (i + delta >= n)
{
delta = n - 1 - i;
}
m[midx] = (cl_int)delta;
break;
default: break;
}
cl_ulong number = genrand_int64(gMTdata); cl_ulong number = genrand_int64(gMTdata);
set_value(t[ii + i], number); set_value(t[ii + i], number);
} }
@@ -542,7 +519,8 @@ template <typename Ty, ShuffleOp operation> struct SHF
static test_status chk(Ty *x, Ty *y, Ty *mx, Ty *my, cl_int *m, static test_status chk(Ty *x, Ty *y, Ty *mx, Ty *my, cl_int *m,
const WorkGroupParams &test_params) const WorkGroupParams &test_params)
{ {
int ii, i, j, k, l, n; int ii, i, j, k, n;
cl_uint l;
int nw = test_params.local_workgroup_size; int nw = test_params.local_workgroup_size;
int ns = test_params.subgroup_size; int ns = test_params.subgroup_size;
int ng = test_params.global_workgroup_size; int ng = test_params.global_workgroup_size;
@@ -567,32 +545,42 @@ template <typename Ty, ShuffleOp operation> struct SHF
{ // inside the subgroup { // inside the subgroup
// shuffle index storage // shuffle index storage
int midx = 4 * ii + 4 * i + 2; int midx = 4 * ii + 4 * i + 2;
l = (int)m[midx]; l = m[midx];
rr = my[ii + i]; rr = my[ii + i];
cl_uint tr_idx;
bool skip = false;
switch (operation) switch (operation)
{ {
// shuffle basic - treat l as index // shuffle basic - treat l as index
case ShuffleOp::shuffle: tr = mx[ii + l]; break; case ShuffleOp::shuffle: tr_idx = l; break;
// shuffle up - treat l as delta
case ShuffleOp::shuffle_up: tr = mx[ii + i - l]; break;
// shuffle up - treat l as delta
case ShuffleOp::shuffle_down:
tr = mx[ii + i + l];
break;
// shuffle xor - treat l as mask // shuffle xor - treat l as mask
case ShuffleOp::shuffle_xor: case ShuffleOp::shuffle_xor: tr_idx = i ^ l; break;
tr = mx[ii + (i ^ l)]; // shuffle up - treat l as delta
case ShuffleOp::shuffle_up:
if (l >= ns) skip = true;
tr_idx = i - l;
break;
// shuffle down - treat l as delta
case ShuffleOp::shuffle_down:
if (l >= ns) skip = true;
tr_idx = i + l;
break; break;
default: break; default: break;
} }
if (!compare(rr, tr)) if (!skip && tr_idx < n)
{ {
log_error("ERROR: sub_group_%s(%s) mismatch for " tr = mx[ii + tr_idx];
"local id %d in sub group %d in group %d\n",
operation_names(operation), if (!compare(rr, tr))
TypeManager<Ty>::name(), i, j, k); {
return TEST_FAIL; log_error("ERROR: sub_group_%s(%s) mismatch for "
"local id %d in sub group %d in group "
"%d\n",
operation_names(operation),
TypeManager<Ty>::name(), i, j, k);
return TEST_FAIL;
}
} }
} }
} }