ahmed/OpenCL-CTS
1
0
Fork 0
mirror of https://github.com/KhronosGroup/OpenCL-CTS.git synced 2026-03-18 22:09:01 +00:00
Code Issues Packages Projects Releases Wiki Activity

Migrate workgroups suite to the new test registration framework (#2308)

Browse Source 
Contributes to #2181

Signed-off-by: Kévin Petit <kpet@free.fr>
This commit is contained in:
Kévin Petit
2025-03-11 22:00:38 +00:00
committed by GitHub
parent 641ffeaa67
commit bb604702a5
7 changed files with 133 additions and 261 deletions
Download Patch File Download Diff File Expand all files Collapse all files

24
test_conformance/workgroups/main.cpp
View File

@@ -16,32 +16,12 @@
#include "harness/compat.h"
#include "harness/testHarness.h"
#include "procs.h"
#include <stdio.h>
#include <string.h>
#if !defined(_WIN32)
#include <unistd.h>
#endif
test_definition test_list[] = {
ADD_TEST_VERSION(work_group_all, Version(2, 0)),
ADD_TEST_VERSION(work_group_any, Version(2, 0)),
ADD_TEST_VERSION(work_group_reduce_add, Version(2, 0)),
ADD_TEST_VERSION(work_group_reduce_min, Version(2, 0)),
ADD_TEST_VERSION(work_group_reduce_max, Version(2, 0)),
ADD_TEST_VERSION(work_group_scan_inclusive_add, Version(2, 0)),
ADD_TEST_VERSION(work_group_scan_inclusive_min, Version(2, 0)),
ADD_TEST_VERSION(work_group_scan_inclusive_max, Version(2, 0)),
ADD_TEST_VERSION(work_group_scan_exclusive_add, Version(2, 0)),
ADD_TEST_VERSION(work_group_scan_exclusive_min, Version(2, 0)),
ADD_TEST_VERSION(work_group_scan_exclusive_max, Version(2, 0)),
ADD_TEST_VERSION(work_group_broadcast_1D, Version(2, 0)),
ADD_TEST_VERSION(work_group_broadcast_2D, Version(2, 0)),
ADD_TEST_VERSION(work_group_broadcast_3D, Version(2, 0)),
};
const int test_num = ARRAY_SIZE(test_list);
test_status InitCL(cl_device_id device) {
auto version = get_device_cl_version(device);
auto expected_min_version = Version(1, 2);
@@ -77,6 +57,8 @@ test_status InitCL(cl_device_id device) {
}
int main(int argc, const char *argv[]) {
return runTestHarnessWithCheck(argc, argv, test_num, test_list, false, 0, InitCL);
return runTestHarnessWithCheck(
argc, argv, test_registry::getInstance().num_tests(),
test_registry::getInstance().definitions(), false, 0, InitCL);
}

74
test_conformance/workgroups/procs.h
View File

@@ -1,74 +0,0 @@
//
// Copyright (c) 2017, 2021 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 "harness/testHarness.h"
#include "harness/kernelHelpers.h"
#include "harness/errorHelpers.h"
#include "harness/typeWrappers.h"
#include "harness/conversions.h"
#include "harness/mt19937.h"
extern int create_program_and_kernel(const char *source,
const char *kernel_name,
cl_program *program_ret,
cl_kernel *kernel_ret);
extern int test_work_group_all(cl_device_id deviceID, cl_context context,
cl_command_queue queue, int num_elements);
extern int test_work_group_any(cl_device_id deviceID, cl_context context,
cl_command_queue queue, int num_elements);
extern int test_work_group_broadcast_1D(cl_device_id deviceID,
cl_context context,
cl_command_queue queue,
int num_elements);
extern int test_work_group_broadcast_2D(cl_device_id deviceID,
cl_context context,
cl_command_queue queue,
int num_elements);
extern int test_work_group_broadcast_3D(cl_device_id deviceID,
cl_context context,
cl_command_queue queue,
int num_elements);
extern int test_work_group_reduce_add(cl_device_id deviceID, cl_context context,
cl_command_queue queue, int num_elements);
extern int test_work_group_reduce_min(cl_device_id deviceID, cl_context context,
cl_command_queue queue, int num_elements);
extern int test_work_group_reduce_max(cl_device_id deviceID, cl_context context,
cl_command_queue queue, int num_elements);
extern int test_work_group_scan_exclusive_add(cl_device_id deviceID,
cl_context context,
cl_command_queue queue,
int num_elements);
extern int test_work_group_scan_exclusive_min(cl_device_id deviceID,
cl_context context,
cl_command_queue queue,
int num_elements);
extern int test_work_group_scan_exclusive_max(cl_device_id deviceID,
cl_context context,
cl_command_queue queue,
int num_elements);
extern int test_work_group_scan_inclusive_add(cl_device_id deviceID,
cl_context context,
cl_command_queue queue,
int num_elements);
extern int test_work_group_scan_inclusive_min(cl_device_id deviceID,
cl_context context,
cl_command_queue queue,
int num_elements);
extern int test_work_group_scan_inclusive_max(cl_device_id deviceID,
cl_context context,
cl_command_queue queue,
int num_elements);

11
test_conformance/workgroups/testBase.h
View File

@@ -16,6 +16,12 @@
#ifndef _testBase_h
#define _testBase_h
#include "harness/testHarness.h"
#include "harness/kernelHelpers.h"
#include "harness/errorHelpers.h"
#include "harness/typeWrappers.h"
#include "harness/conversions.h"
#include "harness/mt19937.h"
#include "harness/compat.h"
#include <stdio.h>
@@ -23,9 +29,4 @@
#include <sys/types.h>
#include <sys/stat.h>
#include "procs.h"
#endif // _testBase_h

15
test_conformance/workgroups/test_wg_all.cpp
View File

@@ -20,8 +20,7 @@
#include <sys/types.h>
#include <sys/stat.h>
#include "procs.h"
#include "testBase.h"
const char *wg_all_kernel_code =
"__kernel void test_wg_all(global float *input, global int *output)\n"
@@ -65,8 +64,7 @@ verify_wg_all(float *inptr, int *outptr, size_t n, size_t wg_size)
return 0;
}
int
test_work_group_all(cl_device_id device, cl_context context, cl_command_queue queue, int n_elems)
REGISTER_TEST_VERSION(work_group_all, Version(2, 0))
{
cl_mem streams[2];
cl_float *input_ptr[1], *p;
@@ -74,8 +72,7 @@ test_work_group_all(cl_device_id device, cl_context context, cl_command_queue qu
cl_program program;
cl_kernel kernel;
size_t threads[1];
size_t wg_size[1];
size_t num_elements;
size_t wg_size[1];
int err;
MTdata d;
@@ -88,8 +85,6 @@ test_work_group_all(cl_device_id device, cl_context context, cl_command_queue qu
err = get_max_allowed_1d_work_group_size_on_device(device, kernel, wg_size);
test_error(err, "get_max_allowed_1d_work_group_size_on_device failed");
num_elements = n_elems;
input_ptr[0] = (cl_float*)malloc(sizeof(cl_float) * (num_elements+1));
output_ptr = (cl_int*)malloc(sizeof(cl_int) * (num_elements+1));
streams[0] =
@@ -111,7 +106,7 @@ test_work_group_all(cl_device_id device, cl_context context, cl_command_queue qu
p = input_ptr[0];
d = init_genrand( gRandomSeed );
for (size_t i = 0; i < (num_elements + 1); i++)
for (int i = 0; i < (num_elements + 1); i++)
{
p[i] = get_random_float((float)(-100000.f * M_PI), (float)(100000.f * M_PI) ,d);
}
@@ -133,7 +128,7 @@ test_work_group_all(cl_device_id device, cl_context context, cl_command_queue qu
}
// Line below is troublesome...
threads[0] = (size_t)n_elems;
threads[0] = (size_t)num_elements;
err = clEnqueueNDRangeKernel( queue, kernel, 1, NULL, threads, wg_size, 0, NULL, NULL );
if (err != CL_SUCCESS)
{

15
test_conformance/workgroups/test_wg_any.cpp
View File

@@ -20,8 +20,7 @@
#include <sys/types.h>
#include <sys/stat.h>
#include "procs.h"
#include "testBase.h"
const char *wg_any_kernel_code =
"__kernel void test_wg_any(global float *input, global int *output)\n"
@@ -65,8 +64,7 @@ verify_wg_any(float *inptr, int *outptr, size_t n, size_t wg_size)
return 0;
}
int
test_work_group_any(cl_device_id device, cl_context context, cl_command_queue queue, int n_elems)
REGISTER_TEST_VERSION(work_group_any, Version(2, 0))
{
cl_mem streams[2];
cl_float *input_ptr[1], *p;
@@ -74,8 +72,7 @@ test_work_group_any(cl_device_id device, cl_context context, cl_command_queue qu
cl_program program;
cl_kernel kernel;
size_t threads[1];
size_t wg_size[1];
size_t num_elements;
size_t wg_size[1];
int err;
MTdata d;
@@ -88,8 +85,6 @@ test_work_group_any(cl_device_id device, cl_context context, cl_command_queue qu
err = get_max_allowed_1d_work_group_size_on_device(device, kernel, wg_size);
test_error(err, "get_max_allowed_1d_work_group_size_on_device failed");
num_elements = n_elems;
input_ptr[0] = (cl_float*)malloc(sizeof(cl_float) * (num_elements+1));
output_ptr = (cl_int*)malloc(sizeof(cl_int) * (num_elements+1));
streams[0] =
@@ -111,7 +106,7 @@ test_work_group_any(cl_device_id device, cl_context context, cl_command_queue qu
p = input_ptr[0];
d = init_genrand( gRandomSeed );
for (size_t i = 0; i < (num_elements + 1); i++)
for (int i = 0; i < (num_elements + 1); i++)
{
p[i] = get_random_float((float)(-100000.f * M_PI), (float)(100000.f * M_PI) ,d);
}
@@ -133,7 +128,7 @@ test_work_group_any(cl_device_id device, cl_context context, cl_command_queue qu
}
// Line below is troublesome...
threads[0] = (size_t)n_elems;
threads[0] = (size_t)num_elements;
err = clEnqueueNDRangeKernel( queue, kernel, 1, NULL, threads, wg_size, 0, NULL, NULL );
if (err != CL_SUCCESS)
{

50
test_conformance/workgroups/test_wg_broadcast.cpp
View File

@@ -22,8 +22,7 @@
#include <algorithm>
#include "procs.h"
#include "testBase.h"
const char *wg_broadcast_1D_kernel_code =
"__kernel void test_wg_broadcast_1D(global float *input, global float *output)\n"
@@ -168,8 +167,7 @@ verify_wg_broadcast_3D(float *inptr, float *outptr, size_t nx, size_t ny, size_t
}
int
test_work_group_broadcast_1D(cl_device_id device, cl_context context, cl_command_queue queue, int n_elems)
REGISTER_TEST_VERSION(work_group_broadcast_1D, Version(2, 0))
{
cl_mem streams[2];
cl_float *input_ptr[1], *p;
@@ -177,8 +175,7 @@ test_work_group_broadcast_1D(cl_device_id device, cl_context context, cl_command
cl_program program;
cl_kernel kernel;
size_t globalsize[1];
size_t wg_size[1];
size_t num_elements;
size_t wg_size[1];
int err;
MTdata d;
@@ -192,8 +189,6 @@ test_work_group_broadcast_1D(cl_device_id device, cl_context context, cl_command
err = get_max_allowed_1d_work_group_size_on_device(device, kernel, wg_size);
test_error(err, "get_max_allowed_1d_work_group_size_on_device failed");
num_elements = n_elems;
input_ptr[0] = (cl_float*)malloc(sizeof(cl_float) * num_elements);
output_ptr = (cl_float*)malloc(sizeof(cl_float) * num_elements);
streams[0] = clCreateBuffer(context, CL_MEM_READ_WRITE,
@@ -214,7 +209,7 @@ test_work_group_broadcast_1D(cl_device_id device, cl_context context, cl_command
p = input_ptr[0];
d = init_genrand( gRandomSeed );
for (size_t i = 0; i < num_elements; i++)
for (int i = 0; i < num_elements; i++)
{
p[i] = get_random_float((float)(-100000.f * M_PI), (float)(100000.f * M_PI) ,d);
}
@@ -236,7 +231,7 @@ test_work_group_broadcast_1D(cl_device_id device, cl_context context, cl_command
}
// Line below is troublesome...
globalsize[0] = (size_t)n_elems;
globalsize[0] = (size_t)num_elements;
err = clEnqueueNDRangeKernel( queue, kernel, 1, NULL, globalsize, wg_size, 0, NULL, NULL );
if (err != CL_SUCCESS)
{
@@ -271,8 +266,7 @@ test_work_group_broadcast_1D(cl_device_id device, cl_context context, cl_command
}
int
test_work_group_broadcast_2D(cl_device_id device, cl_context context, cl_command_queue queue, int n_elems)
REGISTER_TEST_VERSION(work_group_broadcast_2D, Version(2, 0))
{
cl_mem streams[2];
cl_float *input_ptr[1], *p;
@@ -282,8 +276,7 @@ test_work_group_broadcast_2D(cl_device_id device, cl_context context, cl_command
size_t globalsize[2];
size_t localsize[2];
size_t wg_size[1];
size_t num_workgroups;
size_t num_elements;
size_t num_workgroups;
int err;
MTdata d;
@@ -314,7 +307,7 @@ test_work_group_broadcast_2D(cl_device_id device, cl_context context, cl_command
localsize[0] = localsize[1] = 1;
}
num_workgroups = std::max(n_elems / wg_size[0], (size_t)16);
num_workgroups = std::max(num_elements / wg_size[0], (size_t)16);
globalsize[0] = num_workgroups * localsize[0];
globalsize[1] = num_workgroups * localsize[1];
num_elements = globalsize[0] * globalsize[1];
@@ -339,7 +332,7 @@ test_work_group_broadcast_2D(cl_device_id device, cl_context context, cl_command
p = input_ptr[0];
d = init_genrand( gRandomSeed );
for (size_t i = 0; i < num_elements; i++)
for (int i = 0; i < num_elements; i++)
{
p[i] = get_random_float((float)(-100000.f * M_PI), (float)(100000.f * M_PI) ,d);
}
@@ -394,8 +387,7 @@ test_work_group_broadcast_2D(cl_device_id device, cl_context context, cl_command
}
int
test_work_group_broadcast_3D(cl_device_id device, cl_context context, cl_command_queue queue, int n_elems)
REGISTER_TEST_VERSION(work_group_broadcast_3D, Version(2, 0))
{
cl_mem streams[2];
cl_float *input_ptr[1], *p;
@@ -405,8 +397,7 @@ test_work_group_broadcast_3D(cl_device_id device, cl_context context, cl_command
size_t globalsize[3];
size_t localsize[3];
size_t wg_size[1];
size_t num_workgroups;
size_t num_elements;
size_t num_workgroups;
int err;
MTdata d;
@@ -437,7 +428,7 @@ test_work_group_broadcast_3D(cl_device_id device, cl_context context, cl_command
localsize[0] = localsize[1] = localsize[2] = 1;
}
num_workgroups = std::max(n_elems / wg_size[0], (size_t)8);
num_workgroups = std::max(num_elements / wg_size[0], (size_t)8);
globalsize[0] = num_workgroups * localsize[0];
globalsize[1] = num_workgroups * localsize[1];
globalsize[2] = num_workgroups * localsize[2];
@@ -463,7 +454,7 @@ test_work_group_broadcast_3D(cl_device_id device, cl_context context, cl_command
p = input_ptr[0];
d = init_genrand( gRandomSeed );
for (size_t i = 0; i < num_elements; i++)
for (int i = 0; i < num_elements; i++)
{
p[i] = get_random_float((float)(-100000.f * M_PI), (float)(100000.f * M_PI) ,d);
}
@@ -516,18 +507,3 @@ test_work_group_broadcast_3D(cl_device_id device, cl_context context, cl_command
return err;
}
int
test_work_group_broadcast(cl_device_id device, cl_context context, cl_command_queue queue, int n_elems)
{
int err;
err = test_work_group_broadcast_1D(device, context, queue, n_elems);
if (err) return err;
err = test_work_group_broadcast_2D(device, context, queue, n_elems);
if (err) return err;
return err;
}

205
test_conformance/workgroups/test_wg_scan_reduce.cpp
View File

@@ -19,7 +19,7 @@
#include <limits>
#include <vector>
#include "procs.h"
#include "testBase.h"
static std::string make_kernel_string(const std::string &type,
const std::string &kernelName,
@@ -272,184 +272,181 @@ static int run_test(cl_device_id device, cl_context context,
return TEST_PASS;
}
int test_work_group_reduce_add(cl_device_id device, cl_context context,
cl_command_queue queue, int n_elems)
{
int result = TEST_PASS;
result |= run_test<Reduce<Add<cl_int>>>(device, context, queue, n_elems);
result |= run_test<Reduce<Add<cl_uint>>>(device, context, queue, n_elems);
if (gHasLong)
{
result |=
run_test<Reduce<Add<cl_long>>>(device, context, queue, n_elems);
result |=
run_test<Reduce<Add<cl_ulong>>>(device, context, queue, n_elems);
}
return result;
}
int test_work_group_reduce_max(cl_device_id device, cl_context context,
cl_command_queue queue, int n_elems)
{
int result = TEST_PASS;
result |= run_test<Reduce<Max<cl_int>>>(device, context, queue, n_elems);
result |= run_test<Reduce<Max<cl_uint>>>(device, context, queue, n_elems);
if (gHasLong)
{
result |=
run_test<Reduce<Max<cl_long>>>(device, context, queue, n_elems);
result |=
run_test<Reduce<Max<cl_ulong>>>(device, context, queue, n_elems);
}
return result;
}
int test_work_group_reduce_min(cl_device_id device, cl_context context,
cl_command_queue queue, int n_elems)
{
int result = TEST_PASS;
result |= run_test<Reduce<Min<cl_int>>>(device, context, queue, n_elems);
result |= run_test<Reduce<Min<cl_uint>>>(device, context, queue, n_elems);
if (gHasLong)
{
result |=
run_test<Reduce<Min<cl_long>>>(device, context, queue, n_elems);
result |=
run_test<Reduce<Min<cl_ulong>>>(device, context, queue, n_elems);
}
return result;
}
int test_work_group_scan_inclusive_add(cl_device_id device, cl_context context,
cl_command_queue queue, int n_elems)
REGISTER_TEST_VERSION(work_group_reduce_add, Version(2, 0))
{
int result = TEST_PASS;
result |=
run_test<ScanInclusive<Add<cl_int>>>(device, context, queue, n_elems);
run_test<Reduce<Add<cl_int>>>(device, context, queue, num_elements);
result |=
run_test<ScanInclusive<Add<cl_uint>>>(device, context, queue, n_elems);
run_test<Reduce<Add<cl_uint>>>(device, context, queue, num_elements);
if (gHasLong)
{
result |= run_test<Reduce<Add<cl_long>>>(device, context, queue,
num_elements);
result |= run_test<Reduce<Add<cl_ulong>>>(device, context, queue,
num_elements);
}
return result;
}
REGISTER_TEST_VERSION(work_group_reduce_max, Version(2, 0))
{
int result = TEST_PASS;
result |=
run_test<Reduce<Max<cl_int>>>(device, context, queue, num_elements);
result |=
run_test<Reduce<Max<cl_uint>>>(device, context, queue, num_elements);
if (gHasLong)
{
result |= run_test<Reduce<Max<cl_long>>>(device, context, queue,
num_elements);
result |= run_test<Reduce<Max<cl_ulong>>>(device, context, queue,
num_elements);
}
return result;
}
REGISTER_TEST_VERSION(work_group_reduce_min, Version(2, 0))
{
int result = TEST_PASS;
result |=
run_test<Reduce<Min<cl_int>>>(device, context, queue, num_elements);
result |=
run_test<Reduce<Min<cl_uint>>>(device, context, queue, num_elements);
if (gHasLong)
{
result |= run_test<Reduce<Min<cl_long>>>(device, context, queue,
num_elements);
result |= run_test<Reduce<Min<cl_ulong>>>(device, context, queue,
num_elements);
}
return result;
}
REGISTER_TEST_VERSION(work_group_scan_inclusive_add, Version(2, 0))
{
int result = TEST_PASS;
result |= run_test<ScanInclusive<Add<cl_int>>>(device, context, queue,
num_elements);
result |= run_test<ScanInclusive<Add<cl_uint>>>(device, context, queue,
num_elements);
if (gHasLong)
{
result |= run_test<ScanInclusive<Add<cl_long>>>(device, context, queue,
n_elems);
num_elements);
result |= run_test<ScanInclusive<Add<cl_ulong>>>(device, context, queue,
n_elems);
num_elements);
}
return result;
}
int test_work_group_scan_inclusive_max(cl_device_id device, cl_context context,
cl_command_queue queue, int n_elems)
REGISTER_TEST_VERSION(work_group_scan_inclusive_max, Version(2, 0))
{
int result = TEST_PASS;
result |=
run_test<ScanInclusive<Max<cl_int>>>(device, context, queue, n_elems);
result |=
run_test<ScanInclusive<Max<cl_uint>>>(device, context, queue, n_elems);
result |= run_test<ScanInclusive<Max<cl_int>>>(device, context, queue,
num_elements);
result |= run_test<ScanInclusive<Max<cl_uint>>>(device, context, queue,
num_elements);
if (gHasLong)
{
result |= run_test<ScanInclusive<Max<cl_long>>>(device, context, queue,
n_elems);
num_elements);
result |= run_test<ScanInclusive<Max<cl_ulong>>>(device, context, queue,
n_elems);
num_elements);
}
return result;
}
int test_work_group_scan_inclusive_min(cl_device_id device, cl_context context,
cl_command_queue queue, int n_elems)
REGISTER_TEST_VERSION(work_group_scan_inclusive_min, Version(2, 0))
{
int result = TEST_PASS;
result |=
run_test<ScanInclusive<Min<cl_int>>>(device, context, queue, n_elems);
result |=
run_test<ScanInclusive<Min<cl_uint>>>(device, context, queue, n_elems);
result |= run_test<ScanInclusive<Min<cl_int>>>(device, context, queue,
num_elements);
result |= run_test<ScanInclusive<Min<cl_uint>>>(device, context, queue,
num_elements);
if (gHasLong)
{
result |= run_test<ScanInclusive<Min<cl_long>>>(device, context, queue,
n_elems);
num_elements);
result |= run_test<ScanInclusive<Min<cl_ulong>>>(device, context, queue,
n_elems);
num_elements);
}
return result;
}
int test_work_group_scan_exclusive_add(cl_device_id device, cl_context context,
cl_command_queue queue, int n_elems)
REGISTER_TEST_VERSION(work_group_scan_exclusive_add, Version(2, 0))
{
int result = TEST_PASS;
result |=
run_test<ScanExclusive<Add<cl_int>>>(device, context, queue, n_elems);
result |=
run_test<ScanExclusive<Add<cl_uint>>>(device, context, queue, n_elems);
result |= run_test<ScanExclusive<Add<cl_int>>>(device, context, queue,
num_elements);
result |= run_test<ScanExclusive<Add<cl_uint>>>(device, context, queue,
num_elements);
if (gHasLong)
{
result |= run_test<ScanExclusive<Add<cl_long>>>(device, context, queue,
n_elems);
num_elements);
result |= run_test<ScanExclusive<Add<cl_ulong>>>(device, context, queue,
n_elems);
num_elements);
}
return result;
}
int test_work_group_scan_exclusive_max(cl_device_id device, cl_context context,
cl_command_queue queue, int n_elems)
REGISTER_TEST_VERSION(work_group_scan_exclusive_max, Version(2, 0))
{
int result = TEST_PASS;
result |=
run_test<ScanExclusive<Max<cl_int>>>(device, context, queue, n_elems);
result |=
run_test<ScanExclusive<Max<cl_uint>>>(device, context, queue, n_elems);
result |= run_test<ScanExclusive<Max<cl_int>>>(device, context, queue,
num_elements);
result |= run_test<ScanExclusive<Max<cl_uint>>>(device, context, queue,
num_elements);
if (gHasLong)
{
result |= run_test<ScanExclusive<Max<cl_long>>>(device, context, queue,
n_elems);
num_elements);
result |= run_test<ScanExclusive<Max<cl_ulong>>>(device, context, queue,
n_elems);
num_elements);
}
return result;
}
int test_work_group_scan_exclusive_min(cl_device_id device, cl_context context,
cl_command_queue queue, int n_elems)
REGISTER_TEST_VERSION(work_group_scan_exclusive_min, Version(2, 0))
{
int result = TEST_PASS;
result |=
run_test<ScanExclusive<Min<cl_int>>>(device, context, queue, n_elems);
result |=
run_test<ScanExclusive<Min<cl_uint>>>(device, context, queue, n_elems);
result |= run_test<ScanExclusive<Min<cl_int>>>(device, context, queue,
num_elements);
result |= run_test<ScanExclusive<Min<cl_uint>>>(device, context, queue,
num_elements);
if (gHasLong)
{
result |= run_test<ScanExclusive<Min<cl_long>>>(device, context, queue,
n_elems);
num_elements);
result |= run_test<ScanExclusive<Min<cl_ulong>>>(device, context, queue,
n_elems);
num_elements);
}
return result;