mirror of
https://github.com/KhronosGroup/OpenCL-CTS.git
synced 2026-03-19 22:19:02 +00:00
85bae70f81
In this test we repeated call create_single_kernel_helper to create different kernels but reuse the same clProgramWrapper. create_single_kernel_helper() creates a new program each time it is called, creating a new reference on the underlying context. When the test ends, the program is released (as it should when using the clProgramWrapper), however there are multiple program objects attached to the context resulting in reference leak errors. Signed-off-by: Chetankumar Mistry <chetan.mistry@arm.com>
402 lines
12 KiB
C++
402 lines
12 KiB
C++
//
|
|
// 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
|
|
//
|
|
// http://www.apache.org/licenses/LICENSE-2.0
|
|
//
|
|
// Unless required by applicable law or agreed to in writing, software
|
|
// distributed under the License is distributed on an "AS IS" BASIS,
|
|
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
|
|
// See the License for the specific language governing permissions and
|
|
// limitations under the License.
|
|
//
|
|
#include "testBase.h"
|
|
#include "harness/typeWrappers.h"
|
|
#include "harness/conversions.h"
|
|
#include <sstream>
|
|
#include <string>
|
|
#include <cmath>
|
|
|
|
using namespace std;
|
|
|
|
const char *clone_kernel_test_img[] =
|
|
{
|
|
"__kernel void img_read_kernel(read_only image2d_t img, sampler_t sampler, __global int* outbuf)\n"
|
|
"{\n"
|
|
" uint4 color;\n"
|
|
"\n"
|
|
" color = read_imageui(img, sampler, (int2)(0,0));\n"
|
|
" \n"
|
|
" // 7, 8, 9, 10th DWORD\n"
|
|
" outbuf[7] = color.x;\n"
|
|
" outbuf[8] = color.y;\n"
|
|
" outbuf[9] = color.z;\n"
|
|
" outbuf[10] = color.w;\n"
|
|
"}\n"
|
|
"\n"
|
|
"__kernel void img_write_kernel(write_only image2d_t img, uint4 color)\n"
|
|
"{\n"
|
|
" write_imageui (img, (int2)(0, 0), color);\n"
|
|
"}\n"
|
|
|
|
};
|
|
|
|
const char *clone_kernel_test_double[] =
|
|
{
|
|
"#pragma OPENCL EXTENSION cl_khr_fp64 : enable\n"
|
|
"__kernel void clone_kernel_test1(double d, __global double* outbuf)\n"
|
|
"{\n"
|
|
" // use the same outbuf as rest of the tests\n"
|
|
" outbuf[2] = d;\n"
|
|
"}\n"
|
|
};
|
|
|
|
const char *clone_kernel_test_kernel[] = {
|
|
"typedef struct\n"
|
|
"{\n"
|
|
" int i;\n"
|
|
" float f;\n"
|
|
"} structArg;\n"
|
|
"\n"
|
|
"// value type test\n"
|
|
"__kernel void clone_kernel_test0(int iarg, float farg, structArg sarg, __local int* localbuf, __global int* outbuf)\n"
|
|
"{\n"
|
|
" int tid = get_global_id(0);\n"
|
|
"\n"
|
|
" outbuf[0] = iarg;\n"
|
|
" outbuf[1] = sarg.i;\n"
|
|
" \n"
|
|
" ((__global float*)outbuf)[2] = farg;\n"
|
|
" ((__global float*)outbuf)[3] = sarg.f;\n"
|
|
"}\n"
|
|
"\n"
|
|
"__kernel void buf_read_kernel(__global int* buf, __global int* outbuf)\n"
|
|
"{\n"
|
|
" // 6th DWORD\n"
|
|
" outbuf[6] = buf[0];\n"
|
|
"}\n"
|
|
"\n"
|
|
"__kernel void buf_write_kernel(__global int* buf, int write_val)\n"
|
|
"{\n"
|
|
" buf[0] = write_val;\n"
|
|
"}\n"
|
|
|
|
};
|
|
|
|
const int BUF_SIZE = 128;
|
|
|
|
struct structArg
|
|
{
|
|
int i;
|
|
float f;
|
|
};
|
|
|
|
static unsigned char *
|
|
generate_8888_image(int w, int h, MTdata d)
|
|
{
|
|
unsigned char *ptr = (unsigned char*)malloc(w * h * 4);
|
|
int i;
|
|
|
|
for (i=0; i<w*h*4; i++)
|
|
ptr[i] = (unsigned char)genrand_int32( d);
|
|
|
|
return ptr;
|
|
}
|
|
|
|
int test_image_arg_shallow_clone(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements, void* pbufRes, clMemWrapper& bufOut)
|
|
{
|
|
int error;
|
|
cl_image_format img_format;
|
|
clSamplerWrapper sampler;
|
|
img_format.image_channel_order = CL_RGBA;
|
|
img_format.image_channel_data_type = CL_UNSIGNED_INT8;
|
|
cl_image_desc imageDesc;
|
|
memset(&imageDesc, 0x0, sizeof(cl_image_desc));
|
|
imageDesc.image_type = CL_MEM_OBJECT_IMAGE2D;
|
|
imageDesc.image_width = 512;
|
|
imageDesc.image_height = 512;
|
|
|
|
cl_uint color[4] = {1,3,5,7};
|
|
|
|
clProgramWrapper program_read;
|
|
clProgramWrapper program_write;
|
|
clKernelWrapper kernel_read;
|
|
clKernelWrapper kernel_write;
|
|
clKernelWrapper kernel_cloned;
|
|
size_t ndrange1 = 1;
|
|
|
|
clMemWrapper img;
|
|
|
|
if (create_single_kernel_helper(context, &program_read, &kernel_read, 1,
|
|
clone_kernel_test_img, "img_read_kernel")
|
|
!= 0)
|
|
{
|
|
return -1;
|
|
}
|
|
|
|
if (create_single_kernel_helper(context, &program_write, &kernel_write, 1,
|
|
clone_kernel_test_img, "img_write_kernel")
|
|
!= 0)
|
|
{
|
|
return -1;
|
|
}
|
|
|
|
img = clCreateImage(context, CL_MEM_READ_WRITE, &img_format, &imageDesc, NULL, &error);
|
|
test_error( error, "clCreateImage failed." );
|
|
|
|
cl_sampler_properties properties[] = {
|
|
CL_SAMPLER_NORMALIZED_COORDS, CL_FALSE,
|
|
CL_SAMPLER_ADDRESSING_MODE, CL_ADDRESS_CLAMP_TO_EDGE,
|
|
CL_SAMPLER_FILTER_MODE, CL_FILTER_NEAREST,
|
|
0 };
|
|
sampler = clCreateSamplerWithProperties(context, properties, &error);
|
|
test_error( error, "clCreateSamplerWithProperties failed." );
|
|
|
|
error = clSetKernelArg(kernel_write, 1, sizeof(int) * 4, color);
|
|
error += clSetKernelArg(kernel_write, 0, sizeof(cl_mem), &img);
|
|
test_error( error, "clSetKernelArg failed." );
|
|
|
|
error = clEnqueueNDRangeKernel(queue, kernel_write, 1, NULL, &ndrange1, NULL, 0, NULL, NULL);
|
|
test_error( error, "clEnqueueNDRangeKernel failed." );
|
|
|
|
error = clSetKernelArg(kernel_read, 0, sizeof(cl_mem), &img);
|
|
error += clSetKernelArg(kernel_read, 1, sizeof(cl_sampler), &sampler);
|
|
error += clSetKernelArg(kernel_read, 2, sizeof(cl_mem), &bufOut);
|
|
|
|
test_error( error, "clSetKernelArg failed." );
|
|
|
|
// clone the kernel
|
|
kernel_cloned = clCloneKernel(kernel_read, &error);
|
|
test_error( error, "clCloneKernel failed." );
|
|
error = clEnqueueNDRangeKernel(queue, kernel_cloned, 1, NULL, &ndrange1, NULL, 0, NULL, NULL);
|
|
test_error( error, "clEnqueueNDRangeKernel failed." );
|
|
|
|
// read result back
|
|
error = clEnqueueReadBuffer(queue, bufOut, CL_TRUE, 0, 128, pbufRes, 0, NULL, NULL);
|
|
test_error( error, "clEnqueueReadBuffer failed." );
|
|
|
|
if (((cl_uint*)pbufRes)[7] != color[0])
|
|
{
|
|
test_error( error, "clCloneKernel test failed." );
|
|
return -1;
|
|
}
|
|
|
|
if (((cl_uint*)pbufRes)[8] != color[1])
|
|
{
|
|
test_error( error, "clCloneKernel test failed." );
|
|
return -1;
|
|
}
|
|
|
|
if (((cl_uint*)pbufRes)[9] != color[2])
|
|
{
|
|
test_error( error, "clCloneKernel test failed." );
|
|
return -1;
|
|
}
|
|
|
|
if (((cl_uint*)pbufRes)[10] != color[3])
|
|
{
|
|
test_error( error, "clCloneKernel test failed." );
|
|
return -1;
|
|
}
|
|
|
|
return 0;
|
|
}
|
|
|
|
int test_double_arg_clone(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements, void* pbufRes, clMemWrapper& bufOut)
|
|
{
|
|
int error = 0;
|
|
clProgramWrapper program;
|
|
clKernelWrapper kernel;
|
|
clKernelWrapper kernel_cloned;
|
|
size_t ndrange1 = 1;
|
|
|
|
if( create_single_kernel_helper( context, &program, &kernel, 1, clone_kernel_test_double, "clone_kernel_test1" ) != 0 )
|
|
{
|
|
return -1;
|
|
}
|
|
|
|
cl_double d = 1.23;
|
|
error = clSetKernelArg(kernel, 0, sizeof(double), &d);
|
|
error += clSetKernelArg(kernel, 1, sizeof(cl_mem), &bufOut);
|
|
test_error( error, "clSetKernelArg failed." );
|
|
|
|
kernel_cloned = clCloneKernel(kernel, &error);
|
|
test_error( error, "clCloneKernel failed." );
|
|
|
|
error = clEnqueueNDRangeKernel(queue, kernel_cloned, 1, NULL, &ndrange1, NULL, 0, NULL, NULL);
|
|
test_error( error, "clEnqueueNDRangeKernel failed." );
|
|
|
|
// read result back
|
|
error = clEnqueueReadBuffer(queue, bufOut, CL_TRUE, 0, BUF_SIZE, pbufRes, 0, NULL, NULL);
|
|
test_error( error, "clEnqueueReadBuffer failed." );
|
|
|
|
if (abs(((cl_double*)pbufRes)[2] - d) > 0.0000001)
|
|
{
|
|
test_error( error, "clCloneKernel test failed." );
|
|
return -1;
|
|
}
|
|
|
|
return 0;
|
|
}
|
|
|
|
int test_clone_kernel(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
|
|
{
|
|
int error;
|
|
clProgramWrapper program;
|
|
clProgramWrapper program_buf_read;
|
|
clProgramWrapper program_buf_write;
|
|
clKernelWrapper kernel;
|
|
clKernelWrapper kernel_pipe_read;
|
|
clKernelWrapper kernel_buf_read;
|
|
clKernelWrapper kernel_pipe_write;
|
|
clKernelWrapper kernel_buf_write;
|
|
|
|
clKernelWrapper kernel_pipe_read_cloned;
|
|
clKernelWrapper kernel_buf_read_cloned;
|
|
size_t ndrange1 = 1;
|
|
|
|
int write_val = 123;
|
|
|
|
|
|
cl_bool bimg = CL_FALSE;
|
|
cl_bool bdouble = CL_FALSE;
|
|
// test image support
|
|
error = clGetDeviceInfo(deviceID, CL_DEVICE_IMAGE_SUPPORT, sizeof(cl_bool), &bimg, NULL);
|
|
test_error( error, "clGetDeviceInfo failed." );
|
|
|
|
// test double support
|
|
if (is_extension_available(deviceID, "cl_khr_fp64"))
|
|
{
|
|
bdouble = CL_TRUE;
|
|
}
|
|
|
|
/* Create kernels to test with */
|
|
if( create_single_kernel_helper( context, &program, &kernel, 1, clone_kernel_test_kernel, "clone_kernel_test0" ) != 0 )
|
|
{
|
|
return -1;
|
|
}
|
|
|
|
if (create_single_kernel_helper(context, &program_buf_read,
|
|
&kernel_buf_read, 1,
|
|
clone_kernel_test_kernel, "buf_read_kernel")
|
|
!= 0)
|
|
{
|
|
return -1;
|
|
}
|
|
|
|
if (create_single_kernel_helper(
|
|
context, &program_buf_write, &kernel_buf_write, 1,
|
|
clone_kernel_test_kernel, "buf_write_kernel")
|
|
!= 0)
|
|
{
|
|
return -1;
|
|
}
|
|
|
|
// Kernel args
|
|
// Value type
|
|
int intarg = 0;
|
|
float farg = 1.0;
|
|
structArg sa = { 1, 1.0f };
|
|
|
|
// cl_mem
|
|
clMemWrapper buf, bufOut;
|
|
|
|
char* pbuf = new char[BUF_SIZE];
|
|
char* pbufRes = new char[BUF_SIZE];
|
|
buf = clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM_USE_HOST_PTR, BUF_SIZE, pbuf, &error);
|
|
test_error( error, "clCreateBuffer failed." );
|
|
|
|
bufOut = clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR, BUF_SIZE, NULL, &error);
|
|
test_error( error, "clCreateBuffer failed." );
|
|
|
|
error = clSetKernelArg(kernel, 0, sizeof(int), &intarg);
|
|
error += clSetKernelArg(kernel, 1, sizeof(float), &farg);
|
|
error += clSetKernelArg(kernel, 2, sizeof(structArg), &sa);
|
|
error += clSetKernelArg(kernel, 3, 128, NULL); // local mem
|
|
|
|
test_error( error, "clSetKernelArg failed." );
|
|
|
|
// clone the kernel
|
|
clKernelWrapper clonek = clCloneKernel(kernel, &error);
|
|
test_error( error, "clCloneKernel failed." );
|
|
|
|
// set the last arg and enqueue
|
|
error = clSetKernelArg(clonek, 4, sizeof(cl_mem), &bufOut);
|
|
test_error( error, "clSetKernelArg failed." );
|
|
error = clEnqueueNDRangeKernel(queue, clonek, 1, NULL, &ndrange1, NULL, 0, NULL, NULL);
|
|
test_error( error, "clEnqueueNDRangeKernel failed." );
|
|
|
|
// shallow clone tests for buffer
|
|
error = clSetKernelArg(kernel_buf_write, 0, sizeof(cl_mem), &buf);
|
|
error += clSetKernelArg(kernel_buf_write, 1, sizeof(int), &write_val);
|
|
test_error( error, "clSetKernelArg failed." );
|
|
error = clEnqueueNDRangeKernel(queue, kernel_buf_write, 1, NULL, &ndrange1, NULL, 0, NULL, NULL);
|
|
test_error( error, "clEnqueueNDRangeKernel failed." );
|
|
|
|
error = clSetKernelArg(kernel_buf_read, 0, sizeof(cl_mem), &buf);
|
|
error += clSetKernelArg(kernel_buf_read, 1, sizeof(cl_mem), &bufOut);
|
|
test_error( error, "clSetKernelArg failed." );
|
|
|
|
// clone the kernel
|
|
kernel_buf_read_cloned = clCloneKernel(kernel_buf_read, &error);
|
|
test_error( error, "clCloneKernel API call failed." );
|
|
error = clEnqueueNDRangeKernel(queue, kernel_buf_read_cloned, 1, NULL, &ndrange1, NULL, 0, NULL, NULL);
|
|
test_error( error, "clEnqueueNDRangeKernel failed." );
|
|
|
|
// read result back
|
|
error = clEnqueueReadBuffer(queue, bufOut, CL_TRUE, 0, BUF_SIZE, pbufRes, 0, NULL, NULL);
|
|
test_error( error, "clEnqueueReadBuffer failed." );
|
|
|
|
// Compare the results
|
|
if (((int*)pbufRes)[0] != intarg)
|
|
{
|
|
test_error( error, "clCloneKernel test failed. Failed to clone integer type argument." );
|
|
return -1;
|
|
}
|
|
|
|
if (((int*)pbufRes)[1] != sa.i)
|
|
{
|
|
test_error( error, "clCloneKernel test failed. Failed to clone structure type argument." );
|
|
return -1;
|
|
}
|
|
|
|
if (((float*)pbufRes)[2] != farg)
|
|
{
|
|
test_error( error, "clCloneKernel test failed. Failed to clone structure type argument." );
|
|
return -1;
|
|
}
|
|
|
|
if (((float*)pbufRes)[3] != sa.f)
|
|
{
|
|
test_error( error, "clCloneKernel test failed. Failed to clone float type argument." );
|
|
return -1;
|
|
}
|
|
|
|
if (((int*)pbufRes)[6] != write_val)
|
|
{
|
|
test_error( error, "clCloneKernel test failed. Failed to clone cl_mem argument." );
|
|
return -1;
|
|
}
|
|
|
|
if (bimg)
|
|
{
|
|
error = test_image_arg_shallow_clone(deviceID, context, queue, num_elements, pbufRes, bufOut);
|
|
test_error( error, "image arg shallow clone test failed." );
|
|
}
|
|
|
|
if (bdouble)
|
|
{
|
|
error = test_double_arg_clone(deviceID, context, queue, num_elements, pbufRes, bufOut);
|
|
test_error( error, "double arg clone test failed." );
|
|
}
|
|
|
|
delete [] pbuf;
|
|
delete [] pbufRes;
|
|
|
|
return 0;
|
|
}
|
|
|