Use CTS type wrappers for test_constant. (#1543)

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

test_conformance/basic/test_constant.cpp

@@ -21,41 +21,44 @@
 #include <sys/types.h>
 #include <sys/stat.h>
 
+#include <algorithm>
+#include <vector>
 
 #include "procs.h"
 
-const char *constant_kernel_code =
+namespace {
-"__kernel void constant_kernel(__global float *out, __constant float *tmpF, __constant int *tmpI)\n"
+const char* constant_kernel_code = R"(
-"{\n"
+__kernel void constant_kernel(__global float *out, __constant float *tmpF, __constant int *tmpI)
-"    int  tid = get_global_id(0);\n"
-"\n"
-"    float ftmp = tmpF[tid]; \n"
-"    float Itmp = tmpI[tid]; \n"
-"    out[tid] = ftmp * Itmp; \n"
-"}\n";
-
-const char *loop_constant_kernel_code =
-"kernel void loop_constant_kernel(global float *out, constant float *i_pos, int num)\n"
-"{\n"
-"    int tid = get_global_id(0);\n"
-"    float sum = 0;\n"
-"    for (int i = 0; i < num; i++) {\n"
-"        float  pos  = i_pos[i*3];\n"
-"        sum += pos;\n"
-"    }\n"
-"    out[tid] = sum;\n"
-"}\n";
-
-
-static int
-verify(cl_float *tmpF, cl_int *tmpI, cl_float *out, int n)
 {
-    int         i;
+    int  tid = get_global_id(0);
 
-    for (i=0; i < n; i++)
+    float ftmp = tmpF[tid];
+    float Itmp = tmpI[tid];
+    out[tid] = ftmp * Itmp;
+}
+)";
+
+const char* loop_constant_kernel_code = R"(
+kernel void loop_constant_kernel(global float *out, constant float *i_pos, int num)
+{
+    int tid = get_global_id(0);
+    float sum = 0;
+    for (int i = 0; i < num; i++) {
+        float  pos  = i_pos[i*3];
+        sum += pos;
+    }
+    out[tid] = sum;
+}
+)";
+
+
+int verify(std::vector<cl_float>& tmpF, std::vector<cl_int>& tmpI,
+           std::vector<cl_float>& out)
+{
+    for (int i = 0; i < out.size(); i++)
     {
         float f = tmpF[i] * tmpI[i];
-        if( out[i] != f )
+        if (out[i] != f)
         {
             log_error("CONSTANT test failed\n");
             return -1;
@@ -66,214 +69,172 @@ verify(cl_float *tmpF, cl_int *tmpI, cl_float *out, int n)
     return 0;
 }
 
-
+int verify_loop_constant(const std::vector<cl_float>& tmp,
-static int
+                         std::vector<cl_float>& out, cl_int l)
-verify_loop_constant(const cl_float *tmp, cl_float *out, cl_int l, int n)
 {
-    int i;
+    float sum = 0;
-    cl_int j;
+    for (int j = 0; j < l; ++j) sum += tmp[j * 3];
-    for (i=0; i < n; i++)
-    {
-        float sum = 0;
-        for (j=0; j < l; ++j)
-            sum += tmp[j*3];
 
-        if( out[i] != sum )
+    auto predicate = [&sum](cl_float elem) { return sum != elem; };
-        {
+
-            log_error("loop CONSTANT test failed\n");
+    if (std::any_of(out.cbegin(), out.cend(), predicate))
-            return -1;
+    {
-        }
+        log_error("loop CONSTANT test failed\n");
+        return -1;
     }
 
     log_info("loop CONSTANT test passed\n");
     return 0;
 }
 
-int
+template <typename T> void generate_random_inputs(std::vector<T>& v)
-test_constant(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
 {
-    cl_mem            streams[3];
+    RandomSeed seed(gRandomSeed);
-    cl_int            *tmpI;
+
-    cl_float        *tmpF, *out;
+    auto random_generator = [&seed]() {
-    cl_program        program;
+        return static_cast<T>(get_random_float(-0x02000000, 0x02000000, seed));
-    cl_kernel        kernel;
+    };
-    size_t    global_threads[3];
+
-    int                err;
+    std::generate(v.begin(), v.end(), random_generator);
-    unsigned int                i;
+}
+}
+
+int test_constant(cl_device_id device, cl_context context,
+                  cl_command_queue queue, int num_elements)
+{
+    clMemWrapper streams[3];
+    clProgramWrapper program;
+    clKernelWrapper kernel;
+
+    size_t global_threads[3];
+    int err;
     cl_ulong maxSize, maxGlobalSize, maxAllocSize;
     size_t num_floats, num_ints, constant_values;
-    MTdata          d;
+    RoundingMode oldRoundMode;
-    RoundingMode     oldRoundMode;
     int isRTZ = 0;
 
-  /* Verify our test buffer won't be bigger than allowed */
+    /* Verify our test buffer won't be bigger than allowed */
-    err = clGetDeviceInfo( device, CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE, sizeof( maxSize ), &maxSize, 0 );
+    err = clGetDeviceInfo(device, CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE,
-    test_error( err, "Unable to get max constant buffer size" );
+                          sizeof(maxSize), &maxSize, 0);
+    test_error(err, "Unable to get max constant buffer size");
+    log_info("Device reports CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE %llu bytes.\n",
+             maxSize);
 
-  log_info("Device reports CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE %llu bytes.\n", maxSize);
+    // Limit test buffer size to 1/4 of CL_DEVICE_GLOBAL_MEM_SIZE
+    err = clGetDeviceInfo(device, CL_DEVICE_GLOBAL_MEM_SIZE,
+                          sizeof(maxGlobalSize), &maxGlobalSize, 0);
+    test_error(err, "Unable to get CL_DEVICE_GLOBAL_MEM_SIZE");
 
-  // Limit test buffer size to 1/4 of CL_DEVICE_GLOBAL_MEM_SIZE
+    maxSize = std::min(maxSize, maxGlobalSize / 4);
-  err = clGetDeviceInfo(device, CL_DEVICE_GLOBAL_MEM_SIZE, sizeof(maxGlobalSize), &maxGlobalSize, 0);
-  test_error(err, "Unable to get CL_DEVICE_GLOBAL_MEM_SIZE");
 
-  if (maxSize > maxGlobalSize / 4)
+    err = clGetDeviceInfo(device, CL_DEVICE_MAX_MEM_ALLOC_SIZE,
-    maxSize = maxGlobalSize / 4;
+                          sizeof(maxAllocSize), &maxAllocSize, 0);
+    test_error(err, "Unable to get CL_DEVICE_MAX_MEM_ALLOC_SIZE");
 
-  err = clGetDeviceInfo(device, CL_DEVICE_MAX_MEM_ALLOC_SIZE , sizeof(maxAllocSize), &maxAllocSize, 0);
+    maxSize = std::min(maxSize, maxAllocSize);
-  test_error(err, "Unable to get CL_DEVICE_MAX_MEM_ALLOC_SIZE ");
 
-  if (maxSize > maxAllocSize)
+    maxSize /= 4;
-    maxSize = maxAllocSize;
+    num_ints = static_cast<size_t>(maxSize / sizeof(cl_int));
+    num_floats = static_cast<size_t>(maxSize / sizeof(cl_float));
+    constant_values = std::min(num_floats, num_ints);
 
-  maxSize/=4;
-  num_ints = (size_t)maxSize/sizeof(cl_int);
-  num_floats = (size_t)maxSize/sizeof(cl_float);
-  if (num_ints >= num_floats) {
-    constant_values = num_floats;
-  } else {
-    constant_values = num_ints;
-  }
 
-  log_info("Test will attempt to use %lu bytes with one %lu byte constant int buffer and one %lu byte constant float buffer.\n",
+    log_info(
-           constant_values*sizeof(cl_int) + constant_values*sizeof(cl_float), constant_values*sizeof(cl_int), constant_values*sizeof(cl_float));
+        "Test will attempt to use %lu bytes with one %lu byte constant int "
+        "buffer and one %lu byte constant float buffer.\n",
+        constant_values * sizeof(cl_int) + constant_values * sizeof(cl_float),
+        constant_values * sizeof(cl_int), constant_values * sizeof(cl_float));
 
-    tmpI = (cl_int*)malloc(sizeof(cl_int) * constant_values);
+    std::vector<cl_int> tmpI(constant_values);
-    tmpF = (cl_float*)malloc(sizeof(cl_float) * constant_values);
+    std::vector<cl_float> tmpF(constant_values);
-    out  = (cl_float*)malloc(sizeof(cl_float) * constant_values);
+    std::vector<cl_float> out(constant_values);
-    streams[0] = clCreateBuffer(context, CL_MEM_READ_WRITE,
-                                sizeof(cl_float) * constant_values, NULL, NULL);
-    if (!streams[0])
-    {
-        log_error("clCreateBuffer failed\n");
-        return -1;
-    }
-    streams[1] = clCreateBuffer(context, CL_MEM_READ_WRITE,
-                                sizeof(cl_float) * constant_values, NULL, NULL);
-    if (!streams[1])
-    {
-        log_error("clCreateBuffer failed\n");
-        return -1;
-    }
-    streams[2] = clCreateBuffer(context, CL_MEM_READ_WRITE,
-                                sizeof(cl_int) * constant_values, NULL, NULL);
-    if (!streams[2])
-    {
-        log_error("clCreateBuffer failed\n");
-        return -1;
-    }
 
-    d = init_genrand( gRandomSeed );
-    for (i=0; i<constant_values; i++) {
-        tmpI[i] = (int)get_random_float(-0x02000000, 0x02000000, d);
-        tmpF[i] = get_random_float(-0x02000000, 0x02000000, d);
-    }
-    free_mtdata(d); d = NULL;
 
-    err = clEnqueueWriteBuffer(queue, streams[1], CL_TRUE, 0, sizeof(cl_float)*constant_values, (void *)tmpF, 0, NULL, NULL);
+    streams[0] =
-    if (err != CL_SUCCESS)
+        clCreateBuffer(context, CL_MEM_READ_WRITE,
-    {
+                       sizeof(cl_float) * constant_values, nullptr, &err);
-        log_error("clWriteArray failed\n");
+    test_error(err, "clCreateBuffer failed");
-        return -1;
-    }
-  err = clEnqueueWriteBuffer(queue, streams[2], CL_TRUE, 0, sizeof(cl_int)*constant_values, (void *)tmpI, 0, NULL, NULL);
-    if (err != CL_SUCCESS)
-    {
-        log_error("clWriteArray failed\n");
-        return -1;
-    }
 
-  err = create_single_kernel_helper(context, &program, &kernel, 1, &constant_kernel_code, "constant_kernel" );
+    streams[1] =
-    if (err) {
+        clCreateBuffer(context, CL_MEM_READ_WRITE,
-    log_error("Failed to create kernel and program: %d\n", err);
+                       sizeof(cl_float) * constant_values, nullptr, &err);
-    return -1;
+    test_error(err, "clCreateBuffer failed");
-  }
+
+    streams[2] =
+        clCreateBuffer(context, CL_MEM_READ_WRITE,
+                       sizeof(cl_int) * constant_values, nullptr, &err);
+    test_error(err, "clCreateBuffer failed");
+
+    generate_random_inputs(tmpI);
+    generate_random_inputs(tmpF);
+
+    err = clEnqueueWriteBuffer(queue, streams[1], CL_TRUE, 0,
+                               sizeof(cl_float) * constant_values, tmpF.data(),
+                               0, nullptr, nullptr);
+    test_error(err, "clEnqueueWriteBuffer failed");
+    err = clEnqueueWriteBuffer(queue, streams[2], CL_TRUE, 0,
+                               sizeof(cl_int) * constant_values, tmpI.data(), 0,
+                               nullptr, nullptr);
+    test_error(err, "clEnqueueWriteBuffer faile.");
+
+    err = create_single_kernel_helper(context, &program, &kernel, 1,
+                                      &constant_kernel_code, "constant_kernel");
+    test_error(err, "Failed to create kernel and program");
 
 
     err = clSetKernelArg(kernel, 0, sizeof streams[0], &streams[0]);
     err |= clSetKernelArg(kernel, 1, sizeof streams[1], &streams[1]);
     err |= clSetKernelArg(kernel, 2, sizeof streams[2], &streams[2]);
-    if (err != CL_SUCCESS)
+    test_error(err, "clSetKernelArgs failed");
-    {
-        log_error("clSetKernelArgs failed\n");
-        return -1;
-    }
 
     global_threads[0] = constant_values;
-    err = clEnqueueNDRangeKernel( queue, kernel, 1, NULL, global_threads, NULL, 0, NULL, NULL );
+    err = clEnqueueNDRangeKernel(queue, kernel, 1, nullptr, global_threads,
-    if (err != CL_SUCCESS)
+                                 nullptr, 0, nullptr, nullptr);
-    {
+    test_error(err, "clEnqueueNDRangeKernel failed");
-        log_error("clEnqueueNDRangeKernel failed: %d\n", err);
-        return -1;
-    }
-    err = clEnqueueReadBuffer( queue, streams[0], CL_TRUE, 0, sizeof(cl_float)*constant_values, (void *)out, 0, NULL, NULL );
-    if (err != CL_SUCCESS)
-    {
-        log_error("clEnqueueReadBuffer failed\n");
-        return -1;
-    }
 
-    //If we only support rtz mode
+    err = clEnqueueReadBuffer(queue, streams[0], CL_TRUE, 0,
-    if( CL_FP_ROUND_TO_ZERO == get_default_rounding_mode(device) && gIsEmbedded)
+                              sizeof(cl_float) * constant_values, out.data(), 0,
+                              nullptr, nullptr);
+    test_error(err, "clEnqueueReadBuffer failed");
+
+    // If we only support rtz mode
+    if (CL_FP_ROUND_TO_ZERO == get_default_rounding_mode(device) && gIsEmbedded)
     {
         oldRoundMode = set_round(kRoundTowardZero, kfloat);
         isRTZ = 1;
     }
 
-    err = verify(tmpF, tmpI, out, (int)constant_values);
+    err = verify(tmpF, tmpI, out);
 
-    if (isRTZ)
+    if (isRTZ) (void)set_round(oldRoundMode, kfloat);
-        (void)set_round(oldRoundMode, kfloat);
 
     // Loop constant buffer test
-    cl_program loop_program;
+    clProgramWrapper loop_program;
-    cl_kernel  loop_kernel;
+    clKernelWrapper loop_kernel;
     cl_int limit = 2;
 
-    memset(out, 0, sizeof(cl_float) * constant_values);
+    memset(out.data(), 0, sizeof(cl_float) * constant_values);
     err = create_single_kernel_helper(context, &loop_program, &loop_kernel, 1,
-                                      &loop_constant_kernel_code, "loop_constant_kernel" );
+                                      &loop_constant_kernel_code,
-    if (err) {
+                                      "loop_constant_kernel");
-        log_error("Failed to create loop kernel and program: %d\n", err);
+    test_error(err, "Failed to create kernel and program");
-        return -1;
-    }
 
     err = clSetKernelArg(loop_kernel, 0, sizeof streams[0], &streams[0]);
     err |= clSetKernelArg(loop_kernel, 1, sizeof streams[1], &streams[1]);
     err |= clSetKernelArg(loop_kernel, 2, sizeof(limit), &limit);
-    if (err != CL_SUCCESS) {
+    test_error(err, "clSetKernelArgs failed");
-        log_error("clSetKernelArgs for loop kernel failed\n");
-        return -1;
-    }
 
-    err = clEnqueueNDRangeKernel( queue, loop_kernel, 1, NULL, global_threads, NULL, 0, NULL, NULL );
+    err = clEnqueueNDRangeKernel(queue, loop_kernel, 1, nullptr, global_threads,
-    if (err != CL_SUCCESS) {
+                                 nullptr, 0, nullptr, nullptr);
-        log_error("clEnqueueNDRangeKernel failed: %d\n", err);
+    test_error(err, "clEnqueueNDRangeKernel failed");
-        return -1;
-    }
-    err = clEnqueueReadBuffer( queue, streams[0], CL_TRUE, 0, sizeof(cl_float)*constant_values, (void *)out, 0, NULL, NULL );
-    if (err != CL_SUCCESS) {
-        log_error("clEnqueueReadBuffer failed\n");
-        return -1;
-    }
 
-    err = verify_loop_constant(tmpF, out, limit, (int)constant_values);
+    err = clEnqueueReadBuffer(queue, streams[0], CL_TRUE, 0,
+                              sizeof(cl_float) * constant_values, out.data(), 0,
+                              nullptr, nullptr);
+    test_error(err, "clEnqueueReadBuffer failed");
+
+    err = verify_loop_constant(tmpF, out, limit);
 
-    // cleanup
-    clReleaseMemObject(streams[0]);
-    clReleaseMemObject(streams[1]);
-    clReleaseMemObject(streams[2]);
-    clReleaseKernel(kernel);
-    clReleaseProgram(program);
-    clReleaseKernel(loop_kernel);
-    clReleaseProgram(loop_program);
-    free(tmpI);
-    free(tmpF);
-    free(out);
 
     return err;
 }
-
-
-
-
-