Mem-leaks from conformance pipes (#772)
Fix various memory leaks around events.
Convert test to use supplied typewrappers
to avoid memory leaks. Also use error helper
functions to reduce code size.
Use stringstreams to synthesize kernel sources, and
raw c+11 string literals.
Signed-off-by: John Kesapides <john.kesapides@arm.com>
diff --git a/test_conformance/pipes/test_pipe_info.cpp b/test_conformance/pipes/test_pipe_info.cpp
index 5d3e3a4..7543c6c 100644
--- a/test_conformance/pipes/test_pipe_info.cpp
+++ b/test_conformance/pipes/test_pipe_info.cpp
@@ -21,29 +21,25 @@
int test_pipe_info( cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements )
{
- cl_mem pipe;
+ clMemWrapper pipe;
cl_int err;
cl_uint pipe_width = 512;
cl_uint pipe_depth = 1024;
cl_uint returnVal;
- cl_program program;
- cl_kernel kernel;
+ clProgramWrapper program;
+ clKernelWrapper kernel;
- pipe = clCreatePipe(context, CL_MEM_HOST_NO_ACCESS, pipe_width, pipe_depth, NULL, &err);
+ pipe = clCreatePipe(context, CL_MEM_HOST_NO_ACCESS, pipe_width, pipe_depth,
+ NULL, &err);
test_error(err, "clCreatePipe failed.");
- err = clGetPipeInfo(pipe, CL_PIPE_PACKET_SIZE, sizeof(pipe_width), (void *)&returnVal, NULL);
- if ( err )
- {
- log_error( "Error calling clGetPipeInfo(): %d\n", err );
- clReleaseMemObject(pipe);
- return -1;
- }
+ err = clGetPipeInfo(pipe, CL_PIPE_PACKET_SIZE, sizeof(pipe_width),
+ (void *)&returnVal, NULL);
+ test_error(err, "clGetPipeInfo failed.");
- if(pipe_width != returnVal)
+ if (pipe_width != returnVal)
{
- log_error( "Error in clGetPipeInfo() check of pipe packet size\n" );
- clReleaseMemObject(pipe);
+ log_error("Error in clGetPipeInfo() check of pipe packet size\n");
return -1;
}
else
@@ -52,17 +48,11 @@
}
err = clGetPipeInfo(pipe, CL_PIPE_MAX_PACKETS, sizeof(pipe_depth), (void *)&returnVal, NULL);
- if ( err )
- {
- log_error( "Error calling clGetPipeInfo(): %d\n", err );
- clReleaseMemObject(pipe);
- return -1;
- }
+ test_error(err, "clGetPipeInfo failed.");
if(pipe_depth != returnVal)
{
log_error( "Error in clGetPipeInfo() check of pipe max packets\n" );
- clReleaseMemObject(pipe);
return -1;
}
else
@@ -71,39 +61,20 @@
}
err = create_single_kernel_helper_with_build_options(context, &program, &kernel, 1, (const char**)&pipe_kernel_code, "pipe_kernel", "-cl-std=CL2.0 -cl-kernel-arg-info");
- if(err)
- {
- clReleaseMemObject(pipe);
- print_error(err, "Error creating program\n");
- return -1;
- }
+ test_error_ret(err, " Error creating program", -1);
cl_kernel_arg_type_qualifier arg_type_qualifier = 0;
cl_kernel_arg_type_qualifier expected_type_qualifier = CL_KERNEL_ARG_TYPE_PIPE;
err = clGetKernelArgInfo( kernel, 0, CL_KERNEL_ARG_TYPE_QUALIFIER, sizeof(arg_type_qualifier), &arg_type_qualifier, NULL );
- if(err)
- {
- clReleaseMemObject(pipe);
- clReleaseKernel(kernel);
- clReleaseProgram(program);
- print_error(err, "clSetKernelArg failed\n");
- return -1;
- }
+ test_error_ret(err, " clSetKernelArgInfo failed", -1);
err = (arg_type_qualifier != expected_type_qualifier);
+
if(err)
{
- clReleaseMemObject(pipe);
- clReleaseKernel(kernel);
- clReleaseProgram(program);
print_error(err, "ERROR: Bad type qualifier\n");
return -1;
}
- // cleanup
- clReleaseMemObject(pipe);
- clReleaseKernel(kernel);
- clReleaseProgram(program);
-
return err;
}
diff --git a/test_conformance/pipes/test_pipe_limits.cpp b/test_conformance/pipes/test_pipe_limits.cpp
index 29199df..85247f8 100644
--- a/test_conformance/pipes/test_pipe_limits.cpp
+++ b/test_conformance/pipes/test_pipe_limits.cpp
@@ -15,63 +15,103 @@
//
#include "harness/compat.h"
+#include <assert.h>
+#include <iomanip>
+#include <iostream>
+#include <sstream>
#include <stdio.h>
#include <string.h>
-#include <assert.h>
-#include <sys/types.h>
+#include <string>
#include <sys/stat.h>
+#include <sys/types.h>
#include "procs.h"
#include "harness/errorHelpers.h"
#define STRING_LENGTH 1024
-void createKernelSourceCode(char *source, int num_pipes)
+void createKernelSourceCode(std::stringstream &stream, int num_pipes)
{
int i;
- char str[256];
- int str_length;
- strcpy(source, "__kernel void test_multiple_pipe_write(__global int *src, ");
+ stream << "__kernel void test_multiple_pipe_write(__global int *src, ";
+ for (i = 0; i < num_pipes; i++)
+ {
+ stream << "__write_only pipe int pipe" << i << ", ";
+ }
+ stream << R"(int num_pipes )
+ {
+ int gid = get_global_id(0);
+ reserve_id_t res_id;
- for(i = 0; i < num_pipes; i++) {
- sprintf(str, "__write_only pipe int pipe%d, ", i);
- strcat(source, str);
- }
- sprintf(str, "int num_pipes ) \n{\n int gid = get_global_id(0);\n reserve_id_t res_id;\n\n");
- strcat(source, str);
- sprintf(str, " if(gid < (get_global_size(0))/num_pipes)\n {\n res_id = reserve_write_pipe(pipe0, 1);\n if(is_valid_reserve_id(res_id))\n {\n");
- strcat(source, str);
- sprintf(str, " write_pipe(pipe0, res_id, 0, &src[gid]);\n commit_write_pipe(pipe0, res_id);\n }\n }\n");
- strcat(source, str);
- for(i = 1; i < num_pipes; i++){
- sprintf(str, " else if(gid < (%d*get_global_size(0))/num_pipes)\n {\n res_id = reserve_write_pipe(pipe%d, 1);\n if(is_valid_reserve_id(res_id))\n {\n", i+1, i);
- strcat(source, str);
- sprintf(str, " write_pipe(pipe%d, res_id, 0, &src[gid]);\n commit_write_pipe(pipe%d, res_id);\n }\n }\n", i, i);
- strcat(source, str);
- }
- strcat(source, "}\n\n__kernel void test_multiple_pipe_read(__global int *dst, ");
- for(i = 0; i < num_pipes; i++) {
- sprintf(str, "__read_only pipe int pipe%d, ", i);
- strcat(source, str);
- }
- sprintf(str, "int num_pipes ) \n{\n int gid = get_global_id(0);\n reserve_id_t res_id;\n\n");
- strcat(source, str);
- sprintf(str, " if(gid < (get_global_size(0))/num_pipes)\n {\n res_id = reserve_read_pipe(pipe0, 1);\n if(is_valid_reserve_id(res_id))\n {\n");
- strcat(source, str);
- sprintf(str, " read_pipe(pipe0, res_id, 0, &dst[gid]);\n commit_read_pipe(pipe0, res_id);\n }\n }\n");
- strcat(source, str);
- for(i = 1; i < num_pipes; i++){
- sprintf(str, " else if(gid < (%d*get_global_size(0))/num_pipes)\n {\n res_id = reserve_read_pipe(pipe%d, 1);\n if(is_valid_reserve_id(res_id))\n {\n", i+1, i);
- strcat(source, str);
- sprintf(str, " read_pipe(pipe%d, res_id, 0, &dst[gid]);\n commit_read_pipe(pipe%d, res_id);\n }\n }\n", i, i);
- strcat(source, str);
- }
- strcat(source, "}");
+ if(gid < (get_global_size(0))/num_pipes)
+ {
+ res_id = reserve_write_pipe(pipe0, 1);
+ if(is_valid_reserve_id(res_id))
+ {
+ write_pipe(pipe0, res_id, 0, &src[gid]);
+ commit_write_pipe(pipe0, res_id);
+ }
+ })";
- str_length = strlen(source);
- assert(str_length <= STRING_LENGTH*num_pipes);
+ for (i = 1; i < num_pipes; i++)
+ {
+ // clang-format off
+ stream << R"(
+ else if(gid < ()" << (i + 1) << R"(*get_global_size(0))/num_pipes)
+ {
+ res_id = reserve_write_pipe(pipe)" << i << R"(, 1);
+ if(is_valid_reserve_id(res_id))
+ {
+ write_pipe(pipe)" << i << R"(, res_id, 0, &src[gid]);
+ commit_write_pipe(pipe)" << i << R"(, res_id);
+ }
+ }
+ )";
+ // clang-format om
+ }
+ stream << R"(
+ }
+
+ __kernel void test_multiple_pipe_read(__global int *dst, )";
+
+ for (i = 0; i < num_pipes; i++)
+ {
+ stream << "__read_only pipe int pipe" << i << ", ";
+ }
+ stream << R"(int num_pipes )
+ {
+ int gid = get_global_id(0);
+ reserve_id_t res_id;
+
+
+ if(gid < (get_global_size(0))/num_pipes)
+ {
+ res_id = reserve_read_pipe(pipe0, 1);
+ if(is_valid_reserve_id(res_id))
+ {
+ read_pipe(pipe0, res_id, 0, &dst[gid]);
+ commit_read_pipe(pipe0, res_id);
+ }
+ })";
+
+ for (i = 1; i < num_pipes; i++)
+ {
+ // clang-format off
+ stream << R"(
+ else if(gid < ()" << (i + 1) << R"(*get_global_size(0))/num_pipes)
+ {
+ res_id = reserve_read_pipe(pipe)" << i << R"(, 1);
+ if(is_valid_reserve_id(res_id))
+ {
+ read_pipe(pipe)" << i << R"(, res_id, 0, &dst[gid]);
+ commit_read_pipe(pipe)" << i << R"(, res_id);
+ }
+ })";
+ // clang-format on
+ }
+ stream << "}";
}
static int verify_result(void *ptr1, void *ptr2, int n)
@@ -113,28 +153,34 @@
int test_pipe_max_args(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
{
- cl_mem pipes[1024];
- cl_mem buffers[2];
- void *outptr;
- cl_int *inptr;
- cl_program program;
- cl_kernel kernel[2];
- size_t global_work_size[3];
- cl_int err;
- cl_int size;
- int num_pipe_elements = 1024;
- int i, j;
- int max_pipe_args;
- char *source;
- cl_event producer_sync_event = NULL;
- cl_event consumer_sync_event = NULL;
- MTdata d = init_genrand( gRandomSeed );
- const char* kernelName[] = {"test_multiple_pipe_write", "test_multiple_pipe_read"};
+ clMemWrapper pipes[1024];
+ clMemWrapper buffers[2];
+ void *outptr;
+ cl_int *inptr;
+ clProgramWrapper program;
+ clKernelWrapper kernel[2];
+ size_t global_work_size[3];
+ cl_int err;
+ cl_int size;
+ int num_pipe_elements = 1024;
+ int i, j;
+ int max_pipe_args;
+ std::stringstream source;
+ clEventWrapper producer_sync_event = NULL;
+ clEventWrapper consumer_sync_event = NULL;
+ BufferOwningPtr<cl_int> BufferInPtr;
+ BufferOwningPtr<cl_int> BufferOutPtr;
- size_t min_alignment = get_min_alignment(context);
+ MTdataHolder d(gRandomSeed);
+ const char *kernelName[] = { "test_multiple_pipe_write",
+ "test_multiple_pipe_read" };
- err = clGetDeviceInfo(deviceID, CL_DEVICE_MAX_PIPE_ARGS, sizeof(max_pipe_args), (void*)&max_pipe_args, NULL);
- if(err){
+ size_t min_alignment = get_min_alignment(context);
+
+ err = clGetDeviceInfo(deviceID, CL_DEVICE_MAX_PIPE_ARGS,
+ sizeof(max_pipe_args), (void *)&max_pipe_args, NULL);
+ if (err)
+ {
print_error(err, " clGetDeviceInfo failed\n");
return -1;
}
@@ -145,76 +191,41 @@
global_work_size[0] = (cl_uint)num_pipe_elements * max_pipe_args;
size = sizeof(int) * num_pipe_elements * max_pipe_args;
- source = (char *)malloc(STRING_LENGTH * sizeof(char) * max_pipe_args);
inptr = (cl_int *)align_malloc(size, min_alignment);
for(i = 0; i < num_pipe_elements * max_pipe_args; i++){
inptr[i] = (int)genrand_int32(d);
}
+ BufferInPtr.reset(inptr, nullptr, 0, size, true);
buffers[0] = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR, size, inptr, &err);
- if(err){
- clReleaseMemObject(buffers[0]);
- free(source);
- print_error(err, " clCreateBuffer failed\n");
- return -1;
- }
+ test_error_ret(err, " clCreateBuffer failed", -1);
+
outptr = align_malloc(size, min_alignment);
+ BufferOutPtr.reset(outptr, nullptr, 0, size, true);
buffers[1] = clCreateBuffer(context, CL_MEM_USE_HOST_PTR, size, outptr, &err);
- if ( err ){
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- align_free( outptr );
- free(source);
- print_error(err, " clCreateBuffer failed\n" );
- return -1;
- }
+ test_error_ret(err, " clCreateBuffer failed", -1);
for(i = 0; i < max_pipe_args; i++){
pipes[i] = clCreatePipe(context, CL_MEM_HOST_NO_ACCESS, sizeof(int), num_pipe_elements, NULL, &err);
- if(err){
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- align_free( outptr );
- free(source);
- for(j = 0; j < i; j++) {
- clReleaseMemObject(pipes[j]);
- }
- print_error(err, " clCreatePipe failed\n");
- return -1;
- }
+ test_error_ret(err, " clCreatePipe failed", -1);
}
createKernelSourceCode(source, max_pipe_args);
+ std::string kernel_source = source.str();
+ const char *sources[] = { kernel_source.c_str() };
+
// Create producer kernel
- err = create_single_kernel_helper_with_build_options(context, &program, &kernel[0], 1, (const char**)&source, kernelName[0], "-cl-std=CL2.0");
- if(err){
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- for(j = 0; j < max_pipe_args; j++) {
- clReleaseMemObject(pipes[j]);
- }
- align_free(outptr);
- free(source);
- print_error(err, "Error creating program\n");
- return -1;
- }
+ err = create_single_kernel_helper_with_build_options(
+ context, &program, &kernel[0], 1, sources, kernelName[0],
+ "-cl-std=CL2.0");
+ test_error_ret(err, " Error creating program", -1);
+
//Create consumer kernel
kernel[1] = clCreateKernel(program, kernelName[1], &err);
- if( kernel[1] == NULL || err != CL_SUCCESS)
- {
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- for(j = 0; j < max_pipe_args; j++) {
- clReleaseMemObject(pipes[j]);
- }
- align_free(outptr);
- free(source);
- print_error(err, " Error creating kernel\n");
- return -1;
- }
+ test_error_ret(err, " Error creating kernel", -1);
err = clSetKernelArg(kernel[0], 0, sizeof(cl_mem), (void*)&buffers[0]);
for( i = 0; i < max_pipe_args; i++){
@@ -226,127 +237,28 @@
err |= clSetKernelArg(kernel[1], i+1, sizeof(cl_mem), (void*)&pipes[i]);
}
err |= clSetKernelArg(kernel[1], max_pipe_args + 1, sizeof(int), (void*)&max_pipe_args);
- if ( err != CL_SUCCESS ){
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- for(j = 0; j < max_pipe_args; j++) {
- clReleaseMemObject(pipes[j]);
- }
- clReleaseKernel(kernel[0]);
- clReleaseKernel(kernel[1]);
- clReleaseProgram(program);
- align_free(outptr);
- free(source);
- print_error(err, " clSetKernelArg failed");
- return -1;
- }
+ test_error_ret(err, " clSetKernelArg failed", -1);
// Launch Producer kernel
err = clEnqueueNDRangeKernel( queue, kernel[0], 1, NULL, global_work_size, NULL, 0, NULL, &producer_sync_event );
- if ( err != CL_SUCCESS ){
- print_error( err, " clEnqueueNDRangeKernel failed" );
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- for(j = 0; j < max_pipe_args; j++) {
- clReleaseMemObject(pipes[j]);
- }
- clReleaseKernel(kernel[0]);
- clReleaseKernel(kernel[1]);
- clReleaseEvent(producer_sync_event);
- clReleaseEvent(consumer_sync_event);
- clReleaseProgram(program);
- align_free(outptr);
- free(source);
- return -1;
- }
+ test_error_ret(err, " clEnqueueNDRangeKernel failed", -1);
// Launch Consumer kernel
err = clEnqueueNDRangeKernel( queue, kernel[1], 1, NULL, global_work_size, NULL, 1, &producer_sync_event, &consumer_sync_event );
- if ( err != CL_SUCCESS ){
- print_error( err, " clEnqueueNDRangeKernel failed" );
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- for(j = 0; j < max_pipe_args; j++) {
- clReleaseMemObject(pipes[j]);
- }
- clReleaseKernel(kernel[0]);
- clReleaseKernel(kernel[1]);
- clReleaseEvent(producer_sync_event);
- clReleaseEvent(consumer_sync_event);
- clReleaseProgram(program);
- align_free(outptr);
- free(source);
- return -1;
- }
+ test_error_ret(err, " clEnqueueNDRangeKernel failed", -1);
err = clEnqueueReadBuffer(queue, buffers[1], true, 0, size, outptr, 1, &consumer_sync_event, NULL);
- if ( err != CL_SUCCESS ){
- print_error( err, " clEnqueueReadBuffer failed" );
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- for(j = 0; j < max_pipe_args; j++) {
- clReleaseMemObject(pipes[j]);
- }
- clReleaseKernel(kernel[0]);
- clReleaseKernel(kernel[1]);
- clReleaseEvent(producer_sync_event);
- clReleaseEvent(consumer_sync_event);
- clReleaseProgram(program);
- align_free(outptr);
- free(source);
- return -1;
- }
+ test_error_ret(err, " clEnqueueNDRangeKernel failed", -1);
err = clWaitForEvents(1, &consumer_sync_event);
- if ( err != CL_SUCCESS ){
- print_error( err, " clWaitForEvents failed" );
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- for(j = 0; j < max_pipe_args; j++) {
- clReleaseMemObject(pipes[j]);
- }
- clReleaseKernel(kernel[0]);
- clReleaseKernel(kernel[1]);
- clReleaseEvent(producer_sync_event);
- clReleaseEvent(consumer_sync_event);
- clReleaseProgram(program);
- align_free(outptr);
- free(source);
- return -1;
- }
+ test_error_ret(err, " clWaitForEvents failed", -1);
if( verify_result( inptr, outptr, num_pipe_elements*sizeof(cl_int))){
log_error("test_pipe_max_args failed\n");
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- for(j = 0; j < max_pipe_args; j++) {
- clReleaseMemObject(pipes[j]);
- }
- clReleaseKernel(kernel[0]);
- clReleaseKernel(kernel[1]);
- clReleaseEvent(producer_sync_event);
- clReleaseEvent(consumer_sync_event);
- clReleaseProgram(program);
- align_free(outptr);
- free(source);
- return -1;
}
else {
log_info("test_pipe_max_args passed\n");
}
- //cleanup
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- for(j = 0; j < max_pipe_args; j++) {
- clReleaseMemObject(pipes[j]);
- }
- clReleaseKernel(kernel[0]);
- clReleaseKernel(kernel[1]);
- clReleaseEvent(producer_sync_event);
- clReleaseEvent(consumer_sync_event);
- clReleaseProgram(program);
- align_free(outptr);
- free(source);
return 0;
}
@@ -354,39 +266,41 @@
int test_pipe_max_packet_size(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
{
- cl_mem pipe;
- cl_mem buffers[2];
- void *outptr;
- cl_char *inptr;
- cl_program program;
- cl_kernel kernel[2];
- size_t global_work_size[3];
- cl_int err;
- size_t size;
- int num_pipe_elements = 1024;
- int i;
- cl_uint max_pipe_packet_size;
- char *source;
- char str[256];
- int str_length;
- cl_event producer_sync_event = NULL;
- cl_event consumer_sync_event = NULL;
- MTdata d = init_genrand( gRandomSeed );
- const char* kernelName[] = {"test_pipe_max_packet_size_write", "test_pipe_max_packet_size_read"};
+ clMemWrapper pipe;
+ clMemWrapper buffers[2];
+ void *outptr;
+ cl_char *inptr;
+ clProgramWrapper program;
+ clKernelWrapper kernel[2];
+ size_t global_work_size[3];
+ cl_int err;
+ size_t size;
+ int num_pipe_elements = 1024;
+ int i;
+ cl_uint max_pipe_packet_size;
+ clEventWrapper producer_sync_event = NULL;
+ clEventWrapper consumer_sync_event = NULL;
+ BufferOwningPtr<cl_int> BufferInPtr;
+ BufferOwningPtr<cl_int> BufferOutPtr;
+ MTdataHolder d(gRandomSeed);
+ const char *kernelName[] = { "test_pipe_max_packet_size_write",
+ "test_pipe_max_packet_size_read" };
- size_t min_alignment = get_min_alignment(context);
+ size_t min_alignment = get_min_alignment(context);
global_work_size[0] = (cl_uint)num_pipe_elements;
- source = (char*)malloc(STRING_LENGTH*sizeof(char));
+ std::stringstream source;
- err = clGetDeviceInfo(deviceID, CL_DEVICE_PIPE_MAX_PACKET_SIZE, sizeof(max_pipe_packet_size), (void*)&max_pipe_packet_size, NULL);
- if(err){
- print_error(err, " clGetDeviceInfo failed\n");
- return -1;
- }
- if(max_pipe_packet_size < 1024){
- log_error("The device should support minimum packet size of 1024 bytes");
+ err = clGetDeviceInfo(deviceID, CL_DEVICE_PIPE_MAX_PACKET_SIZE,
+ sizeof(max_pipe_packet_size),
+ (void *)&max_pipe_packet_size, NULL);
+ test_error_ret(err, " clCreatePipe failed", -1);
+
+ if (max_pipe_packet_size < 1024)
+ {
+ log_error(
+ "The device should support minimum packet size of 1024 bytes");
return -1;
}
@@ -402,354 +316,248 @@
for(i = 0; i < size; i++){
inptr[i] = (char)genrand_int32(d);
}
+ BufferInPtr.reset(inptr, nullptr, 0, size, true);
buffers[0] = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR, size, inptr, &err);
- if(err){
- clReleaseMemObject(buffers[0]);
- free(source);
- print_error(err, " clCreateBuffer failed\n");
- return -1;
- }
+ test_error_ret(err, " clCreateBuffer failed", -1);
+
outptr = align_malloc(size, min_alignment);
+ BufferOutPtr.reset(outptr, nullptr, 0, size, true);
+
buffers[1] = clCreateBuffer(context, CL_MEM_USE_HOST_PTR, size, outptr, &err);
- if ( err ){
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- align_free( outptr );
- free(source);
- print_error(err, " clCreateBuffer failed\n" );
- return -1;
- }
+ test_error_ret(err, " clCreateBuffer failed", -1);
pipe = clCreatePipe(context, CL_MEM_HOST_NO_ACCESS, max_pipe_packet_size, num_pipe_elements, NULL, &err);
- if(err){
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- align_free( outptr );
- free(source);
- clReleaseMemObject(pipe);
- print_error(err, " clCreatePipe failed\n");
- return -1;
- }
+ test_error_ret(err, " clCreatePipe failed", -1);
- sprintf(str, "typedef struct{\n char a[%d];\n}TestStruct;\n\n__kernel void test_pipe_max_packet_size_write(__global TestStruct *src, __write_only pipe TestStruct out_pipe)\n{\n", max_pipe_packet_size);
- strcpy(source,str);
- strcat(source, " int gid = get_global_id(0);\n reserve_id_t res_id;\n\n");
- sprintf(str, " res_id = reserve_write_pipe(out_pipe, 1);\n if(is_valid_reserve_id(res_id))\n {\n");
- strcat(source, str);
- sprintf(str, " write_pipe(out_pipe, res_id, 0, &src[gid]);\n commit_write_pipe(out_pipe, res_id);\n }\n}\n\n");
- strcat(source, str);
- sprintf(str, "__kernel void test_pipe_max_packet_size_read(__read_only pipe TestStruct in_pipe, __global TestStruct *dst)\n{\n");
- strcat(source, str);
- strcat(source, " int gid = get_global_id(0);\n reserve_id_t res_id;\n\n");
- sprintf(str, " res_id = reserve_read_pipe(in_pipe, 1);\n if(is_valid_reserve_id(res_id))\n {\n");
- strcat(source, str);
- sprintf(str, " read_pipe(in_pipe, res_id, 0, &dst[gid]);\n commit_read_pipe(in_pipe, res_id);\n }\n}\n\n");
- strcat(source, str);
+ // clang-format off
+ source << R"(
+ typedef struct{
+ char a[)" << max_pipe_packet_size << R"(];
+ }TestStruct;
- str_length = strlen(source);
- assert(str_length <= STRING_LENGTH);
+ __kernel void test_pipe_max_packet_size_write(__global TestStruct *src, __write_only pipe TestStruct out_pipe)
+ {
+ int gid = get_global_id(0);
+ reserve_id_t res_id;
+
+ res_id = reserve_write_pipe(out_pipe, 1);
+ if(is_valid_reserve_id(res_id))
+ {
+ write_pipe(out_pipe, res_id, 0, &src[gid]);
+ commit_write_pipe(out_pipe, res_id);
+ }
+ }
+
+ __kernel void test_pipe_max_packet_size_read(__read_only pipe TestStruct in_pipe, __global TestStruct *dst)
+ {
+ int gid = get_global_id(0);
+ reserve_id_t res_id;
+
+ res_id = reserve_read_pipe(in_pipe, 1);
+ if(is_valid_reserve_id(res_id))
+ {
+ read_pipe(in_pipe, res_id, 0, &dst[gid]);
+ commit_read_pipe(in_pipe, res_id);
+ }
+ }
+ )";
+ // clang-format on
+
+ std::string kernel_source = source.str();
+ const char *sources[] = { kernel_source.c_str() };
// Create producer kernel
- err = create_single_kernel_helper_with_build_options(context, &program, &kernel[0], 1, (const char**)&source, kernelName[0], "-cl-std=CL2.0");
- if(err){
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- clReleaseMemObject(pipe);
- align_free(outptr);
- free(source);
- print_error(err, "Error creating program\n");
- return -1;
- }
+ err = create_single_kernel_helper_with_build_options(
+ context, &program, &kernel[0], 1, sources, kernelName[0],
+ "-cl-std=CL2.0");
+ test_error_ret(err, " Error creating program", -1);
+
//Create consumer kernel
kernel[1] = clCreateKernel(program, kernelName[1], &err);
- if( kernel[1] == NULL || err != CL_SUCCESS)
- {
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- clReleaseMemObject(pipe);
- align_free(outptr);
- free(source);
- print_error(err, "Error creating kernel\n");
- return -1;
- }
+ test_error_ret(err, " Error creating kernel", -1);
err = clSetKernelArg(kernel[0], 0, sizeof(cl_mem), (void*)&buffers[0]);
err |= clSetKernelArg(kernel[0], 1, sizeof(cl_mem), (void*)&pipe);
err |= clSetKernelArg(kernel[1], 0, sizeof(cl_mem), (void*)&pipe);
err |= clSetKernelArg(kernel[1], 1, sizeof(cl_mem), (void*)&buffers[1]);
- if ( err != CL_SUCCESS ){
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- clReleaseMemObject(pipe);
- clReleaseKernel(kernel[0]);
- clReleaseKernel(kernel[1]);
- clReleaseProgram(program);
- align_free(outptr);
- free(source);
- print_error(err, " clSetKernelArg failed");
- return -1;
- }
+ test_error_ret(err, " clSetKernelArg failed", -1);
// Launch Producer kernel
err = clEnqueueNDRangeKernel( queue, kernel[0], 1, NULL, global_work_size, NULL, 0, NULL, &producer_sync_event );
- if ( err != CL_SUCCESS ){
- print_error( err, " clEnqueueNDRangeKernel failed" );
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- clReleaseMemObject(pipe);
- clReleaseKernel(kernel[0]);
- clReleaseKernel(kernel[1]);
- clReleaseEvent(producer_sync_event);
- clReleaseEvent(consumer_sync_event);
- clReleaseProgram(program);
- align_free(outptr);
- free(source);
- return -1;
- }
+ test_error_ret(err, " clEnqueueNDRangeKernel failed", -1);
// Launch Consumer kernel
err = clEnqueueNDRangeKernel( queue, kernel[1], 1, NULL, global_work_size, NULL, 1, &producer_sync_event, &consumer_sync_event );
- if ( err != CL_SUCCESS ){
- print_error( err, " clEnqueueNDRangeKernel failed" );
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- clReleaseMemObject(pipe);
- clReleaseKernel(kernel[0]);
- clReleaseKernel(kernel[1]);
- clReleaseEvent(producer_sync_event);
- clReleaseEvent(consumer_sync_event);
- clReleaseProgram(program);
- align_free(outptr);
- free(source);
- return -1;
- }
+ test_error_ret(err, " clEnqueueNDRangeKernel failed", -1);
err = clEnqueueReadBuffer(queue, buffers[1], true, 0, size, outptr, 1, &consumer_sync_event, NULL);
- if ( err != CL_SUCCESS ){
- print_error( err, " clEnqueueReadBuffer failed" );
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- clReleaseMemObject(pipe);
- clReleaseKernel(kernel[0]);
- clReleaseKernel(kernel[1]);
- clReleaseEvent(producer_sync_event);
- clReleaseEvent(consumer_sync_event);
- clReleaseProgram(program);
- align_free(outptr);
- free(source);
- return -1;
- }
+ test_error_ret(err, " clEnqueueReadBuffer failed", -1);
if( verify_result( inptr, outptr, size)){
log_error("test_pipe_max_packet_size failed\n");
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- clReleaseMemObject(pipe);
- clReleaseKernel(kernel[0]);
- clReleaseKernel(kernel[1]);
- clReleaseEvent(producer_sync_event);
- clReleaseEvent(consumer_sync_event);
- clReleaseProgram(program);
- align_free(outptr);
- free(source);
- return -1;
}
else {
log_info("test_pipe_max_packet_size passed\n");
}
- //cleanup
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- clReleaseMemObject(pipe);
- clReleaseKernel(kernel[0]);
- clReleaseKernel(kernel[1]);
- clReleaseEvent(producer_sync_event);
- clReleaseEvent(consumer_sync_event);
- clReleaseProgram(program);
- align_free(outptr);
- free(source);
return 0;
}
int test_pipe_max_active_reservations(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
{
- cl_mem pipe;
- cl_mem buffers[2];
- cl_mem buf_reservations;
- cl_mem buf_status;
- cl_mem buf_reserve_id_t_size;
- cl_mem buf_reserve_id_t_size_aligned;
- cl_int *inptr;
- void *outptr;
- int size, i;
- cl_program program;
- cl_kernel kernel[3];
- size_t global_work_size[3];
- cl_int err;
- int status = 0;
- cl_uint max_active_reservations = 0;
- cl_ulong max_global_size = 0;
- int reserve_id_t_size;
- int temp;
- char *source;
- char str[256];
- int str_length;
- cl_event sync_event = NULL;
- cl_event read_event = NULL;
- MTdata d = init_genrand( gRandomSeed );
- const char* kernelName[3] = {"test_pipe_max_active_reservations_write", "test_pipe_max_active_reservations_read", "pipe_get_reserve_id_t_size"};
+ clMemWrapper pipe;
+ clMemWrapper buffers[2];
+ clMemWrapper buf_reservations;
+ clMemWrapper buf_status;
+ clMemWrapper buf_reserve_id_t_size;
+ clMemWrapper buf_reserve_id_t_size_aligned;
+ cl_int *inptr;
+ void *outptr;
+ int size, i;
+ clProgramWrapper program;
+ clKernelWrapper kernel[3];
+ size_t global_work_size[3];
+ cl_int err;
+ int status = 0;
+ cl_uint max_active_reservations = 0;
+ cl_ulong max_global_size = 0;
+ int reserve_id_t_size;
+ int temp;
+ clEventWrapper sync_event = NULL;
+ clEventWrapper read_event = NULL;
+ BufferOwningPtr<cl_int> BufferInPtr;
+ BufferOwningPtr<cl_int> BufferOutPtr;
+ MTdataHolder d(gRandomSeed);
+ const char *kernelName[3] = { "test_pipe_max_active_reservations_write",
+ "test_pipe_max_active_reservations_read",
+ "pipe_get_reserve_id_t_size" };
- size_t min_alignment = get_min_alignment(context);
+ size_t min_alignment = get_min_alignment(context);
- source = (char*)malloc(2*STRING_LENGTH*sizeof(char));
+ std::stringstream source;
global_work_size[0] = 1;
- err = clGetDeviceInfo(deviceID, CL_DEVICE_PIPE_MAX_ACTIVE_RESERVATIONS, sizeof(max_active_reservations), (void*)&max_active_reservations, NULL);
- if(err){
- print_error(err, " clGetDeviceInfo failed\n");
- return -1;
- }
+ err = clGetDeviceInfo(deviceID, CL_DEVICE_PIPE_MAX_ACTIVE_RESERVATIONS,
+ sizeof(max_active_reservations),
+ (void *)&max_active_reservations, NULL);
+ test_error_ret(err, " clGetDeviceInfo failed", -1);
- err = clGetDeviceInfo(deviceID, CL_DEVICE_GLOBAL_MEM_SIZE, sizeof(max_global_size), (void*)&max_global_size, NULL);
- if(err){
- print_error(err, " clGetDeviceInfo failed\n");
- return -1;
- }
+ err = clGetDeviceInfo(deviceID, CL_DEVICE_GLOBAL_MEM_SIZE,
+ sizeof(max_global_size), (void *)&max_global_size,
+ NULL);
+ test_error_ret(err, " clGetDeviceInfo failed", -1);
- max_active_reservations = (max_active_reservations > max_global_size) ? 1<<16 : max_active_reservations;
+ max_active_reservations = (max_active_reservations > max_global_size)
+ ? 1 << 16
+ : max_active_reservations;
- if(max_active_reservations < 1){
+ if (max_active_reservations < 1)
+ {
log_error("The device should support minimum active reservations of 1");
return -1;
}
// To get reserve_id_t size
buf_reserve_id_t_size = clCreateBuffer(context, CL_MEM_HOST_READ_ONLY, sizeof(reserve_id_t_size), NULL, &err);
- if ( err ){
- clReleaseMemObject(buf_reserve_id_t_size);
- print_error(err, " clCreateBuffer failed\n" );
- return -1;
- }
+ test_error_ret(err, " clCreateBuffer failed", -1);
- sprintf(str, "__kernel void test_pipe_max_active_reservations_write(__global int *src, __write_only pipe int out_pipe, __global char *reserve_id, __global int *reserve_id_t_size_aligned, __global int *status)\n{\n");
- strcpy(source,str);
- sprintf(str, " __global reserve_id_t *res_id_ptr;\n int reserve_idx;\n int commit_idx;\n");
- strcat(source, str);
- sprintf(str, " for(reserve_idx = 0; reserve_idx < %d; reserve_idx++)\n {\n", max_active_reservations);
- strcat(source, str);
- sprintf(str, " res_id_ptr = (__global reserve_id_t*)(reserve_id + reserve_idx*reserve_id_t_size_aligned[0]);\n");
- strcat(source, str);
- sprintf(str, " *res_id_ptr = reserve_write_pipe(out_pipe, 1);\n");
- strcat(source, str);
- sprintf(str, " if(is_valid_reserve_id(res_id_ptr[0]))\n {\n write_pipe(out_pipe, res_id_ptr[0], 0, &src[reserve_idx]);\n }\n");
- strcat(source, str);
- sprintf(str, " else\n {\n *status = -1;\n return;\n }\n }\n");
- strcat(source, str);
- sprintf(str, " for(commit_idx = 0; commit_idx < %d; commit_idx++)\n {\n", max_active_reservations);
- strcat(source, str);
- sprintf(str, " res_id_ptr = (__global reserve_id_t*)(reserve_id + commit_idx*reserve_id_t_size_aligned[0]);\n");
- strcat(source, str);
- sprintf(str, " commit_write_pipe(out_pipe, res_id_ptr[0]);\n }\n}\n\n");
- strcat(source, str);
- sprintf(str, "__kernel void test_pipe_max_active_reservations_read(__read_only pipe int in_pipe, __global int *dst, __global char *reserve_id, __global int *reserve_id_t_size_aligned, __global int *status)\n{\n");
- strcat(source, str);
- sprintf(str, " __global reserve_id_t *res_id_ptr;\n int reserve_idx;\n int commit_idx;\n");
- strcat(source, str);
- sprintf(str, " for(reserve_idx = 0; reserve_idx < %d; reserve_idx++)\n {\n", max_active_reservations);
- strcat(source, str);
- sprintf(str, " res_id_ptr = (__global reserve_id_t*)(reserve_id + reserve_idx*reserve_id_t_size_aligned[0]);\n");
- strcat(source, str);
- sprintf(str, " *res_id_ptr = reserve_read_pipe(in_pipe, 1);\n");
- strcat(source, str);
- sprintf(str, " if(is_valid_reserve_id(res_id_ptr[0]))\n {\n read_pipe(in_pipe, res_id_ptr[0], 0, &dst[reserve_idx]);\n }\n");
- strcat(source, str);
- sprintf(str, " else\n {\n *status = -1;\n return;\n }\n }\n");
- strcat(source, str);
- sprintf(str, " for(commit_idx = 0; commit_idx < %d; commit_idx++)\n {\n", max_active_reservations);
- strcat(source, str);
- sprintf(str, " res_id_ptr = (__global reserve_id_t*)(reserve_id + commit_idx*reserve_id_t_size_aligned[0]);\n");
- strcat(source, str);
- sprintf(str, " commit_read_pipe(in_pipe, res_id_ptr[0]);\n }\n}\n\n");
- strcat(source, str);
- sprintf(str, "__kernel void pipe_get_reserve_id_t_size(__global int *reserve_id_t_size) \n");
- strcat(source, str);
- sprintf(str, "{\n *reserve_id_t_size = sizeof(reserve_id_t);\n}\n");
- strcat(source, str);
+ // clang-format off
+ source << R"(
+ __kernel void test_pipe_max_active_reservations_write(__global int *src, __write_only pipe int out_pipe, __global char *reserve_id, __global int *reserve_id_t_size_aligned, __global int *status)
+ {
+ __global reserve_id_t *res_id_ptr;
+ int reserve_idx;
+ int commit_idx;
- str_length = strlen(source);
- assert(str_length <= 2*STRING_LENGTH);
+ for(reserve_idx = 0; reserve_idx < )" << max_active_reservations << R"(; reserve_idx++)
+ {
+ res_id_ptr = (__global reserve_id_t*)(reserve_id + reserve_idx*reserve_id_t_size_aligned[0]);
+ *res_id_ptr = reserve_write_pipe(out_pipe, 1);
+ if(is_valid_reserve_id(res_id_ptr[0]))
+ {
+ write_pipe(out_pipe, res_id_ptr[0], 0, &src[reserve_idx]);
+ }
+ else
+ {
+ *status = -1;
+ return;
+ }
+ }
+
+ for(commit_idx = 0; commit_idx < )" << max_active_reservations << R"(; commit_idx++)
+ {
+ res_id_ptr = (__global reserve_id_t*)(reserve_id + commit_idx*reserve_id_t_size_aligned[0]);
+ commit_write_pipe(out_pipe, res_id_ptr[0]);
+ }
+ }
+
+ __kernel void test_pipe_max_active_reservations_read(__read_only pipe int in_pipe, __global int *dst, __global char *reserve_id, __global int *reserve_id_t_size_aligned, __global int *status)
+ {
+ __global reserve_id_t *res_id_ptr;
+ int reserve_idx;
+ int commit_idx;
+
+ for(reserve_idx = 0; reserve_idx < )" << max_active_reservations << R"(; reserve_idx++)
+ {
+ res_id_ptr = (__global reserve_id_t*)(reserve_id + reserve_idx*reserve_id_t_size_aligned[0]);
+ *res_id_ptr = reserve_read_pipe(in_pipe, 1);
+
+ if(is_valid_reserve_id(res_id_ptr[0]))
+ {
+ read_pipe(in_pipe, res_id_ptr[0], 0, &dst[reserve_idx]);
+ }
+ else
+ {
+ *status = -1;
+ return;
+ }
+ }
+
+ for(commit_idx = 0; commit_idx < )" << max_active_reservations << R"(; commit_idx++)
+ {
+ res_id_ptr = (__global reserve_id_t*)(reserve_id + commit_idx*reserve_id_t_size_aligned[0]);
+ commit_read_pipe(in_pipe, res_id_ptr[0]);
+ }
+ }
+
+ __kernel void pipe_get_reserve_id_t_size(__global int *reserve_id_t_size)
+ {
+ *reserve_id_t_size = sizeof(reserve_id_t);
+ }
+ )";
+ // clang-format on
+
+ std::string kernel_source = source.str();
+ const char *sources[] = { kernel_source.c_str() };
// Create producer kernel
- err = create_single_kernel_helper_with_build_options(context, &program, &kernel[0], 1, (const char**)&source, kernelName[0], "-cl-std=CL2.0");
- if(err){
- clReleaseMemObject(buf_reserve_id_t_size);
- print_error(err, "Error creating program\n");
- return -1;
- }
+ err = create_single_kernel_helper_with_build_options(
+ context, &program, &kernel[0], 1, sources, kernelName[0],
+ "-cl-std=CL2.0");
+ test_error_ret(err, " Error creating program", -1);
// Create consumer kernel
kernel[1] = clCreateKernel(program, kernelName[1], &err);
- if( kernel[1] == NULL || err != CL_SUCCESS)
- {
- clReleaseMemObject(buf_reserve_id_t_size);
- print_error(err, "Error creating kernel\n");
- return -1;
- }
+ test_error_ret(err, " Error creating kernel", -1);
// Create size query kernel for reserve_id_t
kernel[2] = clCreateKernel(program, kernelName[2], &err);
- if( kernel[2] == NULL || err != CL_SUCCESS)
- {
- clReleaseMemObject(buf_reserve_id_t_size);
- print_error(err, "Error creating kernel\n");
- return -1;
- }
+ test_error_ret(err, " Error creating kernel", -1);
+
err = clSetKernelArg(kernel[2], 0, sizeof(cl_mem), (void*)&buf_reserve_id_t_size);
- if(err){
- clReleaseMemObject(buf_reserve_id_t_size);
- print_error(err, "Error creating program\n");
- return -1;
- }
+ test_error_ret(err, " clSetKernelArg failed", -1);
+
//Launch size query kernel for reserve_id_t
err = clEnqueueNDRangeKernel( queue, kernel[2], 1, NULL, global_work_size, NULL, 0, NULL, &sync_event );
- if ( err != CL_SUCCESS ){
- print_error( err, " clEnqueueNDRangeKernel failed" );
- clReleaseMemObject(buf_reserve_id_t_size);
- clReleaseKernel(kernel[0]);
- clReleaseKernel(kernel[1]);
- clReleaseKernel(kernel[2]);
- clReleaseEvent(sync_event);
- clReleaseProgram(program);
- return -1;
- }
+ test_error_ret(err, " clEnqueueNDRangeKernel failed", -1);
err = clEnqueueReadBuffer(queue, buf_reserve_id_t_size, true, 0, sizeof(reserve_id_t_size), &reserve_id_t_size, 1, &sync_event, &read_event);
- if ( err != CL_SUCCESS ){
- print_error( err, " clEnqueueReadBuffer failed" );
- clReleaseMemObject(buf_reserve_id_t_size);
- clReleaseKernel(kernel[0]);
- clReleaseKernel(kernel[1]);
- clReleaseKernel(kernel[2]);
- clReleaseEvent(sync_event);
- clReleaseProgram(program);
- return -1;
- }
+ test_error_ret(err, " clEnqueueReadBuffer failed", -1);
err = clWaitForEvents(1, &read_event);
- if ( err != CL_SUCCESS ){
- print_error( err, " clWaitForEvents failed" );
- clReleaseMemObject(buf_reserve_id_t_size);
- clReleaseKernel(kernel[0]);
- clReleaseKernel(kernel[1]);
- clReleaseKernel(kernel[2]);
- clReleaseEvent(sync_event);
- clReleaseEvent(read_event);
- clReleaseProgram(program);
- return -1;
- }
+ test_error_ret(err, " clWaitForEvents failed", -1);
// Round reserve_id_t_size to the nearest power of 2
temp = 1;
@@ -763,323 +571,84 @@
for(i = 0; i < max_active_reservations; i++){
inptr[i] = (int)genrand_int32(d);
}
+ BufferInPtr.reset(inptr, nullptr, 0, size, true);
buffers[0] = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR, size, inptr, &err);
- if ( err ){
- clReleaseMemObject(buf_reserve_id_t_size);
- clReleaseMemObject(buffers[0]);
- clReleaseKernel(kernel[0]);
- clReleaseKernel(kernel[1]);
- clReleaseKernel(kernel[2]);
- clReleaseEvent(sync_event);
- clReleaseEvent(read_event);
- clReleaseProgram(program);
- print_error(err, " clCreateBuffer failed\n" );
- return -1;
- }
+ test_error_ret(err, " clCreateBuffer failed", -1);
outptr = align_malloc(size, min_alignment);
+ BufferOutPtr.reset(outptr, nullptr, 0, size, true);
+
buffers[1] = clCreateBuffer(context, CL_MEM_HOST_READ_ONLY, size, NULL, &err);
- if ( err ){
- clReleaseMemObject(buf_reserve_id_t_size);
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- clReleaseKernel(kernel[0]);
- clReleaseKernel(kernel[1]);
- clReleaseKernel(kernel[2]);
- clReleaseEvent(sync_event);
- clReleaseEvent(read_event);
- clReleaseProgram(program);
- align_free(outptr);
- print_error(err, " clCreateBuffer failed\n" );
- return -1;
- }
+ test_error_ret(err, " clCreateBuffer failed", -1);
buf_reserve_id_t_size_aligned = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR, sizeof(reserve_id_t_size), &reserve_id_t_size, &err);
- if ( err ){
- clReleaseMemObject(buf_reserve_id_t_size);
- clReleaseMemObject(buf_reserve_id_t_size_aligned);
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- clReleaseKernel(kernel[0]);
- clReleaseKernel(kernel[1]);
- clReleaseKernel(kernel[2]);
- clReleaseEvent(sync_event);
- clReleaseEvent(read_event);
- clReleaseProgram(program);
- print_error(err, " clCreateBuffer failed\n" );
- return -1;
- }
+ test_error_ret(err, " clCreateBuffer failed", -1);
//For error status
buf_status = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR, sizeof(int), &status, &err);
- if ( err ){
- clReleaseMemObject(buf_status);
- clReleaseMemObject(buf_reserve_id_t_size);
- clReleaseMemObject(buf_reserve_id_t_size_aligned);
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- align_free(outptr);
- clReleaseKernel(kernel[0]);
- clReleaseKernel(kernel[1]);
- clReleaseKernel(kernel[2]);
- clReleaseEvent(sync_event);
- clReleaseEvent(read_event);
- clReleaseProgram(program);
- print_error(err, " clCreateBuffer failed\n" );
- return -1;
- }
+ test_error_ret(err, " clCreateBuffer failed", -1);
pipe = clCreatePipe(context, CL_MEM_HOST_NO_ACCESS, sizeof(int), max_active_reservations, NULL, &err);
- if(err){
- clReleaseMemObject(buf_status);
- clReleaseMemObject(buf_reserve_id_t_size);
- clReleaseMemObject(buf_reserve_id_t_size_aligned);
- clReleaseMemObject(pipe);
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- align_free(outptr);
- clReleaseKernel(kernel[0]);
- clReleaseKernel(kernel[1]);
- clReleaseKernel(kernel[2]);
- clReleaseEvent(sync_event);
- clReleaseEvent(read_event);
- clReleaseProgram(program);
- print_error(err, " clCreatePipe failed\n");
- return -1;
- }
+ test_error_ret(err, " clCreatePipe failed", -1);
// Global buffer to hold all active reservation ids
buf_reservations = clCreateBuffer(context, CL_MEM_HOST_NO_ACCESS, reserve_id_t_size*max_active_reservations, NULL, &err);
- if ( err != CL_SUCCESS ){
- print_error( err, " clCreateBuffer failed" );
- clReleaseMemObject(buf_status);
- clReleaseMemObject(buf_reserve_id_t_size);
- clReleaseMemObject(buf_reserve_id_t_size_aligned);
- clReleaseMemObject(buf_reservations);
- clReleaseMemObject(pipe);
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- align_free(outptr);
- clReleaseKernel(kernel[0]);
- clReleaseKernel(kernel[1]);
- clReleaseKernel(kernel[2]);
- clReleaseEvent(sync_event);
- clReleaseEvent(read_event);
- clReleaseProgram(program);
- return -1;
- }
+ test_error_ret(err, " clCreateBuffer failed", -1);
err = clSetKernelArg(kernel[0], 0, sizeof(cl_mem), (void*)&buffers[0]);
err |= clSetKernelArg(kernel[0], 1, sizeof(cl_mem), (void*)&pipe);
err |= clSetKernelArg(kernel[0], 2, sizeof(cl_mem), (void*)&buf_reservations);
err |= clSetKernelArg(kernel[0], 3, sizeof(cl_mem), (void*)&buf_reserve_id_t_size_aligned);
err |= clSetKernelArg(kernel[0], 4, sizeof(cl_mem), (void*)&buf_status);
- if ( err != CL_SUCCESS ){
- clReleaseMemObject(buf_status);
- clReleaseMemObject(buf_reserve_id_t_size);
- clReleaseMemObject(buf_reserve_id_t_size_aligned);
- clReleaseMemObject(buf_reservations);
- clReleaseMemObject(pipe);
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- align_free(outptr);
- clReleaseKernel(kernel[0]);
- clReleaseKernel(kernel[1]);
- clReleaseKernel(kernel[2]);
- clReleaseProgram(program);
- print_error(err, " clSetKernelArg failed");
- return -1;
- }
+ test_error_ret(err, " clSetKernelArg failed", -1);
err = clSetKernelArg(kernel[1], 0, sizeof(cl_mem), (void*)&pipe);
err |= clSetKernelArg(kernel[1], 1, sizeof(cl_mem), (void*)&buffers[1]);
err |= clSetKernelArg(kernel[1], 2, sizeof(cl_mem), (void*)&buf_reservations);
err |= clSetKernelArg(kernel[1], 3, sizeof(cl_mem), (void*)&buf_reserve_id_t_size_aligned);
err |= clSetKernelArg(kernel[1], 4, sizeof(cl_mem), (void*)&buf_status);
- if ( err != CL_SUCCESS ){
- clReleaseMemObject(buf_status);
- clReleaseMemObject(buf_reserve_id_t_size);
- clReleaseMemObject(buf_reserve_id_t_size_aligned);
- clReleaseMemObject(buf_reservations);
- clReleaseMemObject(pipe);
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- align_free(outptr);
- clReleaseKernel(kernel[0]);
- clReleaseKernel(kernel[1]);
- clReleaseKernel(kernel[2]);
- clReleaseProgram(program);
- print_error(err, " clSetKernelArg failed");
- return -1;
- }
+ test_error_ret(err, " clSetKernelArg failed", -1);
+
+ clReleaseEvent(sync_event);
// Launch Producer kernel
err = clEnqueueNDRangeKernel(queue, kernel[0], 1, NULL, global_work_size, NULL, 0, NULL, &sync_event);
- if ( err != CL_SUCCESS ){
- print_error( err, " clEnqueueNDRangeKernel failed" );
- clReleaseMemObject(buf_status);
- clReleaseMemObject(buf_reserve_id_t_size);
- clReleaseMemObject(buf_reserve_id_t_size_aligned);
- clReleaseMemObject(buf_reservations);
- clReleaseMemObject(pipe);
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- align_free(outptr);
- clReleaseKernel(kernel[0]);
- clReleaseKernel(kernel[1]);
- clReleaseKernel(kernel[2]);
- clReleaseEvent(sync_event);
- clReleaseProgram(program);
- return -1;
- }
+ test_error_ret(err, " clEnqueueNDRangeKernel failed", -1);
err = clEnqueueReadBuffer(queue, buf_status, true, 0, sizeof(int), &status, 1, &sync_event, NULL);
- if ( err != CL_SUCCESS ){
- print_error( err, " clEnqueueReadBuffer failed" );
- clReleaseMemObject(buf_status);
- clReleaseMemObject(buf_reserve_id_t_size);
- clReleaseMemObject(buf_reserve_id_t_size_aligned);
- clReleaseMemObject(buf_reservations);
- clReleaseMemObject(pipe);
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- align_free(outptr);
- clReleaseKernel(kernel[0]);
- clReleaseKernel(kernel[1]);
- clReleaseKernel(kernel[2]);
- clReleaseEvent(sync_event);
- clReleaseProgram(program);
- return -1;
- }
+ test_error_ret(err, " clEnqueueReadBuffer failed", -1);
if(status != 0)
{
log_error("test_pipe_max_active_reservations failed\n");
- clReleaseMemObject(buf_status);
- clReleaseMemObject(buf_reserve_id_t_size);
- clReleaseMemObject(buf_reserve_id_t_size_aligned);
- clReleaseMemObject(buf_reservations);
- clReleaseMemObject(pipe);
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- align_free(outptr);
- clReleaseKernel(kernel[0]);
- clReleaseKernel(kernel[1]);
- clReleaseEvent(sync_event);
- clReleaseProgram(program);
return -1;
}
+ clReleaseEvent(sync_event);
// Launch Consumer kernel
err = clEnqueueNDRangeKernel(queue, kernel[1], 1, NULL, global_work_size, NULL, 0, NULL, &sync_event);
- if ( err != CL_SUCCESS ){
- print_error( err, " clEnqueueNDRangeKernel failed" );
- clReleaseMemObject(buf_status);
- clReleaseMemObject(buf_reserve_id_t_size);
- clReleaseMemObject(buf_reserve_id_t_size_aligned);
- clReleaseMemObject(buf_reservations);
- clReleaseMemObject(pipe);
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- align_free(outptr);
- clReleaseKernel(kernel[0]);
- clReleaseKernel(kernel[1]);
- clReleaseKernel(kernel[2]);
- clReleaseEvent(sync_event);
- clReleaseProgram(program);
- return -1;
- }
+ test_error_ret(err, " clEnqueueNDRangeKernel failed", -1);
err = clEnqueueReadBuffer(queue, buf_status, true, 0, sizeof(int), &status, 1, &sync_event, NULL);
- if ( err != CL_SUCCESS ){
- print_error( err, " clEnqueueReadBuffer failed" );
- clReleaseMemObject(buf_status);
- clReleaseMemObject(buf_reserve_id_t_size);
- clReleaseMemObject(buf_reserve_id_t_size_aligned);
- clReleaseMemObject(buf_reservations);
- clReleaseMemObject(pipe);
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- align_free(outptr);
- clReleaseKernel(kernel[0]);
- clReleaseKernel(kernel[1]);
- clReleaseKernel(kernel[2]);
- clReleaseEvent(sync_event);
- clReleaseProgram(program);
- return -1;
- }
+ test_error_ret(err, " clEnqueueReadBuffer failed", -1);
if(status != 0)
{
log_error("test_pipe_max_active_reservations failed\n");
- clReleaseMemObject(buf_status);
- clReleaseMemObject(buf_reserve_id_t_size);
- clReleaseMemObject(buf_reserve_id_t_size_aligned);
- clReleaseMemObject(buf_reservations);
- clReleaseMemObject(pipe);
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- align_free(outptr);
- clReleaseKernel(kernel[0]);
- clReleaseKernel(kernel[1]);
- clReleaseEvent(sync_event);
- clReleaseProgram(program);
return -1;
}
err = clEnqueueReadBuffer(queue, buffers[1], true, 0, size, outptr, 1, &sync_event, NULL);
- if ( err != CL_SUCCESS ){
- print_error( err, " clEnqueueReadBuffer failed" );
- clReleaseMemObject(buf_status);
- clReleaseMemObject(buf_reserve_id_t_size);
- clReleaseMemObject(buf_reserve_id_t_size_aligned);
- clReleaseMemObject(buf_reservations);
- clReleaseMemObject(pipe);
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- align_free(outptr);
- clReleaseKernel(kernel[0]);
- clReleaseKernel(kernel[1]);
- clReleaseKernel(kernel[2]);
- clReleaseEvent(sync_event);
- clReleaseProgram(program);
- return -1;
- }
+ test_error_ret(err, " clEnqueueReadBuffer failed", -1);
if( verify_result_int( inptr, outptr, max_active_reservations)){
log_error("test_pipe_max_active_reservations failed\n");
- clReleaseMemObject(buf_status);
- clReleaseMemObject(buf_reserve_id_t_size);
- clReleaseMemObject(buf_reserve_id_t_size_aligned);
- clReleaseMemObject(buf_reservations);
- clReleaseMemObject(pipe);
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- align_free(outptr);
- clReleaseKernel(kernel[0]);
- clReleaseKernel(kernel[1]);
- clReleaseKernel(kernel[2]);
- clReleaseEvent(sync_event);
- clReleaseProgram(program);
return -1;
}
else {
log_info("test_pipe_max_active_reservations passed\n");
}
- //cleanup
- clReleaseMemObject(buf_status);
- clReleaseMemObject(buf_reserve_id_t_size);
- clReleaseMemObject(buf_reserve_id_t_size_aligned);
- clReleaseMemObject(buf_reservations);
- clReleaseMemObject(pipe);
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- align_free(outptr);
- clReleaseKernel(kernel[0]);
- clReleaseKernel(kernel[1]);
- clReleaseEvent(sync_event);
- clReleaseProgram(program);
return 0;
}
\ No newline at end of file
diff --git a/test_conformance/pipes/test_pipe_query_functions.cpp b/test_conformance/pipes/test_pipe_query_functions.cpp
index c1b4d92..f9c93aa 100644
--- a/test_conformance/pipes/test_pipe_query_functions.cpp
+++ b/test_conformance/pipes/test_pipe_query_functions.cpp
@@ -79,28 +79,32 @@
int test_pipe_query_functions(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
{
- cl_mem pipe;
- cl_mem buffers[4];
- void *outptr1;
- void *outptr2;
- cl_int *inptr;
- cl_program program;
- cl_kernel kernel[3];
- size_t global_work_size[3];
- size_t half_global_work_size[3];
- size_t global_work_size_pipe_query[3];
- cl_int pipe_max_packets, pipe_num_packets;
- cl_int err;
- cl_int size;
- cl_int i;
- cl_event producer_sync_event = NULL;
- cl_event consumer_sync_event = NULL;
- cl_event pipe_query_sync_event = NULL;
- cl_event pipe_read_sync_event = NULL;
- MTdata d = init_genrand( gRandomSeed );
- const char* kernelName[] = {"test_pipe_write", "test_pipe_read", "test_pipe_query_functions"};
+ clMemWrapper pipe;
+ clMemWrapper buffers[4];
+ void *outptr1;
+ void *outptr2;
+ cl_int *inptr;
+ clProgramWrapper program;
+ clKernelWrapper kernel[3];
+ size_t global_work_size[3];
+ size_t half_global_work_size[3];
+ size_t global_work_size_pipe_query[3];
+ cl_int pipe_max_packets, pipe_num_packets;
+ cl_int err;
+ cl_int size;
+ cl_int i;
+ clEventWrapper producer_sync_event = NULL;
+ clEventWrapper consumer_sync_event = NULL;
+ clEventWrapper pipe_query_sync_event = NULL;
+ clEventWrapper pipe_read_sync_event = NULL;
+ BufferOwningPtr<cl_int> BufferInPtr;
+ BufferOwningPtr<cl_int> BufferOutPtr1;
+ BufferOwningPtr<cl_int> BufferOutPtr2;
+ MTdataHolder d(gRandomSeed);
+ const char *kernelName[] = { "test_pipe_write", "test_pipe_read",
+ "test_pipe_query_functions" };
- size_t min_alignment = get_min_alignment(context);
+ size_t min_alignment = get_min_alignment(context);
size = sizeof(int) * num_elements;
global_work_size[0] = (cl_uint)num_elements;
@@ -109,98 +113,43 @@
inptr = (int *)align_malloc(size, min_alignment);
- for(i = 0; i < num_elements; i++){
+ for (i = 0; i < num_elements; i++)
+ {
inptr[i] = TEST_PRIME_INT;
}
+ BufferInPtr.reset(inptr, nullptr, 0, size, true);
buffers[0] = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR, size, inptr, &err);
- if(err){
- clReleaseMemObject(buffers[0]);
- print_error(err, " clCreateBuffer failed\n");
- return -1;
- }
+ test_error_ret(err, " clCreateBuffer failed", -1);
+
outptr1 = align_malloc(size/2, min_alignment);
outptr2 = align_malloc(size, min_alignment);
+ BufferOutPtr1.reset(outptr1, nullptr, 0, size, true);
+ BufferOutPtr2.reset(outptr2, nullptr, 0, size, true);
+
buffers[1] = clCreateBuffer(context, CL_MEM_HOST_READ_ONLY, size, NULL, &err);
- if ( err ){
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- align_free( outptr1 );
- print_error(err, " clCreateBuffer failed\n" );
- return -1;
- }
+ test_error_ret(err, " clCreateBuffer failed", -1);
buffers[2] = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(int), NULL, &err);
- if ( err ){
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- clReleaseMemObject(buffers[2]);
- align_free( outptr1 );
- print_error(err, " clCreateBuffer failed\n" );
- return -1;
- }
+ test_error_ret(err, " clCreateBuffer failed", -1);
buffers[3] = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(int), NULL, &err);
- if ( err ){
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- clReleaseMemObject(buffers[2]);
- clReleaseMemObject(buffers[3]);
- align_free( outptr1 );
- print_error(err, " clCreateBuffer failed\n" );
- return -1;
- }
+ test_error_ret(err, " clCreateBuffer failed", -1);
pipe = clCreatePipe(context, CL_MEM_HOST_NO_ACCESS, sizeof(int), num_elements, NULL, &err);
- if(err){
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- clReleaseMemObject(buffers[2]);
- clReleaseMemObject(buffers[3]);
- align_free( outptr1 );
- clReleaseMemObject(pipe);
- print_error(err, " clCreatePipe failed\n");
- return -1;
- }
+ test_error_ret(err, " clCreatePipe failed", -1);
// Create producer kernel
err = create_single_kernel_helper_with_build_options(context, &program, &kernel[0], 1, (const char**)&pipe_query_functions_kernel_code, kernelName[0], "-cl-std=CL2.0");
- if(err){
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- clReleaseMemObject(buffers[2]);
- clReleaseMemObject(buffers[3]);
- clReleaseMemObject(pipe);
- align_free(outptr1);
- print_error(err, "Error creating program\n");
- return -1;
- }
+ test_error_ret(err, " Error creating program", -1);
+
//Create consumer kernel
kernel[1] = clCreateKernel(program, kernelName[1], &err);
- if( kernel[1] == NULL || err != CL_SUCCESS)
- {
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- clReleaseMemObject(buffers[2]);
- clReleaseMemObject(buffers[3]);
- clReleaseMemObject(pipe);
- align_free(outptr1);
- print_error(err, "Error creating kernel\n");
- return -1;
- }
+ test_error_ret(err, " Error creating kernel", -1);
+
//Create pipe query functions kernel
kernel[2] = clCreateKernel(program, kernelName[2], &err);
- if( kernel[1] == NULL || err != CL_SUCCESS)
- {
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- clReleaseMemObject(buffers[2]);
- clReleaseMemObject(buffers[3]);
- clReleaseMemObject(pipe);
- align_free(outptr1);
- print_error(err, "Error creating kernel\n");
- return -1;
- }
+ test_error_ret(err, " Error creating kernel", -1);
err = clSetKernelArg(kernel[0], 0, sizeof(cl_mem), (void*)&buffers[0]);
err |= clSetKernelArg(kernel[0], 1, sizeof(cl_mem), (void*)&pipe);
@@ -209,104 +158,21 @@
err |= clSetKernelArg(kernel[2], 0, sizeof(cl_mem), (void*)&pipe);
err |= clSetKernelArg(kernel[2], 1, sizeof(cl_mem), (void*)&buffers[2]);
err |= clSetKernelArg(kernel[2], 2, sizeof(cl_mem), (void*)&buffers[3]);
-
- if ( err != CL_SUCCESS ){
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- clReleaseMemObject(buffers[2]);
- clReleaseMemObject(buffers[3]);
- clReleaseMemObject(pipe);
- clReleaseKernel(kernel[0]);
- clReleaseKernel(kernel[1]);
- clReleaseKernel(kernel[2]);
- clReleaseProgram(program);
- align_free(outptr1);
- print_error(err, " clSetKernelArg failed\n");
- return -1;
- }
+ test_error_ret(err, " clSetKernelArg failed", -1);
// Launch Producer kernel
err = clEnqueueNDRangeKernel( queue, kernel[0], 1, NULL, global_work_size, NULL, 0, NULL, &producer_sync_event );
- if ( err != CL_SUCCESS ){
- print_error( err, " clEnqueueNDRangeKernel failed\n" );
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- clReleaseMemObject(buffers[2]);
- clReleaseMemObject(buffers[3]);
- clReleaseMemObject(pipe);
- clReleaseKernel(kernel[0]);
- clReleaseKernel(kernel[1]);
- clReleaseKernel(kernel[2]);
- clReleaseEvent(producer_sync_event);
- clReleaseEvent(consumer_sync_event);
- clReleaseEvent(pipe_query_sync_event);
- clReleaseEvent(pipe_read_sync_event);
- clReleaseProgram(program);
- align_free(outptr1);
- return -1;
- }
+ test_error_ret(err, " clEnqueueNDRangeKernel failed", -1);
// Launch Pipe query kernel
err = clEnqueueNDRangeKernel( queue, kernel[2], 1, NULL, global_work_size_pipe_query, NULL, 1, &producer_sync_event, &pipe_query_sync_event );
- if ( err != CL_SUCCESS ){
- print_error( err, " clEnqueueNDRangeKernel failed\n" );
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- clReleaseMemObject(buffers[2]);
- clReleaseMemObject(buffers[3]);
- clReleaseMemObject(pipe);
- clReleaseKernel(kernel[0]);
- clReleaseKernel(kernel[1]);
- clReleaseKernel(kernel[2]);
- clReleaseEvent(producer_sync_event);
- clReleaseEvent(consumer_sync_event);
- clReleaseEvent(pipe_query_sync_event);
- clReleaseEvent(pipe_read_sync_event);
- clReleaseProgram(program);
- align_free(outptr1);
- return -1;
- }
+ test_error_ret(err, " clEnqueueNDRangeKernel failed", -1);
err = clEnqueueReadBuffer(queue, buffers[2], true, 0, sizeof(cl_int), &pipe_num_packets, 1, &pipe_query_sync_event, NULL);
- if ( err != CL_SUCCESS ){
- print_error( err, " clEnqueueReadBuffer failed\n" );
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- clReleaseMemObject(buffers[2]);
- clReleaseMemObject(buffers[3]);
- clReleaseMemObject(pipe);
- clReleaseKernel(kernel[0]);
- clReleaseKernel(kernel[1]);
- clReleaseKernel(kernel[2]);
- clReleaseEvent(producer_sync_event);
- clReleaseEvent(consumer_sync_event);
- clReleaseEvent(pipe_query_sync_event);
- clReleaseEvent(pipe_read_sync_event);
- clReleaseProgram(program);
- align_free(outptr1);
- return -1;
- }
-
+ test_error_ret(err, " clEnqueueReadBuffer failed", -1);
err = clEnqueueReadBuffer(queue, buffers[3], true, 0, sizeof(cl_int), &pipe_max_packets, 1, &pipe_query_sync_event, NULL);
- if ( err != CL_SUCCESS ){
- print_error( err, " clEnqueueReadBuffer failed\n" );
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- clReleaseMemObject(buffers[2]);
- clReleaseMemObject(buffers[3]);
- clReleaseMemObject(pipe);
- clReleaseKernel(kernel[0]);
- clReleaseKernel(kernel[1]);
- clReleaseKernel(kernel[2]);
- clReleaseEvent(producer_sync_event);
- clReleaseEvent(consumer_sync_event);
- clReleaseEvent(pipe_query_sync_event);
- clReleaseEvent(pipe_read_sync_event);
- clReleaseProgram(program);
- align_free(outptr1);
- return -1;
- }
+ test_error_ret(err, " clEnqueueReadBuffer failed", -1);
if(pipe_num_packets != num_elements || pipe_max_packets != num_elements)
{
@@ -316,85 +182,20 @@
// Launch Consumer kernel with half the previous global size
err = clEnqueueNDRangeKernel( queue, kernel[1], 1, NULL, half_global_work_size, NULL, 1, &producer_sync_event, &consumer_sync_event );
- if ( err != CL_SUCCESS ){
- print_error( err, " clEnqueueNDRangeKernel failed\n" );
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- clReleaseMemObject(buffers[2]);
- clReleaseMemObject(buffers[3]);
- clReleaseMemObject(pipe);
- clReleaseKernel(kernel[0]);
- clReleaseKernel(kernel[1]);
- clReleaseKernel(kernel[2]);
- clReleaseEvent(producer_sync_event);
- clReleaseEvent(consumer_sync_event);
- clReleaseEvent(pipe_query_sync_event);
- clReleaseEvent(pipe_read_sync_event);
- clReleaseProgram(program);
- align_free(outptr1);
- return -1;
- }
+ test_error_ret(err, " clEnqueueNDRangeKernel failed", -1);
err = clEnqueueReadBuffer(queue, buffers[1], true, 0, size / 2, outptr1, 1, &consumer_sync_event, NULL);
- if ( err != CL_SUCCESS ){
- print_error( err, " clEnqueueReadBuffer failed\n" );
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- clReleaseMemObject(buffers[2]);
- clReleaseMemObject(buffers[3]);
- clReleaseMemObject(pipe);
- clReleaseKernel(kernel[0]);
- clReleaseKernel(kernel[1]);
- clReleaseKernel(kernel[2]);
- clReleaseEvent(producer_sync_event);
- clReleaseEvent(consumer_sync_event);
- clReleaseEvent(pipe_query_sync_event);
- clReleaseEvent(pipe_read_sync_event);
- clReleaseProgram(program);
- align_free(outptr1);
- return -1;
- }
+ test_error_ret(err, " clEnqueueReadBuffer failed", -1);
+
+ // We will reuse this variable so release the previous referred event.
+ clReleaseEvent(pipe_query_sync_event);
// Launch Pipe query kernel
err = clEnqueueNDRangeKernel( queue, kernel[2], 1, NULL, global_work_size_pipe_query, NULL, 1, &consumer_sync_event, &pipe_query_sync_event );
- if ( err != CL_SUCCESS ){
- print_error( err, " clEnqueueNDRangeKernel failed\n" );
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- clReleaseMemObject(buffers[2]);
- clReleaseMemObject(buffers[3]);
- clReleaseMemObject(pipe);
- clReleaseKernel(kernel[0]);
- clReleaseKernel(kernel[1]);
- clReleaseKernel(kernel[2]);
- clReleaseEvent(producer_sync_event);
- clReleaseEvent(consumer_sync_event);
- clReleaseEvent(pipe_query_sync_event);
- clReleaseEvent(pipe_read_sync_event);
- clReleaseProgram(program);
- align_free(outptr1);
- return -1;
- }
+ test_error_ret(err, " clEnqueueNDRangeKernel failed", -1);
err = clEnqueueReadBuffer(queue, buffers[2], true, 0, sizeof(cl_int), &pipe_num_packets, 1, &pipe_query_sync_event, &pipe_read_sync_event);
- if ( err != CL_SUCCESS ){
- print_error( err, " clEnqueueReadBuffer failed\n" );
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- clReleaseMemObject(buffers[2]);
- clReleaseMemObject(buffers[3]);
- clReleaseMemObject(pipe);
- clReleaseKernel(kernel[0]);
- clReleaseKernel(kernel[1]);
- clReleaseKernel(kernel[2]);
- clReleaseEvent(producer_sync_event);
- clReleaseEvent(consumer_sync_event);
- clReleaseEvent(pipe_query_sync_event);
- clReleaseEvent(pipe_read_sync_event);
- clReleaseProgram(program);
- align_free(outptr1);
- return -1;
- }
+ test_error_ret(err, " clEnqueueReadBuffer failed", -1);
// After consumer kernel consumes num_elements/2 from the pipe,
// there are (num_elements - num_elements/2) remaining package in the pipe.
@@ -404,68 +205,24 @@
return -1;
}
+ // We will reuse this variable so release the previous referred event.
+ clReleaseEvent(producer_sync_event);
+
// Launch Producer kernel to fill the pipe again
global_work_size[0] = pipe_num_packets;
err = clEnqueueNDRangeKernel( queue, kernel[0], 1, NULL, global_work_size, NULL, 1, &pipe_read_sync_event, &producer_sync_event );
- if ( err != CL_SUCCESS ){
- print_error( err, " clEnqueueNDRangeKernel failed\n" );
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- clReleaseMemObject(buffers[2]);
- clReleaseMemObject(buffers[3]);
- clReleaseMemObject(pipe);
- clReleaseKernel(kernel[0]);
- clReleaseKernel(kernel[1]);
- clReleaseKernel(kernel[2]);
- clReleaseEvent(producer_sync_event);
- clReleaseEvent(consumer_sync_event);
- clReleaseEvent(pipe_query_sync_event);
- clReleaseEvent(pipe_read_sync_event);
- clReleaseProgram(program);
- align_free(outptr1);
- return -1;
- }
+ test_error_ret(err, " clEnqueueNDRangeKernel failed", -1);
+ // We will reuse this variable so release the previous referred event.
+ clReleaseEvent(pipe_query_sync_event);
// Launch Pipe query kernel
err = clEnqueueNDRangeKernel( queue, kernel[2], 1, NULL, global_work_size_pipe_query, NULL, 1, &producer_sync_event, &pipe_query_sync_event );
- if ( err != CL_SUCCESS ){
- print_error( err, " clEnqueueNDRangeKernel failed\n" );
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- clReleaseMemObject(buffers[2]);
- clReleaseMemObject(buffers[3]);
- clReleaseMemObject(pipe);
- clReleaseKernel(kernel[0]);
- clReleaseKernel(kernel[1]);
- clReleaseKernel(kernel[2]);
- clReleaseEvent(producer_sync_event);
- clReleaseEvent(consumer_sync_event);
- clReleaseEvent(pipe_query_sync_event);
- clReleaseEvent(pipe_read_sync_event);
- clReleaseProgram(program);
- align_free(outptr1);
- return -1;
- }
+ test_error_ret(err, " clEnqueueNDRangeKernel failed", -1);
+ // We will reuse this variable so release the previous referred event.
+ clReleaseEvent(pipe_read_sync_event);
err = clEnqueueReadBuffer(queue, buffers[2], true, 0, sizeof(cl_int), &pipe_num_packets, 1, &pipe_query_sync_event, &pipe_read_sync_event);
- if ( err != CL_SUCCESS ){
- print_error( err, " clEnqueueReadBuffer failed\n" );
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- clReleaseMemObject(buffers[2]);
- clReleaseMemObject(buffers[3]);
- clReleaseMemObject(pipe);
- clReleaseKernel(kernel[0]);
- clReleaseKernel(kernel[1]);
- clReleaseKernel(kernel[2]);
- clReleaseEvent(producer_sync_event);
- clReleaseEvent(consumer_sync_event);
- clReleaseEvent(pipe_query_sync_event);
- clReleaseEvent(pipe_read_sync_event);
- clReleaseProgram(program);
- align_free(outptr1);
- return -1;
- }
+ test_error_ret(err, " clEnqueueReadBuffer failed", -1);
if(pipe_num_packets != num_elements)
{
@@ -473,48 +230,16 @@
return -1;
}
+ // We will reuse this variable so release the previous referred event.
+ clReleaseEvent(consumer_sync_event);
+
// Launch Consumer kernel to consume all packets from pipe
global_work_size[0] = pipe_num_packets;
err = clEnqueueNDRangeKernel( queue, kernel[1], 1, NULL, global_work_size, NULL, 1, &pipe_read_sync_event, &consumer_sync_event );
- if ( err != CL_SUCCESS ){
- print_error( err, " clEnqueueNDRangeKernel failed\n" );
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- clReleaseMemObject(buffers[2]);
- clReleaseMemObject(buffers[3]);
- clReleaseMemObject(pipe);
- clReleaseKernel(kernel[0]);
- clReleaseKernel(kernel[1]);
- clReleaseKernel(kernel[2]);
- clReleaseEvent(producer_sync_event);
- clReleaseEvent(consumer_sync_event);
- clReleaseEvent(pipe_query_sync_event);
- clReleaseEvent(pipe_read_sync_event);
- clReleaseProgram(program);
- align_free(outptr1);
- return -1;
- }
+ test_error_ret(err, " clEnqueueNDRangeKernel failed", -1);
err = clEnqueueReadBuffer(queue, buffers[1], true, 0, size, outptr2, 1, &consumer_sync_event, NULL);
- if ( err != CL_SUCCESS ){
- print_error( err, " clEnqueueReadBuffer failed\n" );
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- clReleaseMemObject(buffers[2]);
- clReleaseMemObject(buffers[3]);
- clReleaseMemObject(pipe);
- clReleaseKernel(kernel[0]);
- clReleaseKernel(kernel[1]);
- clReleaseKernel(kernel[2]);
- clReleaseEvent(producer_sync_event);
- clReleaseEvent(consumer_sync_event);
- clReleaseEvent(pipe_query_sync_event);
- clReleaseEvent(pipe_read_sync_event);
- clReleaseProgram(program);
- align_free(outptr1);
- return -1;
- }
-
+ test_error_ret(err, " clEnqueueReadBuffer failed", -1);
if( verify_result(outptr1, outptr2, num_elements )){
log_error("test_pipe_query_functions failed\n");
@@ -523,22 +248,6 @@
else {
log_info("test_pipe_query_functions passed\n");
}
- //cleanup
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- clReleaseMemObject(buffers[2]);
- clReleaseMemObject(buffers[3]);
- clReleaseMemObject(pipe);
- clReleaseKernel(kernel[0]);
- clReleaseKernel(kernel[1]);
- clReleaseKernel(kernel[2]);
- clReleaseEvent(producer_sync_event);
- clReleaseEvent(consumer_sync_event);
- clReleaseEvent(pipe_query_sync_event);
- clReleaseEvent(pipe_read_sync_event);
- clReleaseProgram(program);
- align_free(outptr1);
-
return 0;
}
diff --git a/test_conformance/pipes/test_pipe_read_write.cpp b/test_conformance/pipes/test_pipe_read_write.cpp
index 91a4990..4bb4468 100644
--- a/test_conformance/pipes/test_pipe_read_write.cpp
+++ b/test_conformance/pipes/test_pipe_read_write.cpp
@@ -15,11 +15,15 @@
//
#include "harness/compat.h"
+#include <assert.h>
+#include <iomanip>
+#include <iostream>
+#include <sstream>
#include <stdio.h>
#include <string.h>
-#include <sys/types.h>
+#include <string>
#include <sys/stat.h>
-#include <assert.h>
+#include <sys/types.h>
#include "procs.h"
#include "kernels.h"
@@ -89,113 +93,139 @@
static const char* convenience_half_kernel_name[] = { "test_pipe_convenience_write_half", "test_pipe_convenience_read_half", "test_pipe_convenience_write_half2", "test_pipe_convenience_read_half2", "test_pipe_convenience_write_half4", "test_pipe_convenience_read_half4", "test_pipe_convenience_write_half8", "test_pipe_convenience_read_half8", "test_pipe_convenience_write_half16", "test_pipe_convenience_read_half16" };
static const char* convenience_double_kernel_name[] = { "test_pipe_convenience_write_double", "test_pipe_convenience_read_double", "test_pipe_convenience_write_double2", "test_pipe_convenience_read_double2", "test_pipe_convenience_write_double4", "test_pipe_convenience_read_double4", "test_pipe_convenience_write_double8", "test_pipe_convenience_read_double8", "test_pipe_convenience_write_double16", "test_pipe_convenience_read_double16" };
-static void insertPragmaForHalfType(char *source, char *type)
+static void insertPragmaForHalfType(std::stringstream &stream, char *type)
{
- source[0] = 0;
- if(strncmp(type, "half",4) == 0)
+ if (strncmp(type, "half", 4) == 0)
{
- strcat(source, "#pragma OPENCL EXTENSION cl_khr_fp16 : enable\n");
+ stream << "#pragma OPENCL EXTENSION cl_khr_fp16 : enable\n";
}
}
-void createKernelSource(char *source, char *type)
+void createKernelSource(std::stringstream &stream, char *type)
{
- char str[512];
- int str_length;
+ insertPragmaForHalfType(stream, type);
- insertPragmaForHalfType(source, type);
+ // clang-format off
+ stream << R"(
+ __kernel void test_pipe_write_)" << type << "(__global " << type << " *src, __write_only pipe " << type << R"( out_pipe)
+ {
+ int gid = get_global_id(0);
+ reserve_id_t res_id;
- sprintf(str, "__kernel void test_pipe_write_%s(__global %s *src, __write_only pipe %s out_pipe)\n", type, type, type);
- strcat(source, str);
- sprintf(str, "{\n int gid = get_global_id(0);\n reserve_id_t res_id; \n\n");
- strcat(source, str);
- sprintf(str, " res_id = reserve_write_pipe(out_pipe, 1);\n if(is_valid_reserve_id(res_id))\n {\n");
- strcat(source, str);
- sprintf(str, " write_pipe(out_pipe, res_id, 0, &src[gid]);\n commit_write_pipe(out_pipe, res_id);\n }\n}\n\n");
- strcat(source, str);
- sprintf(str, "__kernel void test_pipe_read_%s(__read_only pipe %s in_pipe, __global %s *dst)\n", type, type, type);
- strcat(source, str);
- sprintf(str, "{\n int gid = get_global_id(0);\n reserve_id_t res_id; \n\n");
- strcat(source, str);
- sprintf(str, " res_id = reserve_read_pipe(in_pipe, 1);\n if(is_valid_reserve_id(res_id))\n {\n");
- strcat(source, str);
- sprintf(str, " read_pipe(in_pipe, res_id, 0, &dst[gid]);\n commit_read_pipe(in_pipe, res_id);\n }\n}\n");
- strcat(source, str);
- str_length = strlen(source);
- assert(str_length <= STRING_LENGTH);
+ res_id = reserve_write_pipe(out_pipe, 1);
+ if(is_valid_reserve_id(res_id))
+ {
+ write_pipe(out_pipe, res_id, 0, &src[gid]);
+ commit_write_pipe(out_pipe, res_id);
+ }
+ }
+
+ __kernel void test_pipe_read_)" << type << "(__read_only pipe " << type << " in_pipe, __global " << type << R"( *dst)
+ {
+ int gid = get_global_id(0);
+ reserve_id_t res_id;
+
+ res_id = reserve_read_pipe(in_pipe, 1);
+ if(is_valid_reserve_id(res_id))
+ {
+ read_pipe(in_pipe, res_id, 0, &dst[gid]);
+ commit_read_pipe(in_pipe, res_id);
+ }
+ }
+ )";
+ // clang-format on
}
-void createKernelSourceWorkGroup(char *source, char *type)
+void createKernelSourceWorkGroup(std::stringstream &stream, char *type)
{
- char str[512];
- int str_length;
+ insertPragmaForHalfType(stream, type);
- insertPragmaForHalfType(source, type);
+ // clang-format off
+ stream << R"(
+ __kernel void test_pipe_workgroup_write_)" << type << "(__global " << type << " *src, __write_only pipe " << type << R"( out_pipe)
+ {
+ int gid = get_global_id(0);
+ __local reserve_id_t res_id;
- sprintf(str, "__kernel void test_pipe_workgroup_write_%s(__global %s *src, __write_only pipe %s out_pipe)\n", type, type, type);
- strcat(source, str);
- sprintf(str, "{\n int gid = get_global_id(0);\n __local reserve_id_t res_id; \n\n");
- strcat(source, str);
- sprintf(str, " res_id = work_group_reserve_write_pipe(out_pipe, get_local_size(0));\n if(is_valid_reserve_id(res_id))\n {\n");
- strcat(source, str);
- sprintf(str, " write_pipe(out_pipe, res_id, get_local_id(0), &src[gid]);\n work_group_commit_write_pipe(out_pipe, res_id);\n }\n}\n\n");
- strcat(source, str);
- sprintf(str, "__kernel void test_pipe_workgroup_read_%s(__read_only pipe %s in_pipe, __global %s *dst)\n", type, type, type);
- strcat(source, str);
- sprintf(str, "{\n int gid = get_global_id(0);\n __local reserve_id_t res_id; \n\n");
- strcat(source, str);
- sprintf(str, " res_id = work_group_reserve_read_pipe(in_pipe, get_local_size(0));\n if(is_valid_reserve_id(res_id))\n {\n");
- strcat(source, str);
- sprintf(str, " read_pipe(in_pipe, res_id, get_local_id(0), &dst[gid]);\n work_group_commit_read_pipe(in_pipe, res_id);\n }\n}\n");
- strcat(source, str);
- str_length = strlen(source);
- assert(str_length <= STRING_LENGTH);
+ res_id = work_group_reserve_write_pipe(out_pipe, get_local_size(0));
+ if(is_valid_reserve_id(res_id))
+ {
+ write_pipe(out_pipe, res_id, get_local_id(0), &src[gid]);
+ work_group_commit_write_pipe(out_pipe, res_id);
+ }
+ }
+
+ __kernel void test_pipe_workgroup_read_)" << type << "(__read_only pipe " << type << " in_pipe, __global " << type << R"( *dst)
+ {
+ int gid = get_global_id(0);
+ __local reserve_id_t res_id;
+
+ res_id = work_group_reserve_read_pipe(in_pipe, get_local_size(0));
+ if(is_valid_reserve_id(res_id))
+ {
+ read_pipe(in_pipe, res_id, get_local_id(0), &dst[gid]);
+ work_group_commit_read_pipe(in_pipe, res_id);
+ }
+ }
+ )";
+ // clang-format on
}
-void createKernelSourceSubGroup(char *source, char *type)
+void createKernelSourceSubGroup(std::stringstream &stream, char *type)
{
- char str[512];
- int str_length;
+ insertPragmaForHalfType(stream, type);
- insertPragmaForHalfType(source, type);
+ // clang-format off
+ stream << R"(
+ #pragma OPENCL EXTENSION cl_khr_subgroups : enable
+ __kernel void test_pipe_subgroup_write_)" << type << "(__global " << type << " *src, __write_only pipe " << type << R"( out_pipe)
+ {
+ int gid = get_global_id(0);
+ reserve_id_t res_id;
- sprintf(str, "#pragma OPENCL EXTENSION cl_khr_subgroups : enable\n__kernel void test_pipe_subgroup_write_%s(__global %s *src, __write_only pipe %s out_pipe)\n", type, type, type);
- strcat(source, str);
- sprintf(str, "{\n int gid = get_global_id(0);\n reserve_id_t res_id; \n\n");
- strcat(source, str);
- sprintf(str, " res_id = sub_group_reserve_write_pipe(out_pipe, get_sub_group_size());\n if(is_valid_reserve_id(res_id))\n {\n");
- strcat(source, str);
- sprintf(str, " write_pipe(out_pipe, res_id, get_sub_group_local_id(), &src[gid]);\n sub_group_commit_write_pipe(out_pipe, res_id);\n }\n}\n\n");
- strcat(source, str);
- sprintf(str, "__kernel void test_pipe_subgroup_read_%s(__read_only pipe %s in_pipe, __global %s *dst)\n", type, type, type);
- strcat(source, str);
- sprintf(str, "{\n int gid = get_global_id(0);\n reserve_id_t res_id; \n\n");
- strcat(source, str);
- sprintf(str, " res_id = sub_group_reserve_read_pipe(in_pipe, get_sub_group_size());\n if(is_valid_reserve_id(res_id))\n {\n");
- strcat(source, str);
- sprintf(str, " read_pipe(in_pipe, res_id, get_sub_group_local_id(), &dst[gid]);\n sub_group_commit_read_pipe(in_pipe, res_id);\n }\n}\n");
- strcat(source, str);
- str_length = strlen(source);
- assert(str_length <= STRING_LENGTH);
+ res_id = sub_group_reserve_write_pipe(out_pipe, get_sub_group_size());
+ if(is_valid_reserve_id(res_id))
+ {
+ write_pipe(out_pipe, res_id, get_sub_group_local_id(), &src[gid]);
+ sub_group_commit_write_pipe(out_pipe, res_id);
+ }
+ }
+
+ __kernel void test_pipe_subgroup_read_)" << type << "(__read_only pipe " << type << " in_pipe, __global " << type << R"( *dst)
+ {
+ int gid = get_global_id(0);
+ reserve_id_t res_id;
+
+ res_id = sub_group_reserve_read_pipe(in_pipe, get_sub_group_size());
+ if(is_valid_reserve_id(res_id))
+ {
+ read_pipe(in_pipe, res_id, get_sub_group_local_id(), &dst[gid]);
+ sub_group_commit_read_pipe(in_pipe, res_id);
+ }
+ }
+ )";
+ // clang-format on
}
-void createKernelSourceConvenience(char *source, char *type)
+void createKernelSourceConvenience(std::stringstream &stream, char *type)
{
- char str[512];
- int str_length;
+ insertPragmaForHalfType(stream, type);
- insertPragmaForHalfType(source, type);
+ // clang-format off
+ stream << R"(
+ __kernel void test_pipe_convenience_write_)" << type << "(__global " << type << " *src, __write_only pipe " << type << R"( out_pipe)
+ {
+ int gid = get_global_id(0);
+ write_pipe(out_pipe, &src[gid]);
+ }
- sprintf(str, "__kernel void test_pipe_convenience_write_%s(__global %s *src, __write_only pipe %s out_pipe)\n", type, type, type);
- strcat(source, str);
- sprintf(str, "{\n int gid = get_global_id(0);\n write_pipe(out_pipe, &src[gid]);\n}\n\n");
- strcat(source, str);
- sprintf(str, "__kernel void test_pipe_convenience_read_%s(__read_only pipe %s in_pipe, __global %s *dst)\n", type, type, type);
- strcat(source, str);
- sprintf(str, "{\n int gid = get_global_id(0);\n read_pipe(in_pipe, &dst[gid]);\n}\n");
- strcat(source, str);
- str_length = strlen(source);
- assert(str_length <= STRING_LENGTH);
+ __kernel void test_pipe_convenience_read_)" << type << "(__read_only pipe " << type << " in_pipe, __global " << type << R"( *dst)
+ {
+ int gid = get_global_id(0);
+ read_pipe(in_pipe, &dst[gid]);
+ }
+ )";
+ // clang-format on
}
// verify functions
@@ -424,23 +454,24 @@
int test_pipe_readwrite( cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements, size_t size, char *type, int loops,
void *inptr[5], const char *kernelName[], int (*fn)(void *, void *, int) )
{
- cl_mem pipes[5];
- cl_mem buffers[10];
- void *outptr[5];
- cl_program program[5];
- cl_kernel kernel[10];
- size_t global_work_size[3];
- size_t local_work_size[3];
- cl_int err;
- int i, ii;
- size_t ptrSizes[5];
- int total_errors = 0;
- cl_event producer_sync_event[5];
- cl_event consumer_sync_event[5];
- char *sourceCode[5];
- char vector_type[10];
+ clMemWrapper pipes[5];
+ clMemWrapper buffers[10];
+ void *outptr[5];
+ BufferOwningPtr<cl_int> BufferOutPtr[5];
+ clProgramWrapper program[5];
+ clKernelWrapper kernel[10];
+ size_t global_work_size[3];
+ size_t local_work_size[3];
+ cl_int err;
+ int i, ii;
+ size_t ptrSizes[5];
+ int total_errors = 0;
+ clEventWrapper producer_sync_event[5];
+ clEventWrapper consumer_sync_event[5];
+ std::stringstream sourceCode[5];
+ char vector_type[10];
- size_t min_alignment = get_min_alignment(context);
+ size_t min_alignment = get_min_alignment(context);
global_work_size[0] = (cl_uint)num_elements;
@@ -450,217 +481,133 @@
ptrSizes[3] = ptrSizes[2] << 1;
ptrSizes[4] = ptrSizes[3] << 1;
- for( i = 0; i < loops; i++)
+ for (i = 0; i < loops; i++)
{
ii = i << 1;
- sourceCode[i] = (char*) malloc(STRING_LENGTH * sizeof(char));
- buffers[ii] = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR, ptrSizes[i] * num_elements, inptr[i], &err);
- if(err){
- clReleaseMemObject(buffers[ii]);
- align_free( outptr[i] );
- print_error(err, " clCreateBuffer failed\n");
- return -1;
- }
- outptr[i] = align_malloc( ptrSizes[i] * num_elements, min_alignment);
- buffers[ii+1] = clCreateBuffer(context, CL_MEM_USE_HOST_PTR, ptrSizes[i] * num_elements, outptr[i], &err);
- if ( err ){
- clReleaseMemObject(buffers[ii]);
- align_free( outptr[i] );
- print_error(err, " clCreateBuffer failed\n" );
- return -1;
- }
+ buffers[ii] =
+ clCreateBuffer(context, CL_MEM_COPY_HOST_PTR,
+ ptrSizes[i] * num_elements, inptr[i], &err);
+ test_error_ret(err, " clCreateBuffer failed", -1);
+
+ outptr[i] = align_malloc(ptrSizes[i] * num_elements, min_alignment);
+ BufferOutPtr[i].reset(outptr[i], nullptr, 0, size, true);
+ buffers[ii + 1] =
+ clCreateBuffer(context, CL_MEM_USE_HOST_PTR,
+ ptrSizes[i] * num_elements, outptr[i], &err);
+ test_error_ret(err, " clCreateBuffer failed", -1);
+
// Creating pipe with non-power of 2 size
- pipes[i] = clCreatePipe(context, CL_MEM_HOST_NO_ACCESS, ptrSizes[i], num_elements+3, NULL, &err);
- if(err){
- clReleaseMemObject(pipes[i]);
- print_error(err, " clCreatePipe failed\n");
- return -1;
- }
+ pipes[i] = clCreatePipe(context, CL_MEM_HOST_NO_ACCESS, ptrSizes[i],
+ num_elements + 3, NULL, &err);
+ test_error_ret(err, " clCreatePipe failed", -1);
- switch(i)
+ switch (i)
{
- case 0:
- sprintf(vector_type, "%s", type);
- break;
- case 1:
- sprintf(vector_type, "%s%d", type, 2);
- break;
- case 2:
- sprintf(vector_type, "%s%d", type, 4);
- break;
- case 3:
- sprintf(vector_type, "%s%d", type, 8);
- break;
- case 4:
- sprintf(vector_type, "%s%d", type, 16);
- break;
+ case 0: sprintf(vector_type, "%s", type); break;
+ case 1: sprintf(vector_type, "%s%d", type, 2); break;
+ case 2: sprintf(vector_type, "%s%d", type, 4); break;
+ case 3: sprintf(vector_type, "%s%d", type, 8); break;
+ case 4: sprintf(vector_type, "%s%d", type, 16); break;
}
- if(useWorkgroupReserve == 1){
+ if (useWorkgroupReserve == 1)
+ {
createKernelSourceWorkGroup(sourceCode[i], vector_type);
}
- else if(useSubgroupReserve == 1){
+ else if (useSubgroupReserve == 1)
+ {
createKernelSourceSubGroup(sourceCode[i], vector_type);
}
- else if(useConvenienceBuiltIn == 1){
+ else if (useConvenienceBuiltIn == 1)
+ {
createKernelSourceConvenience(sourceCode[i], vector_type);
}
- else{
+ else
+ {
createKernelSource(sourceCode[i], vector_type);
}
+ std::string kernel_source = sourceCode[i].str();
+ const char *sources[] = { kernel_source.c_str() };
// Create producer kernel
- err = create_single_kernel_helper_with_build_options(context, &program[i], &kernel[ii], 1, (const char**)&sourceCode[i], kernelName[ii], "-cl-std=CL2.0");
- if(err){
- clReleaseMemObject(buffers[ii]);
- clReleaseMemObject(buffers[ii+1]);
- clReleaseMemObject(pipes[i]);
- align_free( outptr[i] );
- print_error(err, "Error creating program\n");
- return -1;
- }
- //Create consumer kernel
+ err = create_single_kernel_helper_with_build_options(
+ context, &program[i], &kernel[ii], 1, sources, kernelName[ii],
+ "-cl-std=CL2.0");
+
+ test_error_ret(err, " Error creating program", -1);
+
+ // Create consumer kernel
kernel[ii + 1] = clCreateKernel(program[i], kernelName[ii + 1], &err);
- if( kernel[ii + 1] == NULL || err != CL_SUCCESS)
- {
- clReleaseMemObject(buffers[ii]);
- clReleaseMemObject(buffers[ii+1]);
- clReleaseMemObject(pipes[i]);
- align_free( outptr[i] );
- log_error("Creating program for %s\n", type);
- print_error( err, "Unable to create kernel" );
- return -1;
- }
+ test_error_ret(err, " Error creating kernel", -1);
- err = clSetKernelArg(kernel[ii], 0, sizeof(cl_mem), (void*)&buffers[ii]);
- err |= clSetKernelArg(kernel[ii], 1, sizeof(cl_mem), (void*)&pipes[i]);
- err |= clSetKernelArg(kernel[ii + 1], 0, sizeof(cl_mem), (void*)&pipes[i]);
- err |= clSetKernelArg(kernel[ii + 1], 1, sizeof(cl_mem), (void*)&buffers[ii + 1]);
- if ( err != CL_SUCCESS ){
- clReleaseMemObject(buffers[ii]);
- clReleaseMemObject(buffers[ii+1]);
- clReleaseMemObject(pipes[i]);
- clReleaseKernel(kernel[ii]);
- clReleaseKernel(kernel[ii+1]);
- clReleaseProgram(program[i]);
- align_free(outptr[i]);
- print_error(err, " clSetKernelArg failed");
- return -1;
- }
+ err =
+ clSetKernelArg(kernel[ii], 0, sizeof(cl_mem), (void *)&buffers[ii]);
+ err |= clSetKernelArg(kernel[ii], 1, sizeof(cl_mem), (void *)&pipes[i]);
+ err |= clSetKernelArg(kernel[ii + 1], 0, sizeof(cl_mem),
+ (void *)&pipes[i]);
+ err |= clSetKernelArg(kernel[ii + 1], 1, sizeof(cl_mem),
+ (void *)&buffers[ii + 1]);
+ test_error_ret(err, " clSetKernelArg failed", -1);
- if(useWorkgroupReserve == 1 || useSubgroupReserve == 1)
+ if (useWorkgroupReserve == 1 || useSubgroupReserve == 1)
{
- err = get_max_common_work_group_size( context, kernel[ii], global_work_size[0], &local_work_size[0] );
- test_error( err, "Unable to get work group size to use" );
+ err = get_max_common_work_group_size(
+ context, kernel[ii], global_work_size[0], &local_work_size[0]);
+ test_error(err, "Unable to get work group size to use");
// Launch Producer kernel
- err = clEnqueueNDRangeKernel( queue, kernel[ii], 1, NULL, global_work_size, local_work_size, 0, NULL, &producer_sync_event[i] );
- if ( err != CL_SUCCESS ){
- print_error( err, " clEnqueueNDRangeKernel failed" );
- clReleaseMemObject(buffers[ii]);
- clReleaseMemObject(buffers[ii+1]);
- clReleaseMemObject(pipes[i]);
- clReleaseKernel(kernel[ii]);
- clReleaseKernel(kernel[ii+1]);
- clReleaseEvent(producer_sync_event[i]);
- clReleaseEvent(consumer_sync_event[i]);
- clReleaseProgram(program[i]);
- align_free(outptr[i]);
- return -1;
- }
+ err = clEnqueueNDRangeKernel(queue, kernel[ii], 1, NULL,
+ global_work_size, local_work_size, 0,
+ NULL, &producer_sync_event[i]);
+ test_error_ret(err, " clEnqueueNDRangeKernel failed", -1);
}
else
{
// Launch Producer kernel
- err = clEnqueueNDRangeKernel( queue, kernel[ii], 1, NULL, global_work_size, NULL, 0, NULL, &producer_sync_event[i] );
- if ( err != CL_SUCCESS ){
- print_error( err, " clEnqueueNDRangeKernel failed" );
- clReleaseMemObject(buffers[ii]);
- clReleaseMemObject(buffers[ii+1]);
- clReleaseMemObject(pipes[i]);
- clReleaseKernel(kernel[ii]);
- clReleaseKernel(kernel[ii+1]);
- clReleaseEvent(producer_sync_event[i]);
- clReleaseEvent(consumer_sync_event[i]);
- clReleaseProgram(program[i]);
- align_free(outptr[i]);
- return -1;
- }
+ err = clEnqueueNDRangeKernel(queue, kernel[ii], 1, NULL,
+ global_work_size, NULL, 0, NULL,
+ &producer_sync_event[i]);
+ test_error_ret(err, " clEnqueueNDRangeKernel failed", -1);
}
- if(useWorkgroupReserve == 1 || useSubgroupReserve == 1)
+ if (useWorkgroupReserve == 1 || useSubgroupReserve == 1)
{
- err = get_max_common_work_group_size( context, kernel[ii + 1], global_work_size[0], &local_work_size[0] );
- test_error( err, "Unable to get work group size to use" );
+ err = get_max_common_work_group_size(context, kernel[ii + 1],
+ global_work_size[0],
+ &local_work_size[0]);
+ test_error(err, "Unable to get work group size to use");
// Launch Consumer kernel
- err = clEnqueueNDRangeKernel( queue, kernel[ii + 1], 1, NULL, global_work_size, local_work_size, 1, &producer_sync_event[i], &consumer_sync_event[i] );
- if ( err != CL_SUCCESS ){
- print_error( err, " clEnqueueNDRangeKernel failed" );
- clReleaseMemObject(buffers[ii]);
- clReleaseMemObject(buffers[ii+1]);
- clReleaseMemObject(pipes[i]);
- clReleaseKernel(kernel[ii]);
- clReleaseKernel(kernel[ii+1]);
- clReleaseEvent(producer_sync_event[i]);
- clReleaseEvent(consumer_sync_event[i]);
- clReleaseProgram(program[i]);
- align_free(outptr[i]);
- return -1;
- }
+ err = clEnqueueNDRangeKernel(queue, kernel[ii + 1], 1, NULL,
+ global_work_size, local_work_size, 1,
+ &producer_sync_event[i],
+ &consumer_sync_event[i]);
+ test_error_ret(err, " clEnqueueNDRangeKernel failed", -1);
}
else
{
// Launch Consumer kernel
- err = clEnqueueNDRangeKernel( queue, kernel[ii + 1], 1, NULL, global_work_size, NULL, 1, &producer_sync_event[i], &consumer_sync_event[i] );
- if ( err != CL_SUCCESS ){
- print_error( err, " clEnqueueNDRangeKernel failed" );
- clReleaseMemObject(buffers[ii]);
- clReleaseMemObject(buffers[ii+1]);
- clReleaseMemObject(pipes[i]);
- clReleaseKernel(kernel[ii]);
- clReleaseKernel(kernel[ii+1]);
- clReleaseEvent(producer_sync_event[i]);
- clReleaseEvent(consumer_sync_event[i]);
- clReleaseProgram(program[i]);
- align_free(outptr[i]);
- return -1;
- }
+ err = clEnqueueNDRangeKernel(
+ queue, kernel[ii + 1], 1, NULL, global_work_size, NULL, 1,
+ &producer_sync_event[i], &consumer_sync_event[i]);
+ test_error_ret(err, " clEnqueueNDRangeKernel failed", -1);
}
- err = clEnqueueReadBuffer(queue, buffers[ii+1], true, 0, ptrSizes[i]*num_elements, outptr[i], 1, &consumer_sync_event[i], NULL);
- if ( err != CL_SUCCESS ){
- print_error( err, " clEnqueueReadBuffer failed" );
- clReleaseMemObject(buffers[ii]);
- clReleaseMemObject(buffers[ii+1]);
- clReleaseMemObject(pipes[i]);
- clReleaseKernel(kernel[ii]);
- clReleaseKernel(kernel[ii+1]);
- clReleaseEvent(producer_sync_event[i]);
- clReleaseEvent(consumer_sync_event[i]);
- clReleaseProgram(program[i]);
- align_free(outptr[i]);
- return -1;
- }
+ err = clEnqueueReadBuffer(queue, buffers[ii + 1], true, 0,
+ ptrSizes[i] * num_elements, outptr[i], 1,
+ &consumer_sync_event[i], NULL);
+ test_error_ret(err, " clEnqueueReadBuffer failed", -1);
- if( fn( inptr[i], outptr[i], (int)(ptrSizes[i] * (size_t)num_elements / ptrSizes[0]))){
- log_error("%s%d test failed\n", type, 1<<i);
+ if (fn(inptr[i], outptr[i],
+ (int)(ptrSizes[i] * (size_t)num_elements / ptrSizes[0])))
+ {
+ log_error("%s%d test failed\n", type, 1 << i);
total_errors++;
}
- else {
- log_info("%s%d test passed\n", type, 1<<i);
+ else
+ {
+ log_info("%s%d test passed\n", type, 1 << i);
}
- //cleanup
- clReleaseMemObject(buffers[ii]);
- clReleaseMemObject(buffers[ii+1]);
- clReleaseMemObject(pipes[i]);
- clReleaseKernel(kernel[ii]);
- clReleaseKernel(kernel[ii+1]);
- clReleaseEvent(producer_sync_event[i]);
- clReleaseEvent(consumer_sync_event[i]);
- clReleaseProgram(program[i]);
- align_free(outptr[i]);
-
}
return total_errors;
@@ -669,166 +616,80 @@
int test_pipe_readwrite_struct_generic( cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements,
const char *kernelCode, const char *kernelName[])
{
- cl_mem buffers[2];
- cl_mem pipe;
- void *outptr;
- TestStruct *inptr;
- cl_program program;
- cl_kernel kernel[2];
- size_t size = sizeof(TestStruct);
- size_t global_work_size[3];
- cl_int err;
- int total_errors = 0;
- int i;
- MTdata d = init_genrand( gRandomSeed );
- cl_event producer_sync_event = NULL;
- cl_event consumer_sync_event = NULL;
+ clMemWrapper buffers[2];
+ clMemWrapper pipe;
+ void *outptr;
+ TestStruct *inptr;
+ BufferOwningPtr<cl_int> BufferInPtr;
+ BufferOwningPtr<TestStruct> BufferOutPtr;
+ clProgramWrapper program;
+ clKernelWrapper kernel[2];
+ size_t size = sizeof(TestStruct);
+ size_t global_work_size[3];
+ cl_int err;
+ int total_errors = 0;
+ int i;
+ MTdataHolder d(gRandomSeed);
+ clEventWrapper producer_sync_event = NULL;
+ clEventWrapper consumer_sync_event = NULL;
- size_t min_alignment = get_min_alignment(context);
+ size_t min_alignment = get_min_alignment(context);
global_work_size[0] = (size_t)num_elements;
inptr = (TestStruct *)align_malloc(size * num_elements, min_alignment);
- for ( i = 0; i < num_elements; i++ ){
+ for (i = 0; i < num_elements; i++)
+ {
inptr[i].a = (char)genrand_int32(d);
inptr[i].b = genrand_int32(d);
}
+ BufferInPtr.reset(inptr, nullptr, 0, size, true);
+
buffers[0] = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR, size * num_elements, inptr, &err);
- if(err){
- clReleaseMemObject(buffers[0]);
- print_error(err, " clCreateBuffer failed\n");
- return -1;
- }
+ test_error_ret(err, " clCreateBuffer failed", -1);
+
outptr = align_malloc( size * num_elements, min_alignment);
+ BufferOutPtr.reset(outptr, nullptr, 0, size, true);
+
buffers[1] = clCreateBuffer(context, CL_MEM_USE_HOST_PTR, size * num_elements, outptr, &err);
- if (err){
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- align_free( outptr );
- print_error(err, " clCreateBuffer failed\n" );
- return -1;
- }
+ test_error_ret(err, " clCreateBuffer failed", -1);
+
pipe = clCreatePipe(context, CL_MEM_HOST_NO_ACCESS, size, num_elements, NULL, &err);
- if(err){
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- align_free( outptr );
- clReleaseMemObject(pipe);
- print_error(err, " clCreatePipe failed\n");
- return -1;
- }
+ test_error_ret(err, " clCreatePipe failed", -1);
+
// Create producer kernel
err = create_single_kernel_helper_with_build_options(context, &program, &kernel[0], 1, &kernelCode, kernelName[0], "-cl-std=CL2.0");
- if(err){
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- clReleaseMemObject(pipe);
- align_free(outptr);
- log_error(" Error creating program for struct\n");
- print_error(err, "Error creating program\n");
- return -1;
- }
+ test_error_ret(err, " Error creating program", -1);
+
//Create consumer kernel
kernel[1] = clCreateKernel(program, kernelName[1], &err);
- if( kernel[1] == NULL || err != CL_SUCCESS)
- {
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- clReleaseMemObject(pipe);
- align_free(outptr);
- print_error(err, "Error creating kernel\n");
- return -1;
- }
+ test_error_ret(err, " Error creating kernel", -1);
err = clSetKernelArg(kernel[0], 0, sizeof(cl_mem), (void*)&buffers[0]);
err |= clSetKernelArg(kernel[0], 1, sizeof(cl_mem), (void*)&pipe);
err |= clSetKernelArg(kernel[1], 0, sizeof(cl_mem), (void*)&pipe);
err |= clSetKernelArg(kernel[1], 1, sizeof(cl_mem), (void*)&buffers[1]);
- if (err != CL_SUCCESS){
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- clReleaseMemObject(pipe);
- clReleaseKernel(kernel[0]);
- clReleaseKernel(kernel[1]);
- clReleaseProgram(program);
- align_free(outptr);
- print_error(err, " clSetKernelArg failed");
- return -1;
- }
+ test_error_ret(err, " clSetKernelArg failed", -1);
// Launch Producer kernel
err = clEnqueueNDRangeKernel( queue, kernel[0], 1, NULL, global_work_size, NULL, 0, NULL, &producer_sync_event );
- if (err != CL_SUCCESS){
- print_error( err, " clEnqueueNDRangeKernel failed" );
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- clReleaseMemObject(pipe);
- clReleaseKernel(kernel[0]);
- clReleaseKernel(kernel[1]);
- clReleaseEvent(producer_sync_event);
- clReleaseProgram(program);
- align_free(outptr);
- return -1;
- }
+ test_error_ret(err, " clEnqueueNDRangeKernel failed", -1);
// Launch Consumer kernel
err = clEnqueueNDRangeKernel( queue, kernel[1], 1, NULL, global_work_size, NULL, 1, &producer_sync_event, &consumer_sync_event );
- if (err != CL_SUCCESS){
- print_error( err, " clEnqueueNDRangeKernel failed" );
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- clReleaseMemObject(pipe);
- clReleaseKernel(kernel[0]);
- clReleaseKernel(kernel[1]);
- clReleaseEvent(producer_sync_event);
- clReleaseEvent(consumer_sync_event);
- clReleaseProgram(program);
- align_free(outptr);
- return -1;
- }
+ test_error_ret(err, " clEnqueueNDRangeKernel failed", -1);
err = clEnqueueReadBuffer(queue, buffers[1], true, 0, size*num_elements, outptr, 1, &consumer_sync_event, NULL);
- if (err != CL_SUCCESS){
- print_error( err, " clEnqueueReadBuffer failed" );
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- clReleaseMemObject(pipe);
- clReleaseKernel(kernel[0]);
- clReleaseKernel(kernel[1]);
- clReleaseEvent(producer_sync_event);
- clReleaseEvent(consumer_sync_event);
- clReleaseProgram(program);
- align_free(outptr);
- return -1;
- }
+ test_error_ret(err, " clEnqueueReadBuffer failed", -1);
if( verify_readwrite_struct( inptr, outptr, num_elements)){
log_error("struct_readwrite test failed\n");
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- clReleaseMemObject(pipe);
- clReleaseKernel(kernel[0]);
- clReleaseKernel(kernel[1]);
- clReleaseEvent(producer_sync_event);
- clReleaseEvent(consumer_sync_event);
- clReleaseProgram(program);
- align_free(outptr);
return -1;
}
else {
log_info("struct_readwrite test passed\n");
}
- //cleanup
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- clReleaseMemObject(pipe);
- clReleaseKernel(kernel[0]);
- clReleaseKernel(kernel[1]);
- clReleaseEvent(producer_sync_event);
- clReleaseEvent(consumer_sync_event);
- clReleaseProgram(program);
- align_free(outptr);
return 0;
}
diff --git a/test_conformance/pipes/test_pipe_readwrite_errors.cpp b/test_conformance/pipes/test_pipe_readwrite_errors.cpp
index cdfaf5e..1b9fc31 100644
--- a/test_conformance/pipes/test_pipe_readwrite_errors.cpp
+++ b/test_conformance/pipes/test_pipe_readwrite_errors.cpp
@@ -64,23 +64,26 @@
int test_pipe_readwrite_errors(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
{
- cl_mem pipe;
- cl_mem buffers[3];
- void *outptr;
- cl_int *inptr;
- cl_program program;
- cl_kernel kernel[2];
- size_t global_work_size[3];
- cl_int err;
- cl_int size;
- cl_int i;
- cl_int status = 0;
- cl_event producer_sync_event;
- cl_event consumer_sync_event;
- MTdata d = init_genrand( gRandomSeed );
- const char* kernelName[] = {"test_pipe_write_error", "test_pipe_read_error"};
+ clMemWrapper pipe;
+ clMemWrapper buffers[3];
+ void *outptr;
+ cl_int *inptr;
+ clProgramWrapper program;
+ clKernelWrapper kernel[2];
+ size_t global_work_size[3];
+ cl_int err;
+ cl_int size;
+ cl_int i;
+ cl_int status = 0;
+ clEventWrapper producer_sync_event;
+ clEventWrapper consumer_sync_event;
+ BufferOwningPtr<cl_int> BufferInPtr;
+ BufferOwningPtr<cl_int> BufferOutPtr;
+ MTdataHolder d(gRandomSeed);
+ const char *kernelName[] = { "test_pipe_write_error",
+ "test_pipe_read_error" };
- size_t min_alignment = get_min_alignment(context);
+ size_t min_alignment = get_min_alignment(context);
global_work_size[0] = num_elements;
@@ -88,69 +91,36 @@
inptr = (cl_int *)align_malloc(size, min_alignment);
- for(i = 0; i < (cl_int)(size / sizeof(int)); i++){
+ for (i = 0; i < num_elements; i++)
+ {
inptr[i] = (int)genrand_int32(d);
}
+ BufferInPtr.reset(inptr, nullptr, 0, size, true);
- buffers[0] = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR, size, inptr, &err);
- if(err){
- clReleaseMemObject(buffers[0]);
- print_error(err, " clCreateBuffer failed\n");
- return -1;
- }
+ buffers[0] =
+ clCreateBuffer(context, CL_MEM_COPY_HOST_PTR, size, inptr, &err);
+ test_error_ret(err, " clCreateBuffer failed", -1);
+
outptr = align_malloc(size, min_alignment);
+ BufferOutPtr.reset(outptr, nullptr, 0, size, true);
+
buffers[1] = clCreateBuffer(context, CL_MEM_USE_HOST_PTR, size, outptr, &err);
- if ( err ){
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- align_free( outptr );
- print_error(err, " clCreateBuffer failed\n" );
- return -1;
- }
+ test_error_ret(err, " clCreateBuffer failed", -1);
+
buffers[2] = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR, sizeof(int), &status, &err);
- if ( err ){
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- clReleaseMemObject(buffers[2]);
- align_free( outptr );
- print_error(err, " clCreateBuffer failed\n" );
- return -1;
- }
+ test_error_ret(err, " clCreateBuffer failed", -1);
+
//Pipe created with max_packets less than global size
pipe = clCreatePipe(context, CL_MEM_HOST_NO_ACCESS, sizeof(int), num_elements - (num_elements/2), NULL, &err);
- if(err){
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- clReleaseMemObject(buffers[2]);
- align_free( outptr );
- clReleaseMemObject(pipe);
- print_error(err, " clCreatePipe failed\n");
- return -1;
- }
+ test_error_ret(err, " clCreatePipe failed", -1);
// Create producer kernel
err = create_single_kernel_helper_with_build_options(context, &program, &kernel[0], 1, (const char**)&pipe_readwrite_errors_kernel_code, kernelName[0], "-cl-std=CL2.0");
- if(err){
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- clReleaseMemObject(buffers[2]);
- clReleaseMemObject(pipe);
- align_free(outptr);
- print_error(err, "Error creating program\n");
- return -1;
- }
+ test_error_ret(err, " Error creating program", -1);
+
//Create consumer kernel
kernel[1] = clCreateKernel(program, kernelName[1], &err);
- if( kernel[1] == NULL || err != CL_SUCCESS)
- {
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- clReleaseMemObject(buffers[2]);
- clReleaseMemObject(pipe);
- align_free(outptr);
- print_error(err, "Error creating kernel\n");
- return -1;
- }
+ test_error_ret(err, " Error creating kernel", -1);
err = clSetKernelArg(kernel[0], 0, sizeof(cl_mem), (void*)&buffers[0]);
err |= clSetKernelArg(kernel[0], 1, sizeof(cl_mem), (void*)&pipe);
@@ -158,49 +128,15 @@
err |= clSetKernelArg(kernel[1], 0, sizeof(cl_mem), (void*)&pipe);
err |= clSetKernelArg(kernel[1], 1, sizeof(cl_mem), (void*)&buffers[1]);
err |= clSetKernelArg(kernel[1], 2, sizeof(cl_mem), (void*)&buffers[2]);
- if ( err != CL_SUCCESS ){
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- clReleaseMemObject(buffers[2]);
- clReleaseMemObject(pipe);
- clReleaseKernel(kernel[0]);
- clReleaseKernel(kernel[1]);
- clReleaseProgram(program);
- align_free(outptr);
- print_error(err, " clSetKernelArg failed");
- return -1;
- }
+
+ test_error_ret(err, " clSetKernelArg failed", -1);
// Launch Consumer kernel for empty pipe
err = clEnqueueNDRangeKernel( queue, kernel[1], 1, NULL, global_work_size, NULL, 0, NULL, &consumer_sync_event );
- if ( err != CL_SUCCESS ){
- print_error( err, " clEnqueueNDRangeKernel failed" );
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- clReleaseMemObject(buffers[2]);
- clReleaseMemObject(pipe);
- clReleaseKernel(kernel[0]);
- clReleaseKernel(kernel[1]);
- clReleaseEvent(consumer_sync_event);
- clReleaseProgram(program);
- align_free(outptr);
- return -1;
- }
+ test_error_ret(err, " clEnqueueNDRangeKernel failed", -1);
err = clEnqueueReadBuffer(queue, buffers[2], true, 0, sizeof(status), &status, 1, &consumer_sync_event, NULL);
- if ( err != CL_SUCCESS ){
- print_error( err, " clEnqueueReadBuffer failed" );
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- clReleaseMemObject(buffers[2]);
- clReleaseMemObject(pipe);
- clReleaseKernel(kernel[0]);
- clReleaseKernel(kernel[1]);
- clReleaseEvent(consumer_sync_event);
- clReleaseProgram(program);
- align_free(outptr);
- return -1;
- }
+ test_error_ret(err, " clEnqueueReadBuffer failed", -1);
if(status == 0){
log_error("test_pipe_readwrite_errors failed\n");
@@ -212,34 +148,13 @@
// Launch Producer kernel
err = clEnqueueNDRangeKernel( queue, kernel[0], 1, NULL, global_work_size, NULL, 0, NULL, &producer_sync_event );
- if ( err != CL_SUCCESS ){
- print_error( err, " clEnqueueNDRangeKernel failed" );
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- clReleaseMemObject(buffers[2]);
- clReleaseMemObject(pipe);
- clReleaseKernel(kernel[0]);
- clReleaseKernel(kernel[1]);
- clReleaseProgram(program);
- align_free(outptr);
- return -1;
- }
+ test_error_ret(err, " clEnqueueNDRangeKernel failed", -1);
err = clEnqueueReadBuffer(queue, buffers[2], true, 0, sizeof(status), &status, 1, &producer_sync_event, NULL);
- if ( err != CL_SUCCESS ){
- print_error( err, " clEnqueueReadBuffer failed" );
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- clReleaseMemObject(buffers[2]);
- clReleaseMemObject(pipe);
- clReleaseKernel(kernel[0]);
- clReleaseKernel(kernel[1]);
- clReleaseProgram(program);
- align_free(outptr);
- return -1;
- }
+ test_error_ret(err, " clEnqueueReadBuffer failed", -1);
- if(status == 0){
+ if (status == 0)
+ {
log_error("test_pipe_readwrite_errors failed\n");
return -1;
}
@@ -247,66 +162,27 @@
status = 0;
}
+ // We will reuse this variable so release the previous referred event.
+ clReleaseEvent(consumer_sync_event);
+
// Launch Consumer kernel
err = clEnqueueNDRangeKernel( queue, kernel[1], 1, NULL, global_work_size, NULL, 1, &producer_sync_event, &consumer_sync_event );
- if ( err != CL_SUCCESS ){
- print_error( err, " clEnqueueNDRangeKernel failed" );
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- clReleaseMemObject(buffers[2]);
- clReleaseMemObject(pipe);
- clReleaseKernel(kernel[0]);
- clReleaseKernel(kernel[1]);
- clReleaseEvent(producer_sync_event);
- clReleaseEvent(consumer_sync_event);
- clReleaseProgram(program);
- align_free(outptr);
- return -1;
- }
+ test_error_ret(err, " clEnqueueNDRangeKernel failed", -1);
err = clEnqueueReadBuffer(queue, buffers[2], true, 0, sizeof(status), &status, 1, &consumer_sync_event, NULL);
- if ( err != CL_SUCCESS ){
- print_error( err, " clEnqueueReadBuffer failed" );
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- clReleaseMemObject(buffers[2]);
- clReleaseMemObject(pipe);
- clReleaseKernel(kernel[0]);
- clReleaseKernel(kernel[1]);
- clReleaseEvent(producer_sync_event);
- clReleaseEvent(consumer_sync_event);
- clReleaseProgram(program);
- align_free(outptr);
- return -1;
- }
+ test_error_ret(err, " clEnqueueReadBuffer failed", -1);
- if(status == 0)
+ if (status == 0)
{
log_error("test_pipe_readwrite_errors failed\n");
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- clReleaseMemObject(buffers[2]);
- clReleaseMemObject(pipe);
- clReleaseKernel(kernel[0]);
- clReleaseKernel(kernel[1]);
- clReleaseEvent(producer_sync_event);
- clReleaseEvent(consumer_sync_event);
- clReleaseProgram(program);
- align_free(outptr);
return -1;
}
+ else
+ {
+ status = 0;
+ }
log_info("test_pipe_readwrite_errors passed\n");
- //cleanup
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- clReleaseMemObject(buffers[2]);
- clReleaseMemObject(pipe);
- clReleaseKernel(kernel[0]);
- clReleaseKernel(kernel[1]);
- clReleaseEvent(producer_sync_event);
- clReleaseEvent(consumer_sync_event);
- clReleaseProgram(program);
- align_free(outptr);
+
return 0;
}
diff --git a/test_conformance/pipes/test_pipe_subgroups.cpp b/test_conformance/pipes/test_pipe_subgroups.cpp
index 35519b0..b41170c 100644
--- a/test_conformance/pipes/test_pipe_subgroups.cpp
+++ b/test_conformance/pipes/test_pipe_subgroups.cpp
@@ -88,30 +88,35 @@
int test_pipe_subgroups_divergence(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
{
- cl_mem pipe;
- cl_mem buffers[3];
- cl_int *outptr;
- cl_int *inptr;
- cl_int *active_work_item_buffer;
- cl_program program;
- cl_kernel kernel[2];
- size_t global_work_size[3];
- size_t local_work_size[3];
- cl_int err;
- cl_int size;
- int i;
- size_t subgroup_count;
- cl_event producer_sync_event = NULL;
- cl_event consumer_sync_event = NULL;
- const char* kernelName[] = {"test_pipe_subgroups_divergence_write", "test_pipe_subgroups_divergence_read"};
+ clMemWrapper pipe;
+ clMemWrapper buffers[3];
+ cl_int *outptr;
+ cl_int *inptr;
+ cl_int *active_work_item_buffer;
+ clProgramWrapper program;
+ clKernelWrapper kernel[2];
+ size_t global_work_size[3];
+ size_t local_work_size[3];
+ cl_int err;
+ cl_int size;
+ int i;
+ size_t subgroup_count;
+ clEventWrapper producer_sync_event = NULL;
+ clEventWrapper consumer_sync_event = NULL;
+ BufferOwningPtr<cl_int> BufferInPtr;
+ BufferOwningPtr<cl_int> BufferOutPtr;
+ const char *kernelName[] = { "test_pipe_subgroups_divergence_write",
+ "test_pipe_subgroups_divergence_read" };
- size_t min_alignment = get_min_alignment(context);
+ size_t min_alignment = get_min_alignment(context);
global_work_size[0] = (cl_uint)num_elements;
- if(!is_extension_available(deviceID, "cl_khr_subgroups"))
+ if (!is_extension_available(deviceID, "cl_khr_subgroups"))
{
- log_info("cl_khr_subgroups is not supported on this platoform. Skipping test.\n");
+ log_info(
+ "cl_khr_subgroups is not supported on this platoform. Skipping "
+ "test.\n");
return CL_SUCCESS;
}
@@ -125,215 +130,77 @@
outptr[i] = 0;
active_work_item_buffer[i] = 0;
}
+ BufferInPtr.reset(inptr, nullptr, 0, size, true);
+ BufferOutPtr.reset(outptr, nullptr, 0, size, true);
buffers[0] = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR, size, inptr, &err);
- if(err){
- clReleaseMemObject(buffers[0]);
- print_error(err, " clCreateBuffer failed\n");
- return -1;
- }
+ test_error_ret(err, " clCreateBuffer failed", -1);
buffers[1] = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR, size, outptr, &err);
- if ( err ){
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- align_free( outptr );
- print_error(err, " clCreateBuffer failed\n" );
- return -1;
- }
+ test_error_ret(err, " clCreateBuffer failed", -1);
buffers[2] = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR, size, active_work_item_buffer, &err);
- if ( err ){
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- clReleaseMemObject(buffers[2]);
- align_free( outptr );
- print_error(err, " clCreateBuffer failed\n" );
- return -1;
- }
+ test_error_ret(err, " clCreateBuffer failed", -1);
pipe = clCreatePipe(context, CL_MEM_HOST_NO_ACCESS, sizeof(int), num_elements, NULL, &err);
- if(err){
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- clReleaseMemObject(buffers[2]);
- align_free( outptr );
- clReleaseMemObject(pipe);
- print_error(err, " clCreatePipe failed\n");
- return -1;
- }
+ test_error_ret(err, " clCreatePipe failed", -1);
// Create producer kernel
err = create_single_kernel_helper_with_build_options(context, &program, &kernel[0], 1, (const char**)&pipe_subgroups_kernel_code, kernelName[0], "-cl-std=CL2.0");
- if(err){
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- clReleaseMemObject(buffers[2]);
- clReleaseMemObject(pipe);
- align_free(outptr);
- print_error(err, "Error creating program\n");
- return -1;
- }
+ test_error_ret(err, " Error creating program", -1);
+
//Create consumer kernel
kernel[1] = clCreateKernel(program, kernelName[1], &err);
- if( kernel[1] == NULL || err != CL_SUCCESS)
- {
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- clReleaseMemObject(buffers[2]);
- clReleaseMemObject(pipe);
- align_free(outptr);
- print_error(err, "Error creating kernel\n");
- return -1;
- }
+ test_error_ret(err, " Error creating kernel", -1);
err = clSetKernelArg(kernel[0], 0, sizeof(cl_mem), (void*)&buffers[0]);
err |= clSetKernelArg(kernel[0], 1, sizeof(cl_mem), (void*)&pipe);
err |= clSetKernelArg(kernel[0], 2, sizeof(cl_mem), (void*)&buffers[2]);
err |= clSetKernelArg(kernel[1], 0, sizeof(cl_mem), (void*)&pipe);
err |= clSetKernelArg(kernel[1], 1, sizeof(cl_mem), (void*)&buffers[1]);
- if ( err != CL_SUCCESS ){
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- clReleaseMemObject(buffers[2]);
- clReleaseMemObject(pipe);
- clReleaseKernel(kernel[0]);
- clReleaseKernel(kernel[1]);
- clReleaseProgram(program);
- align_free(outptr);
- print_error(err, " clSetKernelArg failed");
- return -1;
- }
+ test_error_ret(err, " clSetKernelArg failed", -1);
err = get_max_common_work_group_size( context, kernel[0], global_work_size[0], &local_work_size[0] );
- if( err != CL_SUCCESS)
- {
- test_error( err, "Unable to get work group size to use" );
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- clReleaseMemObject(buffers[2]);
- clReleaseMemObject(pipe);
- clReleaseKernel(kernel[0]);
- clReleaseKernel(kernel[1]);
- clReleaseProgram(program);
- align_free(outptr);
- return -1;
- }
+ test_error_ret(err, " Unable to get work group size to use", -1);
- cl_platform_id platform;
- err = clGetDeviceInfo(deviceID, CL_DEVICE_PLATFORM, sizeof(platform), &platform, NULL);
- clGetKernelSubGroupInfoKHR_fn clGetKernelSubGroupInfoKHR = (clGetKernelSubGroupInfoKHR_fn) clGetExtensionFunctionAddressForPlatform(platform, "clGetKernelSubGroupInfoKHR");
+ cl_platform_id platform;
+ err = clGetDeviceInfo(deviceID, CL_DEVICE_PLATFORM, sizeof(platform),
+ &platform, NULL);
+ test_error_ret(err, " clGetDeviceInfo failed", -1);
+
+ clGetKernelSubGroupInfoKHR_fn clGetKernelSubGroupInfoKHR =
+ (clGetKernelSubGroupInfoKHR_fn)clGetExtensionFunctionAddressForPlatform(
+ platform, "clGetKernelSubGroupInfoKHR");
err = clGetKernelSubGroupInfoKHR(kernel[0], deviceID, CL_KERNEL_SUB_GROUP_COUNT_FOR_NDRANGE_KHR, sizeof(local_work_size[0]), &local_work_size[0], sizeof(subgroup_count), &subgroup_count, NULL);
+ test_error_ret(err, " clGetKernelSubGroupInfoKHR failed", -1);
if(subgroup_count <= 1)
{
log_info("Only 1 subgroup per workgroup for the kernel. Hence no divergence among subgroups possible. Skipping test.\n");
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- clReleaseMemObject(buffers[2]);
- clReleaseMemObject(pipe);
- clReleaseKernel(kernel[0]);
- clReleaseKernel(kernel[1]);
- clReleaseProgram(program);
- align_free(outptr);
return CL_SUCCESS;
}
// Launch Producer kernel
err = clEnqueueNDRangeKernel( queue, kernel[0], 1, NULL, global_work_size, local_work_size, 0, NULL, &producer_sync_event );
- if ( err != CL_SUCCESS ){
- print_error( err, " clEnqueueNDRangeKernel failed" );
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- clReleaseMemObject(buffers[2]);
- clReleaseMemObject(pipe);
- clReleaseKernel(kernel[0]);
- clReleaseKernel(kernel[1]);
- clReleaseEvent(producer_sync_event);
- clReleaseEvent(consumer_sync_event);
- clReleaseProgram(program);
- align_free(outptr);
- return -1;
- }
+ test_error_ret(err, " clEnqueueNDRangeKernel failed", -1);
err = clEnqueueReadBuffer(queue, buffers[2], true, 0, size, active_work_item_buffer, 1, &producer_sync_event, NULL);
- if ( err != CL_SUCCESS ){
- print_error( err, " clEnqueueReadBuffer failed" );
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- clReleaseMemObject(buffers[2]);
- clReleaseMemObject(pipe);
- clReleaseKernel(kernel[0]);
- clReleaseKernel(kernel[1]);
- clReleaseEvent(producer_sync_event);
- clReleaseEvent(consumer_sync_event);
- clReleaseProgram(program);
- align_free(outptr);
- return -1;
- }
-
+ test_error_ret(err, " clEnqueueReadBuffer failed", -1);
// Launch Consumer kernel
err = clEnqueueNDRangeKernel( queue, kernel[1], 1, NULL, global_work_size, local_work_size, 1, &producer_sync_event, &consumer_sync_event );
- if ( err != CL_SUCCESS ){
- print_error( err, " clEnqueueNDRangeKernel failed" );
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- clReleaseMemObject(buffers[2]);
- clReleaseMemObject(pipe);
- clReleaseKernel(kernel[0]);
- clReleaseKernel(kernel[1]);
- clReleaseEvent(producer_sync_event);
- clReleaseEvent(consumer_sync_event);
- clReleaseProgram(program);
- align_free(outptr);
- return -1;
- }
+ test_error_ret(err, " clEnqueueNDRangeKernel failed", -1);
err = clEnqueueReadBuffer(queue, buffers[1], true, 0, size, outptr, 1, &consumer_sync_event, NULL);
- if ( err != CL_SUCCESS ){
- print_error( err, " clEnqueueReadBuffer failed" );
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- clReleaseMemObject(buffers[2]);
- clReleaseMemObject(pipe);
- clReleaseKernel(kernel[0]);
- clReleaseKernel(kernel[1]);
- clReleaseEvent(producer_sync_event);
- clReleaseEvent(consumer_sync_event);
- clReleaseProgram(program);
- align_free(outptr);
- return -1;
- }
+ test_error_ret(err, " clEnqueueReadBuffer failed", -1);
if( verify_result( active_work_item_buffer, outptr, num_elements)){
log_error("test_pipe_subgroups_divergence failed\n");
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- clReleaseMemObject(buffers[2]);
- clReleaseMemObject(pipe);
- clReleaseKernel(kernel[0]);
- clReleaseKernel(kernel[1]);
- clReleaseEvent(producer_sync_event);
- clReleaseEvent(consumer_sync_event);
- clReleaseProgram(program);
- align_free(outptr);
return -1;
}
else {
log_info("test_pipe_subgroups_divergence passed\n");
}
- //cleanup
- clReleaseMemObject(buffers[0]);
- clReleaseMemObject(buffers[1]);
- clReleaseMemObject(buffers[2]);
- clReleaseMemObject(pipe);
- clReleaseKernel(kernel[0]);
- clReleaseKernel(kernel[1]);
- clReleaseEvent(producer_sync_event);
- clReleaseEvent(consumer_sync_event);
- clReleaseProgram(program);
- align_free(outptr);
return 0;
}
