ahmed/OpenCL-CTS
1
0
Fork 0
mirror of https://github.com/KhronosGroup/OpenCL-CTS.git synced 2026-03-25 08:19:02 +00:00
Code Packages Projects Releases Wiki Activity

Use CTS type wrappers for test_global_linear_id. (#1545)

Browse Source 
Signed-off-by: John Kesapides <john.kesapides@arm.com>

---------

Signed-off-by: John Kesapides <john.kesapides@arm.com>
This commit is contained in:
John Kesapides
2024-09-24 18:02:32 +01:00
committed by GitHub
parent e865c1e8dd
commit cd74e02643
1 changed files with 69 additions and 81 deletions
Download Patch File Download Diff File Expand all files Collapse all files

150
test_conformance/basic/test_global_linear_id.cpp
View File

@@ -1,6 +1,6 @@
// //
// Copyright (c) 2017 The Khronos Group Inc. // Copyright (c) 2017 The Khronos Group Inc.
// //
// Licensed under the Apache License, Version 2.0 (the "License"); // Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License. // you may not use this file except in compliance with the License.
// You may obtain a copy of the License at // You may obtain a copy of the License at
@@ -22,105 +22,93 @@
#include <sys/stat.h> #include <sys/stat.h>
#include "harness/rounding_mode.h" #include "harness/rounding_mode.h"
#include <algorithm>
#include <vector>
#include "procs.h" #include "procs.h"
static const char *global_linear_id_2d_code = namespace {
"__kernel void test_global_linear_id_2d(global int *dst)\n" const char *global_linear_id_2d_code = R"(
"{\n" __kernel void test_global_linear_id_2d(global int *dst)
" 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)
{ {
int i; int tid_x = get_global_id(0);
for (i=0; i<n; i++) int tid_y = get_global_id(1);
int linear_id = tid_y * get_global_size(0) + tid_x;
int result = (linear_id == (int)get_global_linear_id()) ? 0x1 : 0x0;
dst[linear_id] = result;
}
)";
const char *global_linear_id_1d_code = R"(
__kernel void test_global_linear_id_1d(global int *dst)
{
int tid_x = get_global_id(0);
int result = (tid_x == (int)get_global_linear_id()) ? 0x1 : 0x0;
dst[tid_x] = result;
}
)";
int verify_global_linear_id(std::vector<cl_int> &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 TEST_FAIL;
log_error("get_global_linear_id failed\n");
return -1;
}
} }
log_info("get_global_linear_id passed\n"); 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,
int cl_command_queue queue, int num_elements)
test_global_linear_id(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
{ {
cl_mem streams; clProgramWrapper program[2];
cl_program program[2]; clKernelWrapper kernel[2];
cl_kernel kernel[2];
int *output_ptr; const char *kernel_names[] = { "test_global_linear_id_1d",
size_t threads[2]; "test_global_linear_id_2d" };
int err; const char *kernel_code[] = { global_linear_id_1d_code,
num_elements = (int)sqrt((float)num_elements); global_linear_id_2d_code };
int length = num_elements * num_elements; int err = CL_SUCCESS;
output_ptr = (int *)malloc(sizeof(int) * length); num_elements = static_cast<int>(sqrt(static_cast<float>(num_elements)));
int length = 1;
size_t threads[] = { static_cast<size_t>(num_elements),
static_cast<size_t>(num_elements) };
streams = clCreateBuffer(context, CL_MEM_READ_WRITE, length * sizeof(int), for (int i = 0; i < ARRAY_SIZE(program) && !err; i++)
NULL, &err); {
test_error( err, "clCreateBuffer failed."); length *= num_elements;
err = create_single_kernel_helper(context, &program[0], &kernel[0], 1, std::vector<cl_int> output(length);
&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); clMemWrapper streams = clCreateBuffer(
test_error( err, "clSetKernelArgs failed."); context, CL_MEM_READ_WRITE, length * sizeof(cl_int), nullptr, &err);
err = clSetKernelArg(kernel[1], 0, sizeof streams, &streams); test_error(err, "clCreateBuffer failed.");
test_error( err, "clSetKernelArgs failed.");
threads[0] = (size_t)num_elements; err = create_single_kernel_helper(context, &program[i], &kernel[i], 1,
threads[1] = (size_t)num_elements; &kernel_code[i], kernel_names[i]);
err = clEnqueueNDRangeKernel(queue, kernel[1], 2, NULL, threads, NULL, 0, NULL, NULL); test_error(err, "create_single_kernel_helper failed");
test_error( err, "clEnqueueNDRangeKernel failed.");
err = clEnqueueReadBuffer(queue, streams, CL_TRUE, 0, length*sizeof(int), output_ptr, 0, NULL, NULL); err = clSetKernelArg(kernel[i], 0, sizeof streams, &streams);
test_error( err, "clEnqueueReadBuffer failed."); test_error(err, "clSetKernelArgs failed.");
err = verify_global_linear_id(output_ptr, length); err = clEnqueueNDRangeKernel(queue, kernel[i], i + 1, nullptr, threads,
nullptr, 0, nullptr, nullptr);
test_error(err, "clEnqueueNDRangeKernel failed.");
threads[0] = (size_t)num_elements; err = clEnqueueReadBuffer(queue, streams, CL_TRUE, 0,
err = clEnqueueNDRangeKernel(queue, kernel[1], 1, NULL, threads, NULL, 0, NULL, NULL); length * sizeof(cl_int), output.data(), 0,
test_error( err, "clEnqueueNDRangeKernel failed."); nullptr, nullptr);
test_error(err, "clEnqueueReadBuffer failed.");
err = clEnqueueReadBuffer(queue, streams, CL_TRUE, 0, num_elements*sizeof(int), output_ptr, 0, NULL, NULL); err = verify_global_linear_id(output, length);
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);
return err; return err;
} }