Skip to content

Commit

Permalink
Use CTS type wrappers for test_global_linear_id. (KhronosGroup#1545)
Browse files Browse the repository at this point in the history 
Signed-off-by: John Kesapides <[email protected]>

---------

Signed-off-by: John Kesapides <[email protected]>
  • Loading branch information
@JohnKesapidesARM
JohnKesapidesARM authored Sep 24, 2024
1 parent e865c1e commit cd74e02
Showing 1 changed file with 79 additions and 91 deletions.
170 changes: 79 additions & 91 deletions test_conformance/basic/test_global_linear_id.cpp
View file
Original file line number Diff line number Diff line change
@@ -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
Expand All @@ -22,105 +22,93 @@
#include <sys/stat.h>
#include "harness/rounding_mode.h"

#include <algorithm>
#include <vector>

#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<n; i++)
int tid_x = get_global_id(0);
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 -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<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) };

for (int i = 0; i < ARRAY_SIZE(program) && !err; i++)
{
length *= num_elements;

std::vector<cl_int> 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;
}

0 comments on commit cd74e02

Please sign in to comment.