633 lines
21 KiB
C++
633 lines
21 KiB
C++
//
|
|
// Copyright (c) 2017 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.
|
|
//
|
|
#ifndef TEST_CONFORMANCE_CLCPP_PIPES_TEST_PIPES_HPP
|
|
#define TEST_CONFORMANCE_CLCPP_PIPES_TEST_PIPES_HPP
|
|
|
|
#include <sstream>
|
|
#include <string>
|
|
#include <tuple>
|
|
#include <vector>
|
|
#include <algorithm>
|
|
|
|
// Common for all OpenCL C++ tests
|
|
#include "../common.hpp"
|
|
|
|
|
|
namespace test_pipes {
|
|
|
|
enum class pipe_source
|
|
{
|
|
param,
|
|
storage
|
|
};
|
|
|
|
enum class pipe_operation
|
|
{
|
|
work_item,
|
|
work_item_reservation,
|
|
work_group_reservation,
|
|
sub_group_reservation
|
|
};
|
|
|
|
struct test_options
|
|
{
|
|
pipe_operation operation;
|
|
pipe_source source;
|
|
int max_packets;
|
|
int num_packets;
|
|
};
|
|
|
|
struct output_type
|
|
{
|
|
cl_uint write_reservation_is_valid;
|
|
cl_uint write_success;
|
|
|
|
cl_uint num_packets;
|
|
cl_uint max_packets;
|
|
cl_uint read_reservation_is_valid;
|
|
cl_uint read_success;
|
|
|
|
cl_uint value;
|
|
};
|
|
|
|
const std::string source_common = R"(
|
|
struct output_type
|
|
{
|
|
uint write_reservation_is_valid;
|
|
uint write_success;
|
|
|
|
uint num_packets;
|
|
uint max_packets;
|
|
uint read_reservation_is_valid;
|
|
uint read_success;
|
|
|
|
uint value;
|
|
};
|
|
)";
|
|
|
|
// -----------------------------------------------------------------------------------
|
|
// ------------- ONLY FOR OPENCL 22 CONFORMANCE TEST 22 DEVELOPMENT ------------------
|
|
// -----------------------------------------------------------------------------------
|
|
#if defined(DEVELOPMENT) && defined(USE_OPENCLC_KERNELS)
|
|
std::string generate_source(test_options options)
|
|
{
|
|
std::stringstream s;
|
|
s << source_common;
|
|
if (options.operation == pipe_operation::work_item)
|
|
{
|
|
s << R"(
|
|
kernel void producer(write_only pipe uint out_pipe, global struct output_type *output)
|
|
{
|
|
const ulong gid = get_global_id(0);
|
|
|
|
output[gid].write_reservation_is_valid = 1;
|
|
|
|
uint value = gid;
|
|
output[gid].write_success = write_pipe(out_pipe, &value) == 0;
|
|
}
|
|
|
|
kernel void consumer(read_only pipe uint in_pipe, global struct output_type *output)
|
|
{
|
|
const ulong gid = get_global_id(0);
|
|
|
|
output[gid].num_packets = get_pipe_num_packets(in_pipe);
|
|
output[gid].max_packets = get_pipe_max_packets(in_pipe);
|
|
|
|
output[gid].read_reservation_is_valid = 1;
|
|
|
|
uint value;
|
|
output[gid].read_success = read_pipe(in_pipe, &value) == 0;
|
|
output[gid].value = value;
|
|
}
|
|
)";
|
|
}
|
|
else if (options.operation == pipe_operation::work_item_reservation)
|
|
{
|
|
s << R"(
|
|
kernel void producer(write_only pipe uint out_pipe, global struct output_type *output)
|
|
{
|
|
const ulong gid = get_global_id(0);
|
|
if (gid % 2 == 1) return;
|
|
|
|
reserve_id_t reservation = reserve_write_pipe(out_pipe, 2);
|
|
output[gid + 0].write_reservation_is_valid = is_valid_reserve_id(reservation);
|
|
output[gid + 1].write_reservation_is_valid = is_valid_reserve_id(reservation);
|
|
|
|
uint value0 = gid + 0;
|
|
uint value1 = gid + 1;
|
|
output[gid + 0].write_success = write_pipe(out_pipe, reservation, 0, &value0) == 0;
|
|
output[gid + 1].write_success = write_pipe(out_pipe, reservation, 1, &value1) == 0;
|
|
commit_write_pipe(out_pipe, reservation);
|
|
}
|
|
|
|
kernel void consumer(read_only pipe uint in_pipe, global struct output_type *output)
|
|
{
|
|
const ulong gid = get_global_id(0);
|
|
if (gid % 2 == 1) return;
|
|
|
|
output[gid + 0].num_packets = get_pipe_num_packets(in_pipe);
|
|
output[gid + 0].max_packets = get_pipe_max_packets(in_pipe);
|
|
output[gid + 1].num_packets = get_pipe_num_packets(in_pipe);
|
|
output[gid + 1].max_packets = get_pipe_max_packets(in_pipe);
|
|
|
|
reserve_id_t reservation = reserve_read_pipe(in_pipe, 2);
|
|
output[gid + 0].read_reservation_is_valid = is_valid_reserve_id(reservation);
|
|
output[gid + 1].read_reservation_is_valid = is_valid_reserve_id(reservation);
|
|
|
|
uint value0;
|
|
uint value1;
|
|
output[gid + 0].read_success = read_pipe(in_pipe, reservation, 1, &value0) == 0;
|
|
output[gid + 1].read_success = read_pipe(in_pipe, reservation, 0, &value1) == 0;
|
|
commit_read_pipe(in_pipe, reservation);
|
|
output[gid + 0].value = value0;
|
|
output[gid + 1].value = value1;
|
|
}
|
|
)";
|
|
}
|
|
else if (options.operation == pipe_operation::work_group_reservation)
|
|
{
|
|
s << R"(
|
|
kernel void producer(write_only pipe uint out_pipe, global struct output_type *output)
|
|
{
|
|
const ulong gid = get_global_id(0);
|
|
|
|
reserve_id_t reservation = work_group_reserve_write_pipe(out_pipe, get_local_size(0));
|
|
output[gid].write_reservation_is_valid = is_valid_reserve_id(reservation);
|
|
|
|
uint value = gid;
|
|
output[gid].write_success = write_pipe(out_pipe, reservation, get_local_id(0), &value) == 0;
|
|
work_group_commit_write_pipe(out_pipe, reservation);
|
|
}
|
|
|
|
kernel void consumer(read_only pipe uint in_pipe, global struct output_type *output)
|
|
{
|
|
const ulong gid = get_global_id(0);
|
|
|
|
output[gid].num_packets = get_pipe_num_packets(in_pipe);
|
|
output[gid].max_packets = get_pipe_max_packets(in_pipe);
|
|
|
|
reserve_id_t reservation = work_group_reserve_read_pipe(in_pipe, get_local_size(0));
|
|
output[gid].read_reservation_is_valid = is_valid_reserve_id(reservation);
|
|
|
|
uint value;
|
|
output[gid].read_success = read_pipe(in_pipe, reservation, get_local_size(0) - 1 - get_local_id(0), &value) == 0;
|
|
work_group_commit_read_pipe(in_pipe, reservation);
|
|
output[gid].value = value;
|
|
}
|
|
)";
|
|
}
|
|
else if (options.operation == pipe_operation::sub_group_reservation)
|
|
{
|
|
s << R"(
|
|
#pragma OPENCL EXTENSION cl_khr_subgroups : enable
|
|
|
|
kernel void producer(write_only pipe uint out_pipe, global struct output_type *output)
|
|
{
|
|
const ulong gid = get_global_id(0);
|
|
|
|
reserve_id_t reservation = sub_group_reserve_write_pipe(out_pipe, get_sub_group_size());
|
|
output[gid].write_reservation_is_valid = is_valid_reserve_id(reservation);
|
|
|
|
uint value = gid;
|
|
output[gid].write_success = write_pipe(out_pipe, reservation, get_sub_group_local_id(), &value) == 0;
|
|
sub_group_commit_write_pipe(out_pipe, reservation);
|
|
}
|
|
|
|
kernel void consumer(read_only pipe uint in_pipe, global struct output_type *output)
|
|
{
|
|
const ulong gid = get_global_id(0);
|
|
|
|
output[gid].num_packets = get_pipe_num_packets(in_pipe);
|
|
output[gid].max_packets = get_pipe_max_packets(in_pipe);
|
|
|
|
reserve_id_t reservation = sub_group_reserve_read_pipe(in_pipe, get_sub_group_size());
|
|
output[gid].read_reservation_is_valid = is_valid_reserve_id(reservation);
|
|
|
|
uint value;
|
|
output[gid].read_success = read_pipe(in_pipe, reservation, get_sub_group_size() - 1 - get_sub_group_local_id(), &value) == 0;
|
|
sub_group_commit_read_pipe(in_pipe, reservation);
|
|
output[gid].value = value;
|
|
}
|
|
)";
|
|
}
|
|
|
|
return s.str();
|
|
}
|
|
#else
|
|
std::string generate_source(test_options options)
|
|
{
|
|
std::stringstream s;
|
|
s << R"(
|
|
#include <opencl_memory>
|
|
#include <opencl_common>
|
|
#include <opencl_work_item>
|
|
#include <opencl_synchronization>
|
|
#include <opencl_pipe>
|
|
using namespace cl;
|
|
)";
|
|
|
|
s << source_common;
|
|
|
|
std::string init_out_pipe;
|
|
std::string init_in_pipe;
|
|
if (options.source == pipe_source::param)
|
|
{
|
|
init_out_pipe = "auto out_pipe = pipe_param;";
|
|
init_in_pipe = "auto in_pipe = pipe_param;";
|
|
}
|
|
else if (options.source == pipe_source::storage)
|
|
{
|
|
s << "pipe_storage<uint, " << std::to_string(options.max_packets) << "> storage;";
|
|
init_out_pipe = "auto out_pipe = storage.get<pipe_access::write>();";
|
|
init_in_pipe = "auto in_pipe = make_pipe(storage);";
|
|
}
|
|
|
|
if (options.operation == pipe_operation::work_item)
|
|
{
|
|
s << R"(
|
|
kernel void producer(pipe<uint, pipe_access::write> pipe_param, global_ptr<output_type[]> output)
|
|
{
|
|
)" << init_out_pipe << R"(
|
|
const ulong gid = get_global_id(0);
|
|
|
|
output[gid].write_reservation_is_valid = 1;
|
|
|
|
uint value = gid;
|
|
output[gid].write_success = out_pipe.write(value);
|
|
}
|
|
|
|
kernel void consumer(pipe<uint, pipe_access::read> pipe_param, global_ptr<output_type[]> output)
|
|
{
|
|
)" << init_in_pipe << R"(
|
|
const ulong gid = get_global_id(0);
|
|
|
|
output[gid].num_packets = in_pipe.num_packets();
|
|
output[gid].max_packets = in_pipe.max_packets();
|
|
|
|
output[gid].read_reservation_is_valid = 1;
|
|
|
|
uint value;
|
|
output[gid].read_success = in_pipe.read(value);
|
|
output[gid].value = value;
|
|
}
|
|
)";
|
|
}
|
|
else if (options.operation == pipe_operation::work_item_reservation)
|
|
{
|
|
s << R"(
|
|
kernel void producer(pipe<uint, pipe_access::write> pipe_param, global_ptr<output_type[]> output)
|
|
{
|
|
)" << init_out_pipe << R"(
|
|
const ulong gid = get_global_id(0);
|
|
if (gid % 2 == 1) return;
|
|
|
|
auto reservation = out_pipe.reserve(2);
|
|
output[gid + 0].write_reservation_is_valid = reservation.is_valid();
|
|
output[gid + 1].write_reservation_is_valid = reservation.is_valid();
|
|
|
|
uint value0 = gid + 0;
|
|
uint value1 = gid + 1;
|
|
output[gid + 0].write_success = reservation.write(0, value0);
|
|
output[gid + 1].write_success = reservation.write(1, value1);
|
|
reservation.commit();
|
|
}
|
|
|
|
kernel void consumer(pipe<uint, pipe_access::read> pipe_param, global_ptr<output_type[]> output)
|
|
{
|
|
)" << init_in_pipe << R"(
|
|
const ulong gid = get_global_id(0);
|
|
if (gid % 2 == 1) return;
|
|
|
|
output[gid + 0].num_packets = in_pipe.num_packets();
|
|
output[gid + 0].max_packets = in_pipe.max_packets();
|
|
output[gid + 1].num_packets = in_pipe.num_packets();
|
|
output[gid + 1].max_packets = in_pipe.max_packets();
|
|
|
|
auto reservation = in_pipe.reserve(2);
|
|
output[gid + 0].read_reservation_is_valid = reservation.is_valid();
|
|
output[gid + 1].read_reservation_is_valid = reservation.is_valid();
|
|
|
|
uint value0;
|
|
uint value1;
|
|
output[gid + 0].read_success = reservation.read(1, value0);
|
|
output[gid + 1].read_success = reservation.read(0, value1);
|
|
reservation.commit();
|
|
output[gid + 0].value = value0;
|
|
output[gid + 1].value = value1;
|
|
}
|
|
)";
|
|
}
|
|
else if (options.operation == pipe_operation::work_group_reservation)
|
|
{
|
|
s << R"(
|
|
kernel void producer(pipe<uint, pipe_access::write> pipe_param, global_ptr<output_type[]> output)
|
|
{
|
|
)" << init_out_pipe << R"(
|
|
const ulong gid = get_global_id(0);
|
|
|
|
auto reservation = out_pipe.work_group_reserve(get_local_size(0));
|
|
output[gid].write_reservation_is_valid = reservation.is_valid();
|
|
|
|
uint value = gid;
|
|
output[gid].write_success = reservation.write(get_local_id(0), value);
|
|
reservation.commit();
|
|
}
|
|
|
|
kernel void consumer(pipe<uint, pipe_access::read> pipe_param, global_ptr<output_type[]> output)
|
|
{
|
|
)" << init_in_pipe << R"(
|
|
const ulong gid = get_global_id(0);
|
|
|
|
output[gid].num_packets = in_pipe.num_packets();
|
|
output[gid].max_packets = in_pipe.max_packets();
|
|
|
|
auto reservation = in_pipe.work_group_reserve(get_local_size(0));
|
|
output[gid].read_reservation_is_valid = reservation.is_valid();
|
|
|
|
uint value;
|
|
output[gid].read_success = reservation.read(get_local_size(0) - 1 - get_local_id(0), value);
|
|
reservation.commit();
|
|
output[gid].value = value;
|
|
}
|
|
)";
|
|
}
|
|
else if (options.operation == pipe_operation::sub_group_reservation)
|
|
{
|
|
s << R"(
|
|
kernel void producer(pipe<uint, pipe_access::write> pipe_param, global_ptr<output_type[]> output)
|
|
{
|
|
)" << init_out_pipe << R"(
|
|
const ulong gid = get_global_id(0);
|
|
|
|
auto reservation = out_pipe.sub_group_reserve(get_sub_group_size());
|
|
output[gid].write_reservation_is_valid = reservation.is_valid();
|
|
|
|
uint value = gid;
|
|
output[gid].write_success = reservation.write(get_sub_group_local_id(), value);
|
|
reservation.commit();
|
|
}
|
|
|
|
kernel void consumer(pipe<uint, pipe_access::read> pipe_param, global_ptr<output_type[]> output)
|
|
{
|
|
)" << init_in_pipe << R"(
|
|
const ulong gid = get_global_id(0);
|
|
|
|
output[gid].num_packets = in_pipe.num_packets();
|
|
output[gid].max_packets = in_pipe.max_packets();
|
|
|
|
auto reservation = in_pipe.sub_group_reserve(get_sub_group_size());
|
|
output[gid].read_reservation_is_valid = reservation.is_valid();
|
|
|
|
uint value;
|
|
output[gid].read_success = reservation.read(get_sub_group_size() - 1 - get_sub_group_local_id(), value);
|
|
reservation.commit();
|
|
output[gid].value = value;
|
|
}
|
|
)";
|
|
}
|
|
|
|
return s.str();
|
|
}
|
|
#endif
|
|
|
|
int test(cl_device_id device, cl_context context, cl_command_queue queue, test_options options)
|
|
{
|
|
int error = CL_SUCCESS;
|
|
|
|
if (options.num_packets % 2 != 0 || options.max_packets < options.num_packets)
|
|
{
|
|
RETURN_ON_ERROR_MSG(-1, "Invalid test options")
|
|
}
|
|
|
|
#if defined(DEVELOPMENT) && defined(USE_OPENCLC_KERNELS)
|
|
if (options.operation == pipe_operation::sub_group_reservation && !is_extension_available(device, "cl_khr_subgroups"))
|
|
{
|
|
log_info("SKIPPED: Extension `cl_khr_subgroups` is not supported. Skipping tests.\n");
|
|
return CL_SUCCESS;
|
|
}
|
|
#endif
|
|
|
|
cl_program program;
|
|
cl_kernel producer_kernel;
|
|
cl_kernel consumer_kernel;
|
|
|
|
std::string producer_kernel_name = "producer";
|
|
std::string consumer_kernel_name = "consumer";
|
|
std::string source = generate_source(options);
|
|
|
|
// -----------------------------------------------------------------------------------
|
|
// ------------- ONLY FOR OPENCL 22 CONFORMANCE TEST 22 DEVELOPMENT ------------------
|
|
// -----------------------------------------------------------------------------------
|
|
// Only OpenCL C++ to SPIR-V compilation
|
|
#if defined(DEVELOPMENT) && defined(ONLY_SPIRV_COMPILATION)
|
|
error = create_opencl_kernel(
|
|
context, &program, &producer_kernel,
|
|
source, producer_kernel_name
|
|
);
|
|
RETURN_ON_ERROR(error)
|
|
return error;
|
|
// Use OpenCL C kernels instead of OpenCL C++ kernels (test C++ host code)
|
|
#elif defined(DEVELOPMENT) && defined(USE_OPENCLC_KERNELS)
|
|
error = create_opencl_kernel(
|
|
context, &program, &producer_kernel,
|
|
source, producer_kernel_name, "-cl-std=CL2.0", false
|
|
);
|
|
RETURN_ON_ERROR(error)
|
|
consumer_kernel = clCreateKernel(program, consumer_kernel_name.c_str(), &error);
|
|
RETURN_ON_CL_ERROR(error, "clCreateKernel")
|
|
// Normal run
|
|
#else
|
|
error = create_opencl_kernel(
|
|
context, &program, &producer_kernel,
|
|
source, producer_kernel_name
|
|
);
|
|
RETURN_ON_ERROR(error)
|
|
consumer_kernel = clCreateKernel(program, consumer_kernel_name.c_str(), &error);
|
|
RETURN_ON_CL_ERROR(error, "clCreateKernel")
|
|
#endif
|
|
|
|
size_t max_work_group_size;
|
|
error = clGetDeviceInfo(device, CL_DEVICE_MAX_WORK_GROUP_SIZE, sizeof(size_t), &max_work_group_size, NULL);
|
|
RETURN_ON_CL_ERROR(error, "clGetDeviceInfo")
|
|
|
|
const size_t count = options.num_packets;
|
|
const size_t local_size = (std::min)((size_t)256, max_work_group_size);
|
|
const size_t global_size = count;
|
|
|
|
const cl_uint packet_size = sizeof(cl_uint);
|
|
|
|
cl_mem pipe = clCreatePipe(context, 0, packet_size, options.max_packets, NULL, &error);
|
|
RETURN_ON_CL_ERROR(error, "clCreatePipe")
|
|
|
|
cl_mem output_buffer;
|
|
output_buffer = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(output_type) * count, NULL, &error);
|
|
RETURN_ON_CL_ERROR(error, "clCreateBuffer")
|
|
|
|
const char pattern = 0;
|
|
error = clEnqueueFillBuffer(queue, output_buffer, &pattern, sizeof(pattern), 0, sizeof(output_type) * count, 0, NULL, NULL);
|
|
RETURN_ON_CL_ERROR(error, "clEnqueueFillBuffer")
|
|
|
|
error = clSetKernelArg(producer_kernel, 0, sizeof(cl_mem), &pipe);
|
|
RETURN_ON_CL_ERROR(error, "clSetKernelArg")
|
|
error = clSetKernelArg(producer_kernel, 1, sizeof(output_buffer), &output_buffer);
|
|
RETURN_ON_CL_ERROR(error, "clSetKernelArg")
|
|
|
|
error = clEnqueueNDRangeKernel(queue, producer_kernel, 1, NULL, &global_size, NULL, 0, NULL, NULL);
|
|
RETURN_ON_CL_ERROR(error, "clEnqueueNDRangeKernel")
|
|
|
|
error = clSetKernelArg(consumer_kernel, 0, sizeof(cl_mem), &pipe);
|
|
RETURN_ON_CL_ERROR(error, "clSetKernelArg")
|
|
error = clSetKernelArg(consumer_kernel, 1, sizeof(output_buffer), &output_buffer);
|
|
RETURN_ON_CL_ERROR(error, "clSetKernelArg")
|
|
|
|
error = clEnqueueNDRangeKernel(queue, consumer_kernel, 1, NULL, &global_size, &local_size, 0, NULL, NULL);
|
|
RETURN_ON_CL_ERROR(error, "clEnqueueNDRangeKernel")
|
|
|
|
std::vector<output_type> output(count);
|
|
error = clEnqueueReadBuffer(
|
|
queue, output_buffer, CL_TRUE,
|
|
0, sizeof(output_type) * count,
|
|
static_cast<void *>(output.data()),
|
|
0, NULL, NULL
|
|
);
|
|
RETURN_ON_CL_ERROR(error, "clEnqueueReadBuffer")
|
|
|
|
std::vector<bool> existing_values(count, false);
|
|
for (size_t gid = 0; gid < count; gid++)
|
|
{
|
|
const output_type &o = output[gid];
|
|
|
|
if (!o.write_reservation_is_valid)
|
|
{
|
|
RETURN_ON_ERROR_MSG(-1, "write reservation is not valid")
|
|
}
|
|
if (!o.write_success)
|
|
{
|
|
RETURN_ON_ERROR_MSG(-1, "write did not succeed")
|
|
}
|
|
|
|
if (o.num_packets == 0 || o.num_packets > options.num_packets)
|
|
{
|
|
RETURN_ON_ERROR_MSG(-1, "num_packets did not return correct value")
|
|
}
|
|
if (o.max_packets != options.max_packets)
|
|
{
|
|
RETURN_ON_ERROR_MSG(-1, "max_packets did not return correct value")
|
|
}
|
|
if (!o.read_reservation_is_valid)
|
|
{
|
|
RETURN_ON_ERROR_MSG(-1, "read reservation is not valid")
|
|
}
|
|
if (!o.read_success)
|
|
{
|
|
RETURN_ON_ERROR_MSG(-1, "read did not succeed")
|
|
}
|
|
|
|
// Every value must be presented once in any order
|
|
if (o.value >= count || existing_values[o.value])
|
|
{
|
|
RETURN_ON_ERROR_MSG(-1, "kernel did not return correct value")
|
|
}
|
|
existing_values[o.value] = true;
|
|
}
|
|
|
|
clReleaseMemObject(pipe);
|
|
clReleaseMemObject(output_buffer);
|
|
clReleaseKernel(producer_kernel);
|
|
clReleaseKernel(consumer_kernel);
|
|
clReleaseProgram(program);
|
|
return error;
|
|
}
|
|
|
|
const pipe_operation pipe_operations[] = {
|
|
pipe_operation::work_item,
|
|
pipe_operation::work_item_reservation,
|
|
pipe_operation::work_group_reservation,
|
|
pipe_operation::sub_group_reservation
|
|
};
|
|
|
|
const std::tuple<int, int> max_and_num_packets[] = {
|
|
std::make_tuple<int, int>(2, 2),
|
|
std::make_tuple<int, int>(10, 8),
|
|
std::make_tuple<int, int>(256, 254),
|
|
std::make_tuple<int, int>(1 << 16, 1 << 16),
|
|
std::make_tuple<int, int>((1 << 16) + 5, 1 << 16),
|
|
std::make_tuple<int, int>(12345, 12344),
|
|
std::make_tuple<int, int>(1 << 18, 1 << 18)
|
|
};
|
|
|
|
AUTO_TEST_CASE(test_pipes_pipe)
|
|
(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
|
|
{
|
|
std::vector<std::tuple<int, int>> ps;
|
|
for (auto p : max_and_num_packets)
|
|
{
|
|
if (std::get<0>(p) < num_elements)
|
|
ps.push_back(p);
|
|
}
|
|
ps.push_back(std::tuple<int, int>(num_elements, num_elements));
|
|
|
|
int error = CL_SUCCESS;
|
|
|
|
for (auto operation : pipe_operations)
|
|
for (auto p : ps)
|
|
{
|
|
test_options options;
|
|
options.source = pipe_source::param;
|
|
options.max_packets = std::get<0>(p);
|
|
options.num_packets = std::get<1>(p);
|
|
options.operation = operation;
|
|
|
|
error = test(device, context, queue, options);
|
|
RETURN_ON_ERROR(error)
|
|
}
|
|
|
|
return error;
|
|
}
|
|
|
|
AUTO_TEST_CASE(test_pipes_pipe_storage)
|
|
(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
|
|
{
|
|
std::vector<std::tuple<int, int>> ps;
|
|
for (auto p : max_and_num_packets)
|
|
{
|
|
if (std::get<0>(p) < num_elements)
|
|
ps.push_back(p);
|
|
}
|
|
ps.push_back(std::tuple<int, int>(num_elements, num_elements));
|
|
|
|
int error = CL_SUCCESS;
|
|
|
|
for (auto operation : pipe_operations)
|
|
for (auto p : ps)
|
|
{
|
|
test_options options;
|
|
options.source = pipe_source::storage;
|
|
options.max_packets = std::get<0>(p);
|
|
options.num_packets = std::get<1>(p);
|
|
options.operation = operation;
|
|
|
|
error = test(device, context, queue, options);
|
|
RETURN_ON_ERROR(error)
|
|
}
|
|
|
|
return error;
|
|
}
|
|
|
|
} // namespace
|
|
|
|
#endif // TEST_CONFORMANCE_CLCPP_PIPES_TEST_PIPES_HPP
|