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

Use CTS type wrappers for test_enqueued_local_size (#1544)

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
2022-10-14 09:55:10 +01:00
committed by GitHub
parent 5e116e7b0d
commit 90a5183ec4
1 changed files with 56 additions and 60 deletions
Download Patch File Download Diff File Expand all files Collapse all files

116
test_conformance/basic/test_enqueued_local_size.cpp
View File

@@ -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
@@ -26,32 +26,33 @@
#include "procs.h"
static const char *enqueued_local_size_2d_code =
"__kernel void test_enqueued_local_size_2d(global int *dst)\n"
"{\n"
" if ((get_global_id(0) == 0) && (get_global_id(1) == 0))\n"
" {\n"
" dst[0] = (int)get_enqueued_local_size(0)\n;"
" dst[1] = (int)get_enqueued_local_size(1)\n;"
" }\n"
"}\n";
static const char *enqueued_local_size_2d_code = R"(
__kernel void test_enqueued_local_size_2d(global int *dst)
{
if ((get_global_id(0) == 0) && (get_global_id(1) == 0))
{
dst[0] = (int)get_enqueued_local_size(0);
dst[1] = (int)get_enqueued_local_size(1);
}
}
)";
static const char *enqueued_local_size_1d_code =
"__kernel void test_enqueued_local_size_1d(global int *dst)\n"
"{\n"
" int tid_x = get_global_id(0);\n"
" if (get_global_id(0) == 0)\n"
" {\n"
" dst[tid_x] = (int)get_enqueued_local_size(0)\n;"
" }\n"
"}\n";
static const char *enqueued_local_size_1d_code = R"(
__kernel void test_enqueued_local_size_1d(global int *dst)
{
int tid_x = get_global_id(0);
if (get_global_id(0) == 0)
{
dst[tid_x] = (int)get_enqueued_local_size(0);
}
}
)";
static int
verify_enqueued_local_size(int *result, size_t *expected, int n)
static int verify_enqueued_local_size(int *result, size_t *expected, int n)
{
int i;
for (i=0; i<n; i++)
for (i = 0; i < n; i++)
{
if (result[i] != (int)expected[i])
{
@@ -64,14 +65,14 @@ verify_enqueued_local_size(int *result, size_t *expected, int n)
}
int
test_enqueued_local_size(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
int test_enqueued_local_size(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];
clMemWrapper stream;
clProgramWrapper program[2];
clKernelWrapper kernel[2];
int *output_ptr;
cl_int output_ptr[2];
size_t globalsize[2];
size_t localsize[2];
int err;
@@ -97,34 +98,33 @@ test_enqueued_local_size(cl_device_id device, cl_context context, cl_command_que
}
}
output_ptr = (int*)malloc(2 * sizeof(int));
streams =
clCreateBuffer(context, CL_MEM_READ_WRITE, 2 * sizeof(int), NULL, &err);
test_error( err, "clCreateBuffer failed.");
stream = clCreateBuffer(context, CL_MEM_READ_WRITE, 2 * sizeof(cl_int),
nullptr, &err);
test_error(err, "clCreateBuffer failed.");
std::string cl_std = "-cl-std=CL";
cl_std += (get_device_cl_version(device) == Version(3, 0)) ? "3.0" : "2.0";
err = create_single_kernel_helper_with_build_options(
context, &program[0], &kernel[0], 1, &enqueued_local_size_1d_code,
"test_enqueued_local_size_1d", cl_std.c_str());
test_error( err, "create_single_kernel_helper failed");
test_error(err, "create_single_kernel_helper failed");
err = create_single_kernel_helper_with_build_options(
context, &program[1], &kernel[1], 1, &enqueued_local_size_2d_code,
"test_enqueued_local_size_2d", cl_std.c_str());
test_error( err, "create_single_kernel_helper failed");
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.");
err = clSetKernelArg(kernel[0], 0, sizeof stream, &stream);
test_error(err, "clSetKernelArgs failed.");
err = clSetKernelArg(kernel[1], 0, sizeof stream, &stream);
test_error(err, "clSetKernelArgs failed.");
globalsize[0] = (size_t)num_elements;
globalsize[1] = (size_t)num_elements;
globalsize[0] = static_cast<size_t>(num_elements);
globalsize[1] = static_cast<size_t>(num_elements);
size_t max_wgs;
err = clGetDeviceInfo(device, CL_DEVICE_MAX_WORK_GROUP_SIZE, sizeof(max_wgs), &max_wgs, NULL);
test_error( err, "clGetDeviceInfo failed.");
err = clGetDeviceInfo(device, CL_DEVICE_MAX_WORK_GROUP_SIZE,
sizeof(max_wgs), &max_wgs, nullptr);
test_error(err, "clGetDeviceInfo failed.");
localsize[0] = std::min<size_t>(16, max_wgs);
localsize[1] = std::min<size_t>(11, max_wgs / localsize[0]);
@@ -143,35 +143,31 @@ test_enqueued_local_size(cl_device_id device, cl_context context, cl_command_que
}
}
err = clEnqueueNDRangeKernel(queue, kernel[1], 2, NULL, globalsize, localsize, 0, NULL, NULL);
test_error( err, "clEnqueueNDRangeKernel failed.");
err = clEnqueueNDRangeKernel(queue, kernel[1], 2, nullptr, globalsize,
localsize, 0, nullptr, nullptr);
test_error(err, "clEnqueueNDRangeKernel failed.");
err = clEnqueueReadBuffer(queue, streams, CL_TRUE, 0, 2*sizeof(int), output_ptr, 0, NULL, NULL);
test_error( err, "clEnqueueReadBuffer failed.");
err = clEnqueueReadBuffer(queue, stream, CL_BLOCKING, 0, 2 * sizeof(int),
output_ptr, 0, nullptr, nullptr);
test_error(err, "clEnqueueReadBuffer failed.");
err = verify_enqueued_local_size(output_ptr, localsize, 2);
globalsize[0] = (size_t)num_elements;
globalsize[0] = static_cast<size_t>(num_elements);
localsize[0] = 9;
if (use_uniform_work_groups && (globalsize[0] % localsize[0]))
{
globalsize[0] += (localsize[0] - (globalsize[0] % localsize[0]));
}
err = clEnqueueNDRangeKernel(queue, kernel[1], 1, NULL, globalsize, localsize, 0, NULL, NULL);
test_error( err, "clEnqueueNDRangeKernel failed.");
err = clEnqueueNDRangeKernel(queue, kernel[1], 1, nullptr, globalsize,
localsize, 0, nullptr, nullptr);
test_error(err, "clEnqueueNDRangeKernel failed.");
err = clEnqueueReadBuffer(queue, streams, CL_TRUE, 0, 2*sizeof(int), output_ptr, 0, NULL, NULL);
test_error( err, "clEnqueueReadBuffer failed.");
err = clEnqueueReadBuffer(queue, stream, CL_BLOCKING, 0, 2 * sizeof(int),
output_ptr, 0, nullptr, nullptr);
test_error(err, "clEnqueueReadBuffer failed.");
err = verify_enqueued_local_size(output_ptr, localsize, 1);
// cleanup
clReleaseMemObject(streams);
clReleaseKernel(kernel[0]);
clReleaseKernel(kernel[1]);
clReleaseProgram(program[0]);
clReleaseProgram(program[1]);
free(output_ptr);
return err;
}