[Piglit] [PATCH V3 1/1] cl: Add buffer-flags test
Tom Stellard
tom at stellard.net
Mon Dec 9 07:25:55 PST 2013
On Tue, Dec 03, 2013 at 02:58:18PM -0500, Jan Vesely wrote:
> v2: Rework to use subtests
> v3: fix printf output
> use piglit_merge_result
> whitespace fixes
> 0 as default was added in OpenCL 1.2
>
> Signed-off-by: Jan Vesely <jan.vesely at rutgers.edu>
> ---
>
> Hi,
>
> thanks for your feedback.
>
> I was wondering what to do about the overlap with cl-api-create-buffer test.
> cl-api-create-buffer includes similar tests but it does not test copying or
> destination buffers, it just creates a buffer and reads it back.
>
> So my question is whether it would make more sense to merge this test with
> api-cl-create-buffers (and make it use subtests), or have a separate test for
> testing data consistency (this one), and remove the overlapping tests from
> api-cl-create-buffers. Or leave both as they test slightly different things.
>
I went ahead and pushed this patch as is, if we decide to re-organize
the tests we can do it later as a separate series.
I'm not quite sure the best way to organize the tests. Maybe this will
become more clear as we add more tests, but right now I'm leaning towards
having one test program per API call and using piglit's subtest feature
for reporting results. The problem with the way the tests are currently
organized is if we're failing 4/5 of the api-cl-create-buffers tests
and then we regress to 3/5 this won't show up in the piglit results,
since 3/5 and 4/5 are both failures.
-Tom
> I ran this test on intel OCL (Corei7) CPU and nvidia CUDA (GF GT 630), both pass.
> Tests that use CL_MEM_USE_HOST_PTR in source buffer fail on mesa/clover
> on my AMD TURKS card.
>
> regards,
> Jan
>
>
> tests/all_cl.tests | 1 +
> tests/cl/custom/CMakeLists.cl.txt | 1 +
> tests/cl/custom/buffer-flags.c | 226 ++++++++++++++++++++++++++++++++++++++
> 3 files changed, 228 insertions(+)
> create mode 100644 tests/cl/custom/buffer-flags.c
>
> diff --git a/tests/all_cl.tests b/tests/all_cl.tests
> index a648e1a..57bfe3e 100644
> --- a/tests/all_cl.tests
> +++ b/tests/all_cl.tests
> @@ -40,6 +40,7 @@ profile.tests['Program'] = program
> add_plain_test(custom, 'Run simple kernel', ['cl-custom-run-simple-kernel'])
> add_plain_test(custom, 'Flush after enqueue kernel', ['cl-custom-flush-after-enqueue-kernel'])
> add_plain_test(custom, 'r600 create release buffer bug', ['cl-custom-r600-create-release-buffer-bug'])
> +add_plain_test(custom, 'Buffer flags', ['cl-custom-buffer-flags'])
>
> # API
> # Platform
> diff --git a/tests/cl/custom/CMakeLists.cl.txt b/tests/cl/custom/CMakeLists.cl.txt
> index 70649ea..778b845 100644
> --- a/tests/cl/custom/CMakeLists.cl.txt
> +++ b/tests/cl/custom/CMakeLists.cl.txt
> @@ -1,3 +1,4 @@
> piglit_cl_add_custom_test (run-simple-kernel run-simple-kernel.c)
> piglit_cl_add_custom_test (flush-after-enqueue-kernel flush-after-enqueue-kernel.c)
> piglit_cl_add_custom_test (r600-create-release-buffer-bug r600-create-release-buffer-bug.c)
> +piglit_cl_add_custom_test (buffer-flags buffer-flags.c)
> diff --git a/tests/cl/custom/buffer-flags.c b/tests/cl/custom/buffer-flags.c
> new file mode 100644
> index 0000000..321d75f
> --- /dev/null
> +++ b/tests/cl/custom/buffer-flags.c
> @@ -0,0 +1,226 @@
> +/*
> + * Copyright 2013 Jan Vesely
> + *
> + * Permission is hereby granted, free of charge, to any person obtaining a
> + * copy of this software and associated documentation files (the "Software"),
> + * to deal in the Software without restriction, including without limitation
> + * the rights to use, copy, modify, merge, publish, distribute, sublicense,
> + * and/or sell copies of the Software, and to permit persons to whom the
> + * Software is furnished to do so, subject to the following conditions:
> + *
> + * The above copyright notice and this permission notice (including the next
> + * paragraph) shall be included in all copies or substantial portions of the
> + * Software.
> + *
> + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
> + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
> + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL
> + * THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
> + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
> + * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER
> + * DEALINGS IN THE SOFTWARE.
> + *
> + */
> +
> +#include "piglit-framework-cl-custom.h"
> +
> +PIGLIT_CL_CUSTOM_TEST_CONFIG_BEGIN
> +
> + config.name = "CL buffer memory flags";
> + config.run_per_device = true;
> +
> +PIGLIT_CL_CUSTOM_TEST_CONFIG_END
> +
> +
> +/* This is a simple copy-kernel, the purpose of this test is to test buffer
> + * data availability, not specific compute function. */
> +char *source =
> +"__kernel void test (global float *out, global float *in) {\n"
> +" int i = get_global_id(0); \n"
> +" out[i] = in[i]; \n"
> +"} \n";
> +
> +#define BUFFER_SIZE 16 /* not too big */
> +
> +static enum piglit_result
> +buffer_test(piglit_cl_context *ctx,
> + cl_program *prg,
> + cl_mem_flags in_flags,
> + cl_mem_flags out_flags,
> + float data)
> +{
> + float in_data[BUFFER_SIZE];
> + float out_data[BUFFER_SIZE];
> + float *result = out_data;
> +
> + cl_mem in_buffer = NULL, out_buffer = NULL;
> + cl_kernel kernel = NULL;
> +
> + piglit_cl_context context = *ctx;
> +
> + cl_int errNo;
> + unsigned i;
> + size_t global = BUFFER_SIZE;
> + size_t local = 1;
> + enum piglit_result ret = PIGLIT_PASS;
> + const char kernel_name[] = "test";
> +
> + printf("> Running kernel test: in-0x%x-out-0x%x\n",
> + (unsigned)in_flags, (unsigned)out_flags);
> + for (i = 0; i < BUFFER_SIZE; ++i) {
> + in_data[i] = data;
> + out_data[i] = 0.0f;
> + }
> + printf("Using kernel %s\n", kernel_name);
> +
> + printf("Creating buffers...\n");
> + /* Create input buffer */
> + if ((in_flags & CL_MEM_USE_HOST_PTR) ||
> + (in_flags & CL_MEM_COPY_HOST_PTR)) {
> + /* Use host side memory */
> + in_buffer = clCreateBuffer(context->cl_ctx, in_flags,
> + sizeof(in_data), in_data, &errNo);
> + } else {
> + /* Use device memory and copy */
> + in_buffer = clCreateBuffer(context->cl_ctx, in_flags,
> + sizeof(in_data), NULL, &errNo);
> + if (errNo == CL_SUCCESS && !piglit_cl_write_buffer(
> + context->command_queues[0], in_buffer, 0,
> + sizeof(in_data), in_data)) {
> + ret = PIGLIT_FAIL;
> + goto cleanup;
> + }
> + }
> +
> + if(!piglit_cl_check_error(errNo, CL_SUCCESS)) {
> + fprintf(stderr,
> + "Could not create in buffer with flags %x: %s\n",
> + (unsigned)in_flags, piglit_cl_get_error_name(errNo));
> + ret = PIGLIT_FAIL;
> + goto cleanup;
> + }
> +
> + /* Create destination buffer */
> + if ((out_flags & CL_MEM_USE_HOST_PTR) ||
> + (out_flags & CL_MEM_COPY_HOST_PTR)) {
> + out_buffer = clCreateBuffer(context->cl_ctx, out_flags,
> + sizeof(out_data), out_data, &errNo);
> + } else {
> + out_buffer = clCreateBuffer(context->cl_ctx, out_flags,
> + sizeof(out_data), NULL, &errNo);
> + }
> + if(!piglit_cl_check_error(errNo, CL_SUCCESS)) {
> + fprintf(stderr,
> + "Could not create out buffer with flags %x: %s\n",
> + (unsigned)out_flags, piglit_cl_get_error_name(errNo));
> + ret = PIGLIT_FAIL;
> + goto cleanup;
> + }
> + kernel = piglit_cl_create_kernel(*prg, kernel_name);
> +
> + printf("Setting kernel arguments...\n");
> + if (!piglit_cl_set_kernel_arg(kernel, 0, sizeof(cl_mem), &out_buffer)) {
> + ret = PIGLIT_FAIL;
> + goto cleanup;
> + }
> + if (!piglit_cl_set_kernel_arg(kernel, 1, sizeof(cl_mem), &in_buffer)) {
> + ret = PIGLIT_FAIL;
> + goto cleanup;
> + }
> +
> + printf("Running the kernel...\n");
> + if (!piglit_cl_enqueue_ND_range_kernel(context->command_queues[0],
> + kernel, 1, &global, &local)) {
> + ret = PIGLIT_FAIL;
> + goto cleanup;
> + }
> +
> + clFlush(context->command_queues[0]);
> +
> + printf("Retrieving results...\n");
> + if ((out_flags & CL_MEM_USE_HOST_PTR) ||
> + (out_flags & CL_MEM_ALLOC_HOST_PTR)) {
> + /* buffer uses host side memory, map it here,
> + * map is also a synchronization point */
> + result = clEnqueueMapBuffer(context->command_queues[0],
> + out_buffer, true, CL_MAP_READ, 0, sizeof(out_data), 0,
> + NULL, NULL, &errNo);
> + if (!piglit_cl_check_error(errNo, CL_SUCCESS)) {
> + fprintf(stderr,
> + "Could not map out buffer with flags %x: %s\n",
> + (unsigned)out_flags,
> + piglit_cl_get_error_name(errNo));
> + ret = PIGLIT_FAIL;
> + goto cleanup;
> + }
> + } else {
> + /* Copy back from device */
> + if (!piglit_cl_read_buffer(context->command_queues[0],
> + out_buffer, 0, sizeof(out_data),
> + out_data)) {
> + ret = PIGLIT_FAIL;
> + goto cleanup;
> + }
> + }
> +
> + for (i = 0; i < BUFFER_SIZE; ++i) {
> + if (!piglit_cl_probe_floating(result[i], in_data[i], 0)) {
> + printf("Error at float[%u]\n", i);
> + ret = PIGLIT_FAIL;
> + goto cleanup;
> + }
> + }
> +
> + /* cleanup */
> +cleanup:
> + clReleaseMemObject(in_buffer);
> + clReleaseMemObject(out_buffer);
> + clReleaseKernel(kernel);
> + piglit_report_subtest_result(ret, "in-0x%x-out-0x%x",
> + (unsigned)in_flags, (unsigned)out_flags);
> + return ret;
> +
> +};
> +
> +enum piglit_result
> +piglit_cl_test(const int argc,
> + const char **argv,
> + const struct piglit_cl_custom_test_config *config,
> + const struct piglit_cl_custom_test_env *env)
> +{
> +
> + piglit_cl_context context = NULL;
> + cl_program program = NULL;
> +
> + unsigned i, j;
> +
> + static const cl_mem_flags possibilities[] = {
> +#ifdef CL_VERSION_1_2
> + 0,
> +#endif
> + CL_MEM_USE_HOST_PTR,
> + CL_MEM_COPY_HOST_PTR,
> + CL_MEM_ALLOC_HOST_PTR,
> + CL_MEM_COPY_HOST_PTR | CL_MEM_ALLOC_HOST_PTR,
> + };
> +
> + const size_t nump = ARRAY_SIZE(possibilities);
> + enum piglit_result part_ret, ret = PIGLIT_PASS;
> + float data = 10;
> +
> + context = piglit_cl_create_context(env->platform_id, &env->device_id, 1);
> +
> + program = piglit_cl_build_program_with_source(context, 1, &source, NULL);
> +
> + for (i = 0; i < nump; ++i)
> + for (j = 0; j < nump; ++j) {
> + part_ret = buffer_test(&context, &program,
> + possibilities[i], possibilities[j], ++data);
> + piglit_merge_result(&ret, part_ret);
> + }
> +
> +out:
> + clReleaseProgram(program);
> + piglit_cl_release_context(context);
> + return ret;
> +}
> --
> 1.8.3.1
>
More information about the Piglit
mailing list