diff --git a/test_conformance/basic/test_global_linear_id.cpp b/test_conformance/basic/test_global_linear_id.cpp index 046d12a1bf..ccd8ce27c8 100644 --- a/test_conformance/basic/test_global_linear_id.cpp +++ b/test_conformance/basic/test_global_linear_id.cpp @@ -1,6 +1,6 @@ // // 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 @@ -22,105 +22,93 @@ #include #include "harness/rounding_mode.h" +#include +#include + #include "procs.h" -static const char *global_linear_id_2d_code = -"__kernel void test_global_linear_id_2d(global int *dst)\n" -"{\n" -" int tid_x = get_global_id(0);\n" -" int tid_y = get_global_id(1);\n" -"\n" -" int linear_id = tid_y * get_global_size(0) + tid_x;\n" -" int result = (linear_id == (int)get_global_linear_id()) ? 0x1 : 0x0;\n" -" dst[linear_id] = result;\n" -"}\n"; - -static const char *global_linear_id_1d_code = -"__kernel void test_global_linear_id_1d(global int *dst)\n" -"{\n" -" int tid_x = get_global_id(0);\n" -"\n" -" int result = (tid_x == (int)get_global_linear_id()) ? 0x1 : 0x0;\n" -" dst[tid_x] = result;\n" -"}\n"; - - -static int -verify_global_linear_id(int *result, int n) +namespace { +const char *global_linear_id_2d_code = R"( +__kernel void test_global_linear_id_2d(global int *dst) { - int i; - for (i=0; i &result, int n) +{ + if (std::any_of(result.begin(), result.begin() + n, + [](cl_int value) { return 0 == value; })) { - if (result[i] == 0) - { - log_error("get_global_linear_id failed\n"); - return -1; - } + log_error("get_global_linear_id failed\n"); + return TEST_FAIL; } log_info("get_global_linear_id passed\n"); - return 0; + return TEST_PASS; +} } - -int -test_global_linear_id(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements) +int test_global_linear_id(cl_device_id device, cl_context context, + cl_command_queue queue, int num_elements) { - cl_mem streams; - cl_program program[2]; - cl_kernel kernel[2]; - - int *output_ptr; - size_t threads[2]; - int err; - num_elements = (int)sqrt((float)num_elements); - int length = num_elements * num_elements; - - output_ptr = (int *)malloc(sizeof(int) * length); - - streams = clCreateBuffer(context, CL_MEM_READ_WRITE, length * sizeof(int), - NULL, &err); - test_error( err, "clCreateBuffer failed."); - - err = create_single_kernel_helper(context, &program[0], &kernel[0], 1, - &global_linear_id_1d_code, - "test_global_linear_id_1d"); - test_error( err, "create_single_kernel_helper failed"); - err = create_single_kernel_helper(context, &program[1], &kernel[1], 1, - &global_linear_id_2d_code, - "test_global_linear_id_2d"); - test_error( err, "create_single_kernel_helper failed"); - - err = clSetKernelArg(kernel[0], 0, sizeof streams, &streams); - test_error( err, "clSetKernelArgs failed."); - err = clSetKernelArg(kernel[1], 0, sizeof streams, &streams); - test_error( err, "clSetKernelArgs failed."); - - threads[0] = (size_t)num_elements; - threads[1] = (size_t)num_elements; - err = clEnqueueNDRangeKernel(queue, kernel[1], 2, NULL, threads, NULL, 0, NULL, NULL); - test_error( err, "clEnqueueNDRangeKernel failed."); - - err = clEnqueueReadBuffer(queue, streams, CL_TRUE, 0, length*sizeof(int), output_ptr, 0, NULL, NULL); - test_error( err, "clEnqueueReadBuffer failed."); - - err = verify_global_linear_id(output_ptr, length); - - threads[0] = (size_t)num_elements; - err = clEnqueueNDRangeKernel(queue, kernel[1], 1, NULL, threads, NULL, 0, NULL, NULL); - test_error( err, "clEnqueueNDRangeKernel failed."); - - err = clEnqueueReadBuffer(queue, streams, CL_TRUE, 0, num_elements*sizeof(int), output_ptr, 0, NULL, NULL); - test_error( err, "clEnqueueReadBuffer failed."); - - err = verify_global_linear_id(output_ptr, num_elements); - - // cleanup - clReleaseMemObject(streams); - clReleaseKernel(kernel[0]); - clReleaseKernel(kernel[1]); - clReleaseProgram(program[0]); - clReleaseProgram(program[1]); - free(output_ptr); + clProgramWrapper program[2]; + clKernelWrapper kernel[2]; + + const char *kernel_names[] = { "test_global_linear_id_1d", + "test_global_linear_id_2d" }; + const char *kernel_code[] = { global_linear_id_1d_code, + global_linear_id_2d_code }; + int err = CL_SUCCESS; + + num_elements = static_cast(sqrt(static_cast(num_elements))); + int length = 1; + size_t threads[] = { static_cast(num_elements), + static_cast(num_elements) }; + + for (int i = 0; i < ARRAY_SIZE(program) && !err; i++) + { + length *= num_elements; + + std::vector output(length); + + clMemWrapper streams = clCreateBuffer( + context, CL_MEM_READ_WRITE, length * sizeof(cl_int), nullptr, &err); + test_error(err, "clCreateBuffer failed."); + + err = create_single_kernel_helper(context, &program[i], &kernel[i], 1, + &kernel_code[i], kernel_names[i]); + test_error(err, "create_single_kernel_helper failed"); + + err = clSetKernelArg(kernel[i], 0, sizeof streams, &streams); + test_error(err, "clSetKernelArgs failed."); + + err = clEnqueueNDRangeKernel(queue, kernel[i], i + 1, nullptr, threads, + nullptr, 0, nullptr, nullptr); + test_error(err, "clEnqueueNDRangeKernel failed."); + + err = clEnqueueReadBuffer(queue, streams, CL_TRUE, 0, + length * sizeof(cl_int), output.data(), 0, + nullptr, nullptr); + test_error(err, "clEnqueueReadBuffer failed."); + + err = verify_global_linear_id(output, length); + } return err; }