ahmed/OpenCL-CTS
1
0
Fork 0
mirror of https://github.com/KhronosGroup/OpenCL-CTS.git synced 2026-03-19 14:09:03 +00:00
Code Packages Projects Releases Wiki Activity
Files
2d124f61ebee42b694534cb3f1adb0a593c840a4
OpenCL-CTS/test_conformance/clcpp/device_queue/test_enqueue.hpp
Kedar Patil 2821bf1323 Initial open source release of OpenCL 2.2 CTS.
2017-05-16 18:44:33 +05:30

700 lines
23 KiB
C++
Raw Blame History

//
// 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_DEVICE_QUEUE_TEST_ENQUEUE_HPP
#define TEST_CONFORMANCE_CLCPP_DEVICE_QUEUE_TEST_ENQUEUE_HPP
#include <sstream>
#include <string>
#include <vector>
#include <algorithm>
// Common for all OpenCL C++ tests
#include "../common.hpp"
namespace test_enqueue {
struct test_options
{
int test;
};
struct output_type
{
cl_int enqueue_kernel1_success;
cl_int enqueue_kernel2_success;
cl_int enqueue_kernel3_success;
cl_int enqueue_marker_success;
cl_int event1_is_valid;
cl_int event2_is_valid;
cl_int event3_is_valid;
cl_int user_event1_is_valid;
cl_int user_event2_is_valid;
cl_int values[10000];
};
const std::string source_common = R"(
struct output_type
{
int enqueue_kernel1_success;
int enqueue_kernel2_success;
int enqueue_kernel3_success;
int enqueue_marker_success;
int event1_is_valid;
int event2_is_valid;
int event3_is_valid;
int user_event1_is_valid;
int user_event2_is_valid;
int values[10000];
};
)";
// -----------------------------------------------------------------------------------
// ------------- 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.test == 0)
{
s << R"(
kernel void test(queue_t queue, global struct output_type *output)
{
const ulong gid = get_global_id(0);
if (gid != 0)
return;
output->enqueue_kernel2_success = 1;
output->enqueue_kernel3_success = 1;
output->enqueue_marker_success = 1;
output->event2_is_valid = 1;
output->event3_is_valid = 1;
output->user_event1_is_valid = 1;
output->user_event2_is_valid = 1;
queue_t default_queue = get_default_queue();
ndrange_t ndrange1 = ndrange_1D(get_global_size(0));
clk_event_t event1;
int status1 = enqueue_kernel(default_queue, CLK_ENQUEUE_FLAGS_NO_WAIT, ndrange1, 0, NULL, &event1,
^{
const ulong gid = get_global_id(0);
output->values[gid] = 1;
});
output->enqueue_kernel1_success = status1 == CLK_SUCCESS;
output->event1_is_valid = is_valid_event(event1);
release_event(event1);
}
)";
}
else if (options.test == 1)
{
s << R"(
kernel void test(queue_t queue, global struct output_type *output)
{
const ulong gid = get_global_id(0);
if (gid != 0)
return;
output->enqueue_kernel3_success = 1;
output->enqueue_marker_success = 1;
output->event3_is_valid = 1;
output->user_event1_is_valid = 1;
output->user_event2_is_valid = 1;
queue_t default_queue = get_default_queue();
ndrange_t ndrange1 = ndrange_1D(get_global_size(0) / 2);
clk_event_t event1;
int status1 = enqueue_kernel(default_queue, CLK_ENQUEUE_FLAGS_WAIT_WORK_GROUP, ndrange1, 0, NULL, &event1,
^{
const ulong gid = get_global_id(0);
output->values[gid * 2] = 1;
});
output->enqueue_kernel1_success = status1 == CLK_SUCCESS;
output->event1_is_valid = is_valid_event(event1);
ndrange_t ndrange2 = ndrange_1D(1, get_global_size(0) / 2, 1);
clk_event_t event2;
int status2 = enqueue_kernel(queue, CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange2, 1, &event1, &event2,
^{
const ulong gid = get_global_id(0);
output->values[(gid - 1) * 2 + 1] = 1;
});
output->enqueue_kernel2_success = status2 == CLK_SUCCESS;
output->event2_is_valid = is_valid_event(event2);
release_event(event1);
release_event(event2);
}
)";
}
else if (options.test == 2)
{
s << R"(
kernel void test(queue_t queue, global struct output_type *output)
{
const ulong gid = get_global_id(0);
if (gid != 0)
return;
output->enqueue_marker_success = 1;
output->event3_is_valid = 1;
output->enqueue_kernel3_success = 1;
queue_t default_queue = get_default_queue();
clk_event_t user_event1 = create_user_event();
retain_event(user_event1);
output->user_event1_is_valid = is_valid_event(user_event1);
ndrange_t ndrange1 = ndrange_1D(get_global_size(0) / 2);
clk_event_t event1;
int status1 = enqueue_kernel(queue, CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange1, 1, &user_event1, &event1,
^{
const ulong gid = get_global_id(0);
output->values[gid * 2] = 1;
});
output->enqueue_kernel1_success = status1 == CLK_SUCCESS;
output->event1_is_valid = is_valid_event(event1);
release_event(user_event1);
clk_event_t user_event2 = create_user_event();
output->user_event2_is_valid = is_valid_event(user_event2);
clk_event_t events[2];
events[0] = user_event2;
events[1] = user_event1;
ndrange_t ndrange2 = ndrange_1D(1, get_global_size(0) / 2, get_local_size(0));
clk_event_t event2;
int status2 = enqueue_kernel(default_queue, CLK_ENQUEUE_FLAGS_NO_WAIT, ndrange2, 2, events, &event2,
^(local void *p0, local void *p1, local void *p2) {
const ulong gid = get_global_id(0);
const ulong lid = get_local_id(0);
local int2 *l0 = (local int2 *)p0;
local int *l1 = (local int *)p1;
local int *l2 = (local int *)p2;
l1[get_local_size(0) - lid - 1] = gid > 0 ? 1 : 0;
work_group_barrier(CLK_LOCAL_MEM_FENCE);
if (lid < 5) l0[lid] = (int2)(3, 4);
if (lid < 3) l2[lid] = 5;
work_group_barrier(CLK_LOCAL_MEM_FENCE);
output->values[(gid - 1) * 2 + 1] = min(l1[lid], min(l0[0].x, l2[0]));
}, sizeof(int2) * 5, sizeof(int) * get_local_size(0), sizeof(int) * 3);
output->enqueue_kernel2_success = status2 == CLK_SUCCESS;
output->event2_is_valid = is_valid_event(event2);
set_user_event_status(user_event1, CL_COMPLETE);
set_user_event_status(user_event2, CL_COMPLETE);
release_event(user_event1);
release_event(user_event2);
release_event(event1);
release_event(event2);
}
)";
}
else if (options.test == 3)
{
s << R"(
kernel void test(queue_t queue, global struct output_type *output)
{
const ulong gid = get_global_id(0);
if (gid != 0)
return;
output->user_event2_is_valid = 1;
queue_t default_queue = get_default_queue();
ndrange_t ndrange1 = ndrange_1D(get_global_size(0) / 2);
clk_event_t event1;
int status1 = enqueue_kernel(default_queue, CLK_ENQUEUE_FLAGS_WAIT_WORK_GROUP, ndrange1, 0, NULL, &event1,
^{
const ulong gid = get_global_id(0);
output->values[gid * 2] = 20;
});
output->enqueue_kernel1_success = status1 == CLK_SUCCESS;
output->event1_is_valid = is_valid_event(event1);
ndrange_t ndrange2 = ndrange_1D(1, get_global_size(0) / 2, 1);
clk_event_t event2;
int status2 = enqueue_kernel(queue, CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange2, 0, NULL, &event2,
^{
const ulong gid = get_global_id(0);
output->values[(gid - 1) * 2 + 1] = 20;
});
output->enqueue_kernel2_success = status2 == CLK_SUCCESS;
output->event2_is_valid = is_valid_event(event2);
clk_event_t user_event1 = create_user_event();
output->user_event1_is_valid = is_valid_event(user_event1);
clk_event_t events[3];
events[0] = event2;
events[1] = user_event1;
events[2] = event1;
clk_event_t event3;
int status3 = enqueue_marker(queue, 3, events, &event3);
output->enqueue_marker_success = status3 == CLK_SUCCESS;
output->event3_is_valid = is_valid_event(event3);
int status4 = enqueue_kernel(default_queue, CLK_ENQUEUE_FLAGS_NO_WAIT, ndrange_1D(get_global_size(0)), 1, &event3, NULL,
^{
const ulong gid = get_global_id(0);
output->values[gid] /= 20;
});
output->enqueue_kernel3_success = status4 == CLK_SUCCESS;
set_user_event_status(user_event1, CL_COMPLETE);
release_event(user_event1);
release_event(event1);
release_event(event2);
release_event(event3);
}
)";
}
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_device_queue>
using namespace cl;
)";
s << source_common;
if (options.test == 0)
{
s << R"(
kernel void test(device_queue queue, global<output_type> *output)
{
const ulong gid = get_global_id(0);
if (gid != 0)
return;
output->enqueue_kernel2_success = 1;
output->enqueue_kernel3_success = 1;
output->enqueue_marker_success = 1;
output->event2_is_valid = 1;
output->event3_is_valid = 1;
output->user_event1_is_valid = 1;
output->user_event2_is_valid = 1;
device_queue default_queue = get_default_device_queue();
ndrange ndrange1(get_global_size(0));
event event1;
enqueue_status status1 = default_queue.enqueue_kernel(enqueue_policy::no_wait, 0, nullptr, &event1, ndrange1,
[](global<output_type> *output) {
const ulong gid = get_global_id(0);
output->values[gid] = 1;
}, output);
output->enqueue_kernel1_success = status1 == enqueue_status::success;
output->event1_is_valid = event1.is_valid();
event1.release();
}
)";
}
else if (options.test == 1)
{
s << R"(
kernel void test(device_queue queue, global<output_type> *output)
{
const ulong gid = get_global_id(0);
if (gid != 0)
return;
output->enqueue_kernel3_success = 1;
output->enqueue_marker_success = 1;
output->event3_is_valid = 1;
output->user_event1_is_valid = 1;
output->user_event2_is_valid = 1;
device_queue default_queue = get_default_device_queue();
ndrange ndrange1(get_global_size(0) / 2);
event event1;
enqueue_status status1 = default_queue.enqueue_kernel(enqueue_policy::wait_work_group, 0, nullptr, &event1, ndrange1,
[](global<output_type> *output) {
const ulong gid = get_global_id(0);
output->values[gid * 2] = 1;
}, output);
output->enqueue_kernel1_success = status1 == enqueue_status::success;
output->event1_is_valid = event1.is_valid();
ndrange ndrange2(1, get_global_size(0) / 2, 1);
event event2;
enqueue_status status2 = queue.enqueue_kernel(enqueue_policy::wait_kernel, 1, &event1, &event2, ndrange2,
[](global<output_type> *output) {
const ulong gid = get_global_id(0);
output->values[(gid - 1) * 2 + 1] = 1;
}, output);
output->enqueue_kernel2_success = status2 == enqueue_status::success;
output->event2_is_valid = event2.is_valid();
event1.release();
event2.release();
}
)";
}
else if (options.test == 2)
{
s << R"(
kernel void test(device_queue queue, global<output_type> *output)
{
const ulong gid = get_global_id(0);
if (gid != 0)
return;
output->enqueue_marker_success = 1;
output->event3_is_valid = 1;
output->enqueue_kernel3_success = 1;
device_queue default_queue = get_default_device_queue();
event user_event1 = make_user_event();
user_event1.retain();
output->user_event1_is_valid = user_event1.is_valid();
ndrange ndrange1(get_global_size(0) / 2);
event event1;
enqueue_status status1 = queue.enqueue_kernel(enqueue_policy::wait_kernel, 1, &user_event1, &event1, ndrange1,
[](global<output_type> *output){
const ulong gid = get_global_id(0);
output->values[gid * 2] = 1;
}, output);
output->enqueue_kernel1_success = status1 == enqueue_status::success;
output->event1_is_valid = event1.is_valid();
user_event1.release();
event user_event2 = make_user_event();
output->user_event2_is_valid = user_event2.is_valid();
event events[2];
events[0] = user_event2;
events[1] = user_event1;
ndrange ndrange2(1, get_global_size(0) / 2, get_local_size(0));
event event2;
enqueue_status status2 = default_queue.enqueue_kernel(enqueue_policy::no_wait, 2, events, &event2, ndrange2,
[](global<output_type> *output, local_ptr<int2[]> l0, local_ptr<int[]> l1, local_ptr<int[]> l2) {
const ulong gid = get_global_id(0);
const ulong lid = get_local_id(0);
l1[get_local_size(0) - lid - 1] = gid > 0 ? 1 : 0;
work_group_barrier(mem_fence::local);
if (lid < 5) l0[lid] = int2(3, 4);
if (lid < 3) l2[lid] = 5;
work_group_barrier(mem_fence::local);
output->values[(gid - 1) * 2 + 1] = min(l1[lid], min(l0[0].x, l2[0]));
}, output, local_ptr<int2[]>::size_type(5), local_ptr<int[]>::size_type(get_local_size(0)), local_ptr<int[]>::size_type(3));
output->enqueue_kernel2_success = status2 == enqueue_status::success;
output->event2_is_valid = event2.is_valid();
user_event1.set_status(event_status::complete);
user_event2.set_status(event_status::complete);
user_event1.release();
user_event2.release();
event1.release();
event2.release();
}
)";
}
else if (options.test == 3)
{
s << R"(
kernel void test(device_queue queue, global<output_type> *output)
{
const ulong gid = get_global_id(0);
if (gid != 0)
return;
output->user_event2_is_valid = 1;
device_queue default_queue = get_default_device_queue();
ndrange ndrange1(get_global_size(0) / 2);
event event1;
enqueue_status status1 = default_queue.enqueue_kernel(enqueue_policy::wait_work_group, 0, nullptr, &event1, ndrange1,
[](global<output_type> *output) {
const ulong gid = get_global_id(0);
output->values[gid * 2] = 20;
}, output);
output->enqueue_kernel1_success = status1 == enqueue_status::success;
output->event1_is_valid = event1.is_valid();
ndrange ndrange2(1, get_global_size(0) / 2, 1);
event event2;
enqueue_status status2 = queue.enqueue_kernel(enqueue_policy::wait_kernel, 0, nullptr, &event2, ndrange2,
[](global<output_type> *output) {
const ulong gid = get_global_id(0);
output->values[(gid - 1) * 2 + 1] = 20;
}, output);
output->enqueue_kernel2_success = status2 == enqueue_status::success;
output->event2_is_valid = event2.is_valid();
event user_event1 = make_user_event();
output->user_event1_is_valid = user_event1.is_valid();
event events[3];
events[0] = event2;
events[1] = user_event1;
events[2] = event1;
event event3;
enqueue_status status3 = queue.enqueue_marker(3, events, &event3);
output->enqueue_marker_success = status3 == enqueue_status::success;
output->event3_is_valid = event3.is_valid();
enqueue_status status4 = default_queue.enqueue_kernel(enqueue_policy::no_wait, 1, &event3, nullptr, ndrange(get_global_size(0)),
[](global<output_type> *output) {
const ulong gid = get_global_id(0);
output->values[gid] /= 20;
}, output);
output->enqueue_kernel3_success = status4 == enqueue_status::success;
user_event1.set_status(event_status::complete);
user_event1.release();
event1.release();
event2.release();
event3.release();
}
)";
}
return s.str();
}
#endif
int test(cl_device_id device, cl_context context, cl_command_queue queue, test_options options)
{
int error = CL_SUCCESS;
cl_program program;
cl_kernel kernel;
std::string kernel_name = "test";
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, &kernel,
source, 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, &kernel,
source, kernel_name, "-cl-std=CL2.0", false
);
RETURN_ON_ERROR(error)
// Normal run
#else
error = create_opencl_kernel(
context, &program, &kernel,
source, kernel_name
);
RETURN_ON_ERROR(error)
#endif
cl_uint max_queues;
error = clGetDeviceInfo(device, CL_DEVICE_MAX_ON_DEVICE_QUEUES, sizeof(cl_uint), &max_queues, NULL);
RETURN_ON_CL_ERROR(error, "clGetDeviceInfo")
cl_uint max_events;
error = clGetDeviceInfo(device, CL_DEVICE_MAX_ON_DEVICE_EVENTS, sizeof(cl_uint), &max_events, NULL);
RETURN_ON_CL_ERROR(error, "clGetDeviceInfo")
cl_command_queue device_queue1 = NULL;
cl_command_queue device_queue2 = NULL;
cl_queue_properties queue_properties1[] =
{
CL_QUEUE_PROPERTIES, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE | CL_QUEUE_ON_DEVICE | CL_QUEUE_ON_DEVICE_DEFAULT,
0
};
device_queue1 = clCreateCommandQueueWithProperties(context, device, queue_properties1, &error);
RETURN_ON_CL_ERROR(error, "clCreateCommandQueueWithProperties")
if (max_queues > 1)
{
cl_queue_properties queue_properties2[] =
{
CL_QUEUE_PROPERTIES, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE | CL_QUEUE_ON_DEVICE,
0
};
device_queue2 = clCreateCommandQueueWithProperties(context, device, queue_properties2, &error);
RETURN_ON_CL_ERROR(error, "clCreateCommandQueueWithProperties")
}
cl_mem output_buffer;
output_buffer = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(output_type), NULL, &error);
RETURN_ON_CL_ERROR(error, "clCreateBuffer")
error = clSetKernelArg(kernel, 0, sizeof(cl_command_queue), device_queue2 != NULL ? &device_queue2 : &device_queue1);
RETURN_ON_CL_ERROR(error, "clSetKernelArg")
error = clSetKernelArg(kernel, 1, sizeof(output_buffer), &output_buffer);
RETURN_ON_CL_ERROR(error, "clSetKernelArg")
const char pattern = 0;
error = clEnqueueFillBuffer(queue, output_buffer, &pattern, sizeof(pattern), 0, sizeof(output_type), 0, NULL, NULL);
RETURN_ON_CL_ERROR(error, "clEnqueueFillBuffer")
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 local_size = (std::min)((size_t)256, max_work_group_size);
const size_t global_size = 10000 / local_size * local_size;
const size_t count = global_size;
error = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &global_size, &local_size, 0, NULL, NULL);
RETURN_ON_CL_ERROR(error, "clEnqueueNDRangeKernel")
output_type output;
error = clEnqueueReadBuffer(
queue, output_buffer, CL_TRUE,
0, sizeof(output_type),
static_cast<void *>(&output),
0, NULL, NULL
);
RETURN_ON_CL_ERROR(error, "clEnqueueReadBuffer")
if (!output.enqueue_kernel1_success)
{
RETURN_ON_ERROR_MSG(-1, "enqueue_kernel did not succeed")
}
if (!output.enqueue_kernel2_success)
{
RETURN_ON_ERROR_MSG(-1, "enqueue_kernel did not succeed")
}
if (!output.enqueue_kernel3_success)
{
RETURN_ON_ERROR_MSG(-1, "enqueue_kernel did not succeed")
}
if (!output.enqueue_marker_success)
{
RETURN_ON_ERROR_MSG(-1, "enqueue_marker did not succeed")
}
if (!output.event1_is_valid)
{
RETURN_ON_ERROR_MSG(-1, "event1 is not valid")
}
if (!output.event2_is_valid)
{
RETURN_ON_ERROR_MSG(-1, "event2 is not valid")
}
if (!output.event3_is_valid)
{
RETURN_ON_ERROR_MSG(-1, "event3 is not valid")
}
if (!output.user_event1_is_valid)
{
RETURN_ON_ERROR_MSG(-1, "user_event1 is not valid")
}
if (!output.user_event2_is_valid)
{
RETURN_ON_ERROR_MSG(-1, "user_event2 is not valid")
}
for (size_t i = 0; i < count; i++)
{
const cl_int result = output.values[i];
const cl_int expected = 1;
if (result != expected)
{
RETURN_ON_ERROR_MSG(-1,
"kernel did not return correct value. Expected: %s, got: %s",
format_value(expected).c_str(), format_value(result).c_str()
)
}
}
clReleaseMemObject(output_buffer);
clReleaseCommandQueue(device_queue1);
if (device_queue2 != NULL)
clReleaseCommandQueue(device_queue2);
clReleaseKernel(kernel);
clReleaseProgram(program);
return error;
}
AUTO_TEST_CASE(test_enqueue_one_kernel)
(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
{
test_options options;
options.test = 0;
return test(device, context, queue, options);
}
AUTO_TEST_CASE(test_enqueue_two_kernels)
(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
{
test_options options;
options.test = 1;
return test(device, context, queue, options);
}
AUTO_TEST_CASE(test_enqueue_user_events_and_locals)
(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
{
test_options options;
options.test = 2;
return test(device, context, queue, options);
}
AUTO_TEST_CASE(test_enqueue_marker)
(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
{
test_options options;
options.test = 3;
return test(device, context, queue, options);
}
} // namespace
#endif // TEST_CONFORMANCE_CLCPP_DEVICE_QUEUE_TEST_ENQUEUE_HPP
View Git Blame Copy Permalink