diff --git a/test_conformance/basic/CMakeLists.txt b/test_conformance/basic/CMakeLists.txt
index 3f84d7a4..d73b84a2 100644
--- a/test_conformance/basic/CMakeLists.txt
+++ b/test_conformance/basic/CMakeLists.txt
@@ -3,8 +3,7 @@ set(MODULE_NAME BASIC)
 set(${MODULE_NAME}_SOURCES
     main.cpp
     test_fpmath_float.cpp test_fpmath_float2.cpp test_fpmath_float4.cpp
-    test_intmath_int.cpp test_intmath_int2.cpp test_intmath_int4.cpp
-    test_intmath_long.cpp test_intmath_long2.cpp test_intmath_long4.cpp
+    test_intmath.cpp
     test_hiloeo.cpp test_local.cpp test_pointercast.cpp
     test_if.cpp test_loop.cpp
     test_readimage.cpp test_readimage_int16.cpp test_readimage_fp32.cpp
diff --git a/test_conformance/basic/test_intmath.cpp b/test_conformance/basic/test_intmath.cpp
new file mode 100644
index 00000000..6fd41abb
--- /dev/null
+++ b/test_conformance/basic/test_intmath.cpp
@@ -0,0 +1,240 @@
+//
+// Copyright (c) 2020 The Khronos Group Inc.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//    http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+//
+
+#include <functional>
+#include <string>
+#include <vector>
+
+#include "procs.h"
+
+template <typename T> struct TestDef
+{
+    const char *name;
+    const char *kernel_code;
+    std::function<T(T, T, T)> ref;
+};
+
+template <typename T, unsigned N>
+int test_intmath(cl_device_id device, cl_context context,
+                 cl_command_queue queue, int num_elements, std::string typestr)
+{
+    TestDef<T> tests[] = {
+        // Test addition
+        {
+            "test_add",
+            R"(
+  __kernel void test_add(__global TYPE *srcA,
+                         __global TYPE *srcB,
+                         __global TYPE *srcC,
+                         __global TYPE *dst)
+  {
+      int  tid = get_global_id(0);
+      dst[tid] = srcA[tid] + srcB[tid];
+  };
+)",
+            [](T a, T b, T c) { return a + b; },
+        },
+
+        // Test subtraction
+        {
+            "test_sub",
+            R"(
+  __kernel void test_sub(__global TYPE *srcA,
+                         __global TYPE *srcB,
+                         __global TYPE *srcC,
+                         __global TYPE *dst)
+  {
+      int  tid = get_global_id(0);
+      dst[tid] = srcA[tid] - srcB[tid];
+  };
+)",
+            [](T a, T b, T c) { return a - b; },
+        },
+
+        // Test multiplication
+        {
+            "test_mul",
+            R"(
+  __kernel void test_mul(__global TYPE *srcA,
+                         __global TYPE *srcB,
+                         __global TYPE *srcC,
+                         __global TYPE *dst)
+  {
+      int  tid = get_global_id(0);
+      dst[tid] = srcA[tid] * srcB[tid];
+  };
+)",
+            [](T a, T b, T c) { return a * b; },
+        },
+
+        // Test multiply-accumulate
+        {
+            "test_mad",
+            R"(
+  __kernel void test_mad(__global TYPE *srcA,
+                         __global TYPE *srcB,
+                         __global TYPE *srcC,
+                         __global TYPE *dst)
+  {
+      int  tid = get_global_id(0);
+      dst[tid] = srcA[tid] * srcB[tid] + srcC[tid];
+  };
+)",
+            [](T a, T b, T c) { return a * b + c; },
+        },
+    };
+
+    clMemWrapper streams[4];
+    cl_int err;
+
+    if (std::is_same<T, cl_ulong>::value && !gHasLong)
+    {
+        log_info("64-bit integers are not supported on this device. Skipping "
+                 "test.\n");
+        return TEST_SKIPPED_ITSELF;
+    }
+
+    // Create host buffers and fill with random data.
+    std::vector<T> inputA(num_elements * N);
+    std::vector<T> inputB(num_elements * N);
+    std::vector<T> inputC(num_elements * N);
+    std::vector<T> output(num_elements * N);
+    MTdataHolder d(gRandomSeed);
+    for (int i = 0; i < num_elements; i++)
+    {
+        inputA[i] = (T)genrand_int64(d);
+        inputB[i] = (T)genrand_int64(d);
+        inputC[i] = (T)genrand_int64(d);
+    }
+
+    size_t datasize = sizeof(T) * num_elements * N;
+
+    // Create device buffers.
+    for (int i = 0; i < ARRAY_SIZE(streams); i++)
+    {
+        streams[i] =
+            clCreateBuffer(context, CL_MEM_READ_WRITE, datasize, NULL, &err);
+        test_error(err, "clCreateBuffer failed");
+    }
+
+    // Copy input data to device.
+    err = clEnqueueWriteBuffer(queue, streams[0], CL_TRUE, 0, datasize,
+                               inputA.data(), 0, NULL, NULL);
+    test_error(err, "clEnqueueWriteBuffer failed\n");
+    err = clEnqueueWriteBuffer(queue, streams[1], CL_TRUE, 0, datasize,
+                               inputB.data(), 0, NULL, NULL);
+    test_error(err, "clEnqueueWriteBuffer failed\n");
+    err = clEnqueueWriteBuffer(queue, streams[2], CL_TRUE, 0, datasize,
+                               inputC.data(), 0, NULL, NULL);
+    test_error(err, "clEnqueueWriteBuffer failed\n");
+
+    std::string build_options = "-DTYPE=";
+    build_options += typestr;
+
+    // Run test for each operation
+    for (auto test : tests)
+    {
+        log_info("%s... ", test.name);
+
+        // Create kernel and set args
+        clProgramWrapper program;
+        clKernelWrapper kernel;
+        err = create_single_kernel_helper(context, &program, &kernel, 1,
+                                          &test.kernel_code, test.name,
+                                          build_options.c_str());
+        test_error(err, "create_single_kernel_helper failed\n");
+
+        err = clSetKernelArg(kernel, 0, sizeof(cl_mem), &streams[0]);
+        err |= clSetKernelArg(kernel, 1, sizeof(cl_mem), &streams[1]);
+        err |= clSetKernelArg(kernel, 2, sizeof(cl_mem), &streams[2]);
+        err |= clSetKernelArg(kernel, 3, sizeof(cl_mem), &streams[3]);
+        test_error(err, "clSetKernelArgs failed\n");
+
+        // Run kernel
+        size_t threads[1] = { static_cast<size_t>(num_elements) };
+        err = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, threads, NULL, 0,
+                                     NULL, NULL);
+        test_error(err, "clEnqueueNDRangeKernel failed\n");
+
+        // Read results
+        err = clEnqueueReadBuffer(queue, streams[3], CL_TRUE, 0, datasize,
+                                  output.data(), 0, NULL, NULL);
+        test_error(err, "clEnqueueReadBuffer failed\n");
+
+        // Verify results
+        for (int i = 0; i < num_elements * N; i++)
+        {
+            T r = test.ref(inputA[i], inputB[i], inputC[i]);
+            if (r != output[i])
+            {
+                log_error("\n\nverification failed at index %d\n", i);
+                log_error("-> inputs: %llu, %llu, %llu\n",
+                          static_cast<cl_uint>(inputA[i]),
+                          static_cast<cl_uint>(inputB[i]),
+                          static_cast<cl_uint>(inputC[i]));
+                log_error("-> expected %llu, got %llu\n\n",
+                          static_cast<cl_uint>(r),
+                          static_cast<cl_uint>(output[i]));
+                return TEST_FAIL;
+            }
+        }
+        log_info("passed\n");
+    }
+
+    return TEST_PASS;
+}
+
+int test_intmath_int(cl_device_id device, cl_context context,
+                     cl_command_queue queue, int num_elements)
+{
+    return test_intmath<cl_uint, 1>(device, context, queue, num_elements,
+                                    "uint");
+}
+
+int test_intmath_int2(cl_device_id device, cl_context context,
+                      cl_command_queue queue, int num_elements)
+{
+    return test_intmath<cl_uint, 2>(device, context, queue, num_elements,
+                                    "uint2");
+}
+
+int test_intmath_int4(cl_device_id device, cl_context context,
+                      cl_command_queue queue, int num_elements)
+{
+    return test_intmath<cl_uint, 4>(device, context, queue, num_elements,
+                                    "uint4");
+}
+
+int test_intmath_long(cl_device_id device, cl_context context,
+                      cl_command_queue queue, int num_elements)
+{
+    return test_intmath<cl_ulong, 1>(device, context, queue, num_elements,
+                                     "ulong");
+}
+
+int test_intmath_long2(cl_device_id device, cl_context context,
+                       cl_command_queue queue, int num_elements)
+{
+    return test_intmath<cl_ulong, 2>(device, context, queue, num_elements,
+                                     "ulong2");
+}
+
+int test_intmath_long4(cl_device_id device, cl_context context,
+                       cl_command_queue queue, int num_elements)
+{
+    return test_intmath<cl_ulong, 4>(device, context, queue, num_elements,
+                                     "ulong4");
+}
diff --git a/test_conformance/basic/test_intmath_int.cpp b/test_conformance/basic/test_intmath_int.cpp
deleted file mode 100644
index 76676617..00000000
--- a/test_conformance/basic/test_intmath_int.cpp
+++ /dev/null
@@ -1,334 +0,0 @@
-//
-// Copyright (c) 2017 The Khronos Group Inc.
-// 
-// Licensed under the Apache License, Version 2.0 (the "License");
-// you may not use this file except in compliance with the License.
-// You may obtain a copy of the License at
-//
-//    http://www.apache.org/licenses/LICENSE-2.0
-//
-// Unless required by applicable law or agreed to in writing, software
-// distributed under the License is distributed on an "AS IS" BASIS,
-// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
-// See the License for the specific language governing permissions and
-// limitations under the License.
-//
-#include "harness/compat.h"
-
-#include <stdio.h>
-#include <stdlib.h>
-#include <string.h>
-#include <sys/types.h>
-#include <sys/stat.h>
-
-
-#include "procs.h"
-
-const char *int_add_kernel_code =
-"__kernel void test_int_add(__global int *srcA, __global int *srcB, __global int *dst)\n"
-"{\n"
-"    int  tid = get_global_id(0);\n"
-"\n"
-"    dst[tid] = srcA[tid] + srcB[tid];\n"
-"}\n";
-
-const char *int_sub_kernel_code =
-"__kernel void test_int_sub(__global int *srcA, __global int *srcB, __global int *dst)\n"
-"{\n"
-"    int  tid = get_global_id(0);\n"
-"\n"
-"    dst[tid] = srcA[tid] - srcB[tid];\n"
-"}\n";
-
-const char *int_mul_kernel_code =
-"__kernel void test_int_mul(__global int *srcA, __global int *srcB, __global int *dst)\n"
-"{\n"
-"    int  tid = get_global_id(0);\n"
-"\n"
-"    dst[tid] = srcA[tid] * srcB[tid];\n"
-"}\n";
-
-const char *int_mad_kernel_code =
-"__kernel void test_int_mad(__global int *srcA, __global int *srcB, __global int *srcC, __global int *dst)\n"
-"{\n"
-"    int  tid = get_global_id(0);\n"
-"\n"
-"    dst[tid] = srcA[tid] * srcB[tid] + srcC[tid];\n"
-"}\n";
-
-static const float    MAX_ERR = 1e-5f;
-
-int
-verify_int_add(int *inptrA, int *inptrB, int *outptr, int n)
-{
-    int            r;
-    int         i;
-
-    for (i=0; i<n; i++)
-    {
-        r = inptrA[i] + inptrB[i];
-        if (r != outptr[i])
-        {
-            log_error("INT_ADD int test failed\n");
-            return -1;
-        }
-    }
-
-    log_info("INT_ADD int test passed\n");
-    return 0;
-}
-
-int
-verify_int_sub(int *inptrA, int *inptrB, int *outptr, int n)
-{
-    int            r;
-    int         i;
-
-    for (i=0; i<n; i++)
-    {
-        r = inptrA[i] - inptrB[i];
-        if (r != outptr[i])
-        {
-            log_error("INT_SUB int test failed\n");
-            return -1;
-        }
-    }
-
-    log_info("INT_SUB int test passed\n");
-    return 0;
-}
-
-int
-verify_int_mul(int *inptrA, int *inptrB, int *outptr, int n)
-{
-    int            r;
-    int         i;
-
-    for (i=0; i<n; i++)
-    {
-        r = inptrA[i] * inptrB[i];
-        if (r != outptr[i])
-        {
-            log_error("INT_MUL int test failed\n");
-            return -1;
-        }
-    }
-
-    log_info("INT_MUL int test passed\n");
-    return 0;
-}
-
-int
-verify_int_mad(int *inptrA, int *inptrB, int *inptrC, int *outptr, int n)
-{
-    int            r;
-    int         i;
-
-    for (i=0; i<n; i++)
-    {
-        r = inptrA[i] * inptrB[i] + inptrC[i];
-        if (r != outptr[i])
-        {
-            log_error("INT_MAD int test failed\n");
-            return -1;
-        }
-    }
-
-    log_info("INT_MAD int test passed\n");
-    return 0;
-}
-
-int
-test_intmath_int(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
-{
-    cl_mem streams[4];
-    cl_program program[4];
-    cl_kernel kernel[4];
-
-    cl_int *input_ptr[3], *output_ptr, *p;
-    size_t threads[1];
-    int err, i;
-    MTdata d = init_genrand( gRandomSeed );
-
-    size_t length = sizeof(cl_int) * num_elements;
-
-    input_ptr[0] = (cl_int*)malloc(length);
-    input_ptr[1] = (cl_int*)malloc(length);
-    input_ptr[2] = (cl_int*)malloc(length);
-    output_ptr   = (cl_int*)malloc(length);
-
-    streams[0] = clCreateBuffer(context, CL_MEM_READ_WRITE, length, NULL, NULL);
-    if (!streams[0])
-    {
-        log_error("clCreateBuffer failed\n");
-        return -1;
-    }
-    streams[1] = clCreateBuffer(context, CL_MEM_READ_WRITE, length, NULL, NULL);
-    if (!streams[1])
-    {
-        log_error("clCreateBuffer failed\n");
-        return -1;
-    }
-    streams[2] = clCreateBuffer(context, CL_MEM_READ_WRITE, length, NULL, NULL);
-    if (!streams[2])
-    {
-        log_error("clCreateBuffer failed\n");
-        return -1;
-    }
-    streams[3] = clCreateBuffer(context, CL_MEM_READ_WRITE, length, NULL, NULL);
-    if (!streams[3])
-    {
-        log_error("clCreateBuffer failed\n");
-        return -1;
-    }
-
-    p = input_ptr[0];
-    for (i=0; i<num_elements; i++)
-        p[i] = (int)genrand_int32(d);
-    p = input_ptr[1];
-    for (i=0; i<num_elements; i++)
-        p[i] = (int)genrand_int32(d);
-    p = input_ptr[2];
-    for (i=0; i<num_elements; i++)
-        p[i] = (int)genrand_int32(d);
-
-    free_mtdata(d); d = NULL;
-
-    err = clEnqueueWriteBuffer(queue, streams[0], CL_TRUE, 0, length, input_ptr[0], 0, NULL, NULL);
-    if (err != CL_SUCCESS)
-    {
-        log_error("clEnqueueWriteBuffer failed\n");
-        return -1;
-    }
-    err = clEnqueueWriteBuffer(queue, streams[1], CL_TRUE, 0, length, input_ptr[1], 0, NULL, NULL);
-    if (err != CL_SUCCESS)
-    {
-        log_error("clEnqueueWriteBuffer failed\n");
-        return -1;
-    }
-    err = clEnqueueWriteBuffer(queue, streams[2], CL_TRUE, 0, length, input_ptr[2], 0, NULL, NULL);
-    if (err != CL_SUCCESS)
-    {
-        log_error("clEnqueueWriteBuffer failed\n");
-        return -1;
-    }
-
-    err = create_single_kernel_helper(context, &program[0], &kernel[0], 1, &int_add_kernel_code, "test_int_add");
-    if (err != CL_SUCCESS)
-    {
-        log_error("create_single_kernel_helper failed\n");
-        return -1;
-    }
-
-    err = create_single_kernel_helper(context, &program[1], &kernel[1], 1, &int_sub_kernel_code, "test_int_sub");
-    if (err != CL_SUCCESS)
-    {
-        log_error("create_single_kernel_helper failed\n");
-        return -1;
-    }
-
-    err = create_single_kernel_helper(context, &program[2], &kernel[2], 1, &int_mul_kernel_code, "test_int_mul");
-    if (err != CL_SUCCESS)
-    {
-        log_error("create_single_kernel_helper failed\n");
-        return -1;
-    }
-
-    err = create_single_kernel_helper(context, &program[3], &kernel[3], 1, &int_mad_kernel_code, "test_int_mad");
-    if (err != CL_SUCCESS)
-    {
-        log_error("create_single_kernel_helper failed\n");
-        return -1;
-    }
-
-  err  = clSetKernelArg(kernel[0], 0, sizeof streams[0], &streams[0]);
-  err |= clSetKernelArg(kernel[0], 1, sizeof streams[1], &streams[1]);
-  err |= clSetKernelArg(kernel[0], 2, sizeof streams[3], &streams[3]);
-    if (err != CL_SUCCESS)
-    {
-        log_error("clSetKernelArgs failed\n");
-        return -1;
-    }
-
-  err  = clSetKernelArg(kernel[1], 0, sizeof streams[0], &streams[0]);
-  err |= clSetKernelArg(kernel[1], 1, sizeof streams[1], &streams[1]);
-  err |= clSetKernelArg(kernel[1], 2, sizeof streams[3], &streams[3]);
-    if (err != CL_SUCCESS)
-    {
-        log_error("clSetKernelArgs failed\n");
-        return -1;
-    }
-
-  err  = clSetKernelArg(kernel[2], 0, sizeof streams[0], &streams[0]);
-  err |= clSetKernelArg(kernel[2], 1, sizeof streams[1], &streams[1]);
-  err |= clSetKernelArg(kernel[2], 2, sizeof streams[3], &streams[3]);
-    if (err != CL_SUCCESS)
-    {
-        log_error("clSetKernelArgs failed\n");
-        return -1;
-    }
-
-  err  = clSetKernelArg(kernel[3], 0, sizeof streams[0], &streams[0]);
-  err |= clSetKernelArg(kernel[3], 1, sizeof streams[1], &streams[1]);
-  err |= clSetKernelArg(kernel[3], 2, sizeof streams[2], &streams[2]);
-  err |= clSetKernelArg(kernel[3], 3, sizeof streams[3], &streams[3]);
-    if (err != CL_SUCCESS)
-    {
-        log_error("clSetKernelArgs failed\n");
-        return -1;
-    }
-
-  threads[0] = (unsigned int)num_elements;
-  for (i=0; i<4; i++)
-  {
-    err = clEnqueueNDRangeKernel(queue, kernel[i], 1, NULL, threads, NULL, 0, NULL, NULL);
-    if (err != CL_SUCCESS)
-    {
-      log_error("clEnqueueNDRangeKernel failed\n");
-      return -1;
-    }
-
-    err = clEnqueueReadBuffer(queue, streams[3], CL_TRUE, 0, length, output_ptr, 0, NULL, NULL);
-    if (err != CL_SUCCESS)
-    {
-      log_error("clEnqueueReadBuffer failed\n");
-      return -1;
-    }
-
-    switch (i)
-    {
-      case 0:
-        err = verify_int_add(input_ptr[0], input_ptr[1], output_ptr, num_elements);
-        break;
-      case 1:
-        err = verify_int_sub(input_ptr[0], input_ptr[1], output_ptr, num_elements);
-        break;
-      case 2:
-        err = verify_int_mul(input_ptr[0], input_ptr[1], output_ptr, num_elements);
-        break;
-      case 3:
-        err = verify_int_mad(input_ptr[0], input_ptr[1], input_ptr[2], output_ptr, num_elements);
-        break;
-    }
-    if (err)
-      break;
-    }
-
-    // cleanup
-    clReleaseMemObject(streams[0]);
-    clReleaseMemObject(streams[1]);
-    clReleaseMemObject(streams[2]);
-    clReleaseMemObject(streams[3]);
-    for (i=0; i<4; i++)
-    {
-        clReleaseKernel(kernel[i]);
-        clReleaseProgram(program[i]);
-    }
-    free(input_ptr[0]);
-    free(input_ptr[1]);
-    free(input_ptr[2]);
-    free(output_ptr);
-    return err;
-}
-
-
diff --git a/test_conformance/basic/test_intmath_int2.cpp b/test_conformance/basic/test_intmath_int2.cpp
deleted file mode 100644
index b4e98e10..00000000
--- a/test_conformance/basic/test_intmath_int2.cpp
+++ /dev/null
@@ -1,333 +0,0 @@
-//
-// Copyright (c) 2017 The Khronos Group Inc.
-// 
-// Licensed under the Apache License, Version 2.0 (the "License");
-// you may not use this file except in compliance with the License.
-// You may obtain a copy of the License at
-//
-//    http://www.apache.org/licenses/LICENSE-2.0
-//
-// Unless required by applicable law or agreed to in writing, software
-// distributed under the License is distributed on an "AS IS" BASIS,
-// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
-// See the License for the specific language governing permissions and
-// limitations under the License.
-//
-#include "harness/compat.h"
-
-#include <stdio.h>
-#include <stdlib.h>
-#include <string.h>
-#include <sys/types.h>
-#include <sys/stat.h>
-
-
-#include "procs.h"
-
-const char *int_add2_kernel_code =
-"__kernel void test_int_add2(__global int2 *srcA, __global int2 *srcB, __global int2 *dst)\n"
-"{\n"
-"    int  tid = get_global_id(0);\n"
-"\n"
-"    dst[tid] = srcA[tid] + srcB[tid];\n"
-"}\n";
-
-const char *int_sub2_kernel_code =
-"__kernel void test_int_sub2(__global int2 *srcA, __global int2 *srcB, __global int2 *dst)\n"
-"{\n"
-"    int  tid = get_global_id(0);\n"
-"\n"
-"    dst[tid] = srcA[tid] - srcB[tid];\n"
-"}\n";
-
-const char *int_mul2_kernel_code =
-"__kernel void test_int_mul2(__global int2 *srcA, __global int2 *srcB, __global int2 *dst)\n"
-"{\n"
-"    int  tid = get_global_id(0);\n"
-"\n"
-"    dst[tid] = srcA[tid] * srcB[tid];\n"
-"}\n";
-
-const char *int_mad2_kernel_code =
-"__kernel void test_int_mad2(__global int2 *srcA, __global int2 *srcB, __global int2 *srcC, __global int2 *dst)\n"
-"{\n"
-"    int  tid = get_global_id(0);\n"
-"\n"
-"    dst[tid] = srcA[tid] * srcB[tid] + srcC[tid];\n"
-"}\n";
-
-int
-verify_int_add2(int *inptrA, int *inptrB, int *outptr, int n)
-{
-    int            r;
-    int         i;
-
-    for (i=0; i<n; i++)
-    {
-        r = inptrA[i] + inptrB[i];
-        if (r != outptr[i])
-        {
-            log_error("INT_ADD int2 test failed\n");
-            return -1;
-        }
-    }
-
-    log_info("INT_ADD int2 test passed\n");
-    return 0;
-}
-
-int
-verify_int_sub2(int *inptrA, int *inptrB, int *outptr, int n)
-{
-    int            r;
-    int         i;
-
-    for (i=0; i<n; i++)
-    {
-        r = inptrA[i] - inptrB[i];
-        if (r != outptr[i])
-        {
-            log_error("INT_SUB int2 test failed\n");
-            return -1;
-        }
-    }
-
-    log_info("INT_SUB int2 test passed\n");
-    return 0;
-}
-
-int
-verify_int_mul2(int *inptrA, int *inptrB, int *outptr, int n)
-{
-    int            r;
-    int         i;
-
-    for (i=0; i<n; i++)
-    {
-        r = inptrA[i] * inptrB[i];
-        if (r != outptr[i])
-        {
-            log_error("INT_MUL int2 test failed\n");
-            return -1;
-        }
-    }
-
-    log_info("INT_MUL int2 test passed\n");
-    return 0;
-}
-
-int
-verify_int_mad2(int *inptrA, int *inptrB, int *inptrC, int *outptr, int n)
-{
-    int            r;
-    int         i;
-
-    for (i=0; i<n; i++)
-    {
-        r = inptrA[i] * inptrB[i] + inptrC[i];
-        if (r != outptr[i])
-        {
-            log_error("INT_MAD int2 test failed\n");
-            return -1;
-        }
-    }
-
-    log_info("INT_MAD int2 test passed\n");
-    return 0;
-}
-
-int
-test_intmath_int2(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
-{
-    cl_mem streams[4];
-    cl_program program[4];
-    cl_kernel kernel[4];
-
-    cl_int *input_ptr[3], *output_ptr, *p;
-    size_t threads[1];
-    int err, i;
-    MTdata d = init_genrand( gRandomSeed );
-
-    size_t length = sizeof(cl_int) * 2 * num_elements;
-
-    input_ptr[0] = (cl_int*)malloc(length);
-    input_ptr[1] = (cl_int*)malloc(length);
-    input_ptr[2] = (cl_int*)malloc(length);
-    output_ptr   = (cl_int*)malloc(length);
-
-    streams[0] = clCreateBuffer(context, CL_MEM_READ_WRITE, length, NULL, NULL);
-    if (!streams[0])
-    {
-        log_error("clCreateBuffer failed\n");
-        return -1;
-    }
-    streams[1] = clCreateBuffer(context, CL_MEM_READ_WRITE, length, NULL, NULL);
-    if (!streams[1])
-    {
-        log_error("clCreateBuffer failed\n");
-        return -1;
-    }
-    streams[2] = clCreateBuffer(context, CL_MEM_READ_WRITE, length, NULL, NULL);
-    if (!streams[2])
-    {
-        log_error("clCreateBuffer failed\n");
-        return -1;
-    }
-    streams[3] = clCreateBuffer(context, CL_MEM_READ_WRITE, length, NULL, NULL);
-    if (!streams[3])
-    {
-        log_error("clCreateBuffer failed\n");
-        return -1;
-    }
-
-    p = input_ptr[0];
-    for (i=0; i<num_elements*2; i++)
-        p[i] = (int)genrand_int32(d);
-    p = input_ptr[1];
-    for (i=0; i<num_elements*2; i++)
-        p[i] = (int)genrand_int32(d);
-    p = input_ptr[2];
-    for (i=0; i<num_elements*2; i++)
-        p[i] = (int)genrand_int32(d);
-
-    free_mtdata( d );
-    d = NULL;
-
-    err = clEnqueueWriteBuffer(queue, streams[0], CL_TRUE, 0, length, input_ptr[0], 0, NULL, NULL);
-    if (err != CL_SUCCESS)
-    {
-        log_error("clEnqueueWriteBuffer failed\n");
-        return -1;
-    }
-    err = clEnqueueWriteBuffer(queue, streams[1], CL_TRUE, 0, length, input_ptr[1], 0, NULL, NULL);
-    if (err != CL_SUCCESS)
-    {
-        log_error("clEnqueueWriteBuffer failed\n");
-        return -1;
-    }
-    err = clEnqueueWriteBuffer(queue, streams[2], CL_TRUE, 0, length, input_ptr[2], 0, NULL, NULL);
-    if (err != CL_SUCCESS)
-    {
-        log_error("clEnqueueWriteBuffer failed\n");
-        return -1;
-    }
-
-    err = create_single_kernel_helper(context, &program[0], &kernel[0], 1, &int_add2_kernel_code, "test_int_add2");
-    if (err != CL_SUCCESS)
-    {
-        log_error("create_single_kernel_helper failed\n");
-        return -1;
-    }
-
-    err = create_single_kernel_helper(context, &program[1], &kernel[1], 1, &int_sub2_kernel_code, "test_int_sub2");
-    if (err != CL_SUCCESS)
-    {
-        log_error("create_single_kernel_helper failed\n");
-        return -1;
-    }
-
-    err = create_single_kernel_helper(context, &program[2], &kernel[2], 1, &int_mul2_kernel_code, "test_int_mul2");
-    if (err != CL_SUCCESS)
-    {
-        log_error("create_single_kernel_helper failed\n");
-        return -1;
-    }
-
-    err = create_single_kernel_helper(context, &program[3], &kernel[3], 1, &int_mad2_kernel_code, "test_int_mad2");
-    if (err != CL_SUCCESS)
-    {
-        log_error("create_single_kernel_helper failed\n");
-        return -1;
-    }
-
-  err  = clSetKernelArg(kernel[0], 0, sizeof streams[0], &streams[0]);
-  err |= clSetKernelArg(kernel[0], 1, sizeof streams[1], &streams[1]);
-  err |= clSetKernelArg(kernel[0], 2, sizeof streams[3], &streams[3]);
-    if (err != CL_SUCCESS)
-    {
-        log_error("clSetKernelArgs failed\n");
-        return -1;
-    }
-
-  err  = clSetKernelArg(kernel[1], 0, sizeof streams[0], &streams[0]);
-  err |= clSetKernelArg(kernel[1], 1, sizeof streams[1], &streams[1]);
-  err |= clSetKernelArg(kernel[1], 2, sizeof streams[3], &streams[3]);
-    if (err != CL_SUCCESS)
-    {
-        log_error("clSetKernelArgs failed\n");
-        return -1;
-    }
-
-  err  = clSetKernelArg(kernel[2], 0, sizeof streams[0], &streams[0]);
-  err |= clSetKernelArg(kernel[2], 1, sizeof streams[1], &streams[1]);
-  err |= clSetKernelArg(kernel[2], 2, sizeof streams[3], &streams[3]);
-    if (err != CL_SUCCESS)
-    {
-        log_error("clSetKernelArgs failed\n");
-        return -1;
-    }
-
-  err  = clSetKernelArg(kernel[3], 0, sizeof streams[0], &streams[0]);
-  err |= clSetKernelArg(kernel[3], 1, sizeof streams[1], &streams[1]);
-  err |= clSetKernelArg(kernel[3], 2, sizeof streams[2], &streams[2]);
-  err |= clSetKernelArg(kernel[3], 3, sizeof streams[3], &streams[3]);
-    if (err != CL_SUCCESS)
-    {
-        log_error("clSetKernelArgs failed\n");
-        return -1;
-    }
-
-  threads[0] = (unsigned int)num_elements;
-  for (i=0; i<4; i++)
-  {
-    err = clEnqueueNDRangeKernel(queue, kernel[i], 1, NULL, threads, NULL, 0, NULL, NULL);
-    if (err != CL_SUCCESS)
-    {
-      log_error("clEnqueueNDRangeKernel failed\n");
-      return -1;
-    }
-
-    err = clEnqueueReadBuffer(queue, streams[3], CL_TRUE, 0, length, output_ptr, 0, NULL, NULL);
-    if (err != CL_SUCCESS)
-    {
-      log_error("clEnqueueReadBuffer failed\n");
-      return -1;
-    }
-
-    switch (i)
-    {
-      case 0:
-        err = verify_int_add2(input_ptr[0], input_ptr[1], output_ptr, num_elements);
-        break;
-      case 1:
-        err = verify_int_sub2(input_ptr[0], input_ptr[1], output_ptr, num_elements);
-        break;
-      case 2:
-        err = verify_int_mul2(input_ptr[0], input_ptr[1], output_ptr, num_elements);
-        break;
-      case 3:
-        err = verify_int_mad2(input_ptr[0], input_ptr[1], input_ptr[2], output_ptr, num_elements);
-        break;
-    }
-    if (err)
-      break;
-    }
-
-    // cleanup
-    clReleaseMemObject(streams[0]);
-    clReleaseMemObject(streams[1]);
-    clReleaseMemObject(streams[2]);
-    clReleaseMemObject(streams[3]);
-    for (i=0; i<4; i++)
-    {
-        clReleaseKernel(kernel[i]);
-        clReleaseProgram(program[i]);
-    }
-    free(input_ptr[0]);
-    free(input_ptr[1]);
-    free(input_ptr[2]);
-    free(output_ptr);
-    return err;
-}
-
-
diff --git a/test_conformance/basic/test_intmath_int4.cpp b/test_conformance/basic/test_intmath_int4.cpp
deleted file mode 100644
index 7be9da8d..00000000
--- a/test_conformance/basic/test_intmath_int4.cpp
+++ /dev/null
@@ -1,332 +0,0 @@
-//
-// Copyright (c) 2017 The Khronos Group Inc.
-// 
-// Licensed under the Apache License, Version 2.0 (the "License");
-// you may not use this file except in compliance with the License.
-// You may obtain a copy of the License at
-//
-//    http://www.apache.org/licenses/LICENSE-2.0
-//
-// Unless required by applicable law or agreed to in writing, software
-// distributed under the License is distributed on an "AS IS" BASIS,
-// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
-// See the License for the specific language governing permissions and
-// limitations under the License.
-//
-#include "harness/compat.h"
-
-#include <stdio.h>
-#include <stdlib.h>
-#include <string.h>
-#include <sys/types.h>
-#include <sys/stat.h>
-
-
-#include "procs.h"
-
-const char *int_add4_kernel_code =
-"__kernel void test_int_add4(__global int4 *srcA, __global int4 *srcB, __global int4 *dst)\n"
-"{\n"
-"    int  tid = get_global_id(0);\n"
-"\n"
-"    dst[tid] = srcA[tid] + srcB[tid];\n"
-"}\n";
-
-const char *int_sub4_kernel_code =
-"__kernel void test_int_sub4(__global int4 *srcA, __global int4 *srcB, __global int4 *dst)\n"
-"{\n"
-"    int  tid = get_global_id(0);\n"
-"\n"
-"    dst[tid] = srcA[tid] - srcB[tid];\n"
-"}\n";
-
-const char *int_mul4_kernel_code =
-"__kernel void test_int_mul4(__global int4 *srcA, __global int4 *srcB, __global int4 *dst)\n"
-"{\n"
-"    int  tid = get_global_id(0);\n"
-"\n"
-"    dst[tid] = srcA[tid] * srcB[tid];\n"
-"}\n";
-
-const char *int_mad4_kernel_code =
-"__kernel void test_int_mad4(__global int4 *srcA, __global int4 *srcB, __global int4 *srcC, __global int4 *dst)\n"
-"{\n"
-"    int  tid = get_global_id(0);\n"
-"\n"
-"    dst[tid] = srcA[tid] * srcB[tid] + srcC[tid];\n"
-"}\n";
-
-int
-verify_int_add4(int *inptrA, int *inptrB, int *outptr, int n)
-{
-    int            r;
-    int         i;
-
-    for (i=0; i<n; i++)
-    {
-        r = inptrA[i] + inptrB[i];
-        if (r != outptr[i])
-        {
-            log_error("INT_ADD int4 test failed\n");
-            return -1;
-        }
-    }
-
-    log_info("INT_ADD int4 test passed\n");
-    return 0;
-}
-
-int
-verify_int_sub4(int *inptrA, int *inptrB, int *outptr, int n)
-{
-    int            r;
-    int         i;
-
-    for (i=0; i<n; i++)
-    {
-        r = inptrA[i] - inptrB[i];
-        if (r != outptr[i])
-        {
-            log_error("INT_SUB int4 test failed\n");
-            return -1;
-        }
-    }
-
-    log_info("INT_SUB int4 test passed\n");
-    return 0;
-}
-
-int
-verify_int_mul4(int *inptrA, int *inptrB, int *outptr, int n)
-{
-    int            r;
-    int         i;
-
-    for (i=0; i<n; i++)
-    {
-        r = inptrA[i] * inptrB[i];
-        if (r != outptr[i])
-        {
-            log_error("INT_MUL int4 test failed\n");
-            return -1;
-        }
-    }
-
-    log_info("INT_MUL int4 test passed\n");
-    return 0;
-}
-
-int
-verify_int_mad4(int *inptrA, int *inptrB, int *inptrC, int *outptr, int n)
-{
-    int            r;
-    int         i;
-
-    for (i=0; i<n; i++)
-    {
-        r = inptrA[i] * inptrB[i] + inptrC[i];
-        if (r != outptr[i])
-        {
-            log_error("INT_MAD int4 test failed\n");
-            return -1;
-        }
-    }
-
-    log_info("INT_MAD int4 test passed\n");
-    return 0;
-}
-
-int
-test_intmath_int4(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
-{
-    cl_mem streams[4];
-    cl_program program[4];
-    cl_kernel kernel[4];
-
-    cl_int *input_ptr[3], *output_ptr, *p;
-    size_t threads[1];
-    int err, i;
-    MTdata d = init_genrand( gRandomSeed );
-
-    size_t length = sizeof(cl_int) * 4 * num_elements;
-
-    input_ptr[0] = (cl_int*)malloc(length);
-    input_ptr[1] = (cl_int*)malloc(length);
-    input_ptr[2] = (cl_int*)malloc(length);
-    output_ptr   = (cl_int*)malloc(length);
-
-    streams[0] = clCreateBuffer(context, CL_MEM_READ_WRITE, length, NULL, NULL);
-    if (!streams[0])
-    {
-        log_error("clCreateBuffer failed\n");
-        return -1;
-    }
-    streams[1] = clCreateBuffer(context, CL_MEM_READ_WRITE, length, NULL, NULL);
-    if (!streams[1])
-    {
-        log_error("clCreateBuffer failed\n");
-        return -1;
-    }
-    streams[2] = clCreateBuffer(context, CL_MEM_READ_WRITE, length, NULL, NULL);
-    if (!streams[2])
-    {
-        log_error("clCreateBuffer failed\n");
-        return -1;
-    }
-    streams[3] = clCreateBuffer(context, CL_MEM_READ_WRITE, length, NULL, NULL);
-    if (!streams[3])
-    {
-        log_error("clCreateBuffer failed\n");
-        return -1;
-    }
-
-    p = input_ptr[0];
-    for (i=0; i<num_elements*4; i++)
-        p[i] = (int)genrand_int32(d);
-    p = input_ptr[1];
-    for (i=0; i<num_elements*4; i++)
-        p[i] = (int)genrand_int32(d);
-    p = input_ptr[2];
-    for (i=0; i<num_elements*4; i++)
-        p[i] = (int)genrand_int32(d);
-
-    free_mtdata(d); d = NULL;
-
-    err = clEnqueueWriteBuffer(queue, streams[0], CL_TRUE, 0, length, input_ptr[0], 0, NULL, NULL);
-    if (err != CL_SUCCESS)
-    {
-        log_error("clEnqueueWriteBuffer failed\n");
-        return -1;
-    }
-    err = clEnqueueWriteBuffer(queue, streams[1], CL_TRUE, 0, length, input_ptr[1], 0, NULL, NULL);
-    if (err != CL_SUCCESS)
-    {
-        log_error("clEnqueueWriteBuffer failed\n");
-        return -1;
-    }
-    err = clEnqueueWriteBuffer(queue, streams[2], CL_TRUE, 0, length, input_ptr[2], 0, NULL, NULL);
-    if (err != CL_SUCCESS)
-    {
-        log_error("clEnqueueWriteBuffer failed\n");
-        return -1;
-    }
-
-    err = create_single_kernel_helper(context, &program[0], &kernel[0], 1, &int_add4_kernel_code, "test_int_add4");
-    if (err != CL_SUCCESS)
-    {
-        log_error("create_single_kernel_helper failed\n");
-        return -1;
-    }
-
-    err = create_single_kernel_helper(context, &program[1], &kernel[1], 1, &int_sub4_kernel_code, "test_int_sub4");
-    if (err != CL_SUCCESS)
-    {
-        log_error("create_single_kernel_helper failed\n");
-        return -1;
-    }
-
-    err = create_single_kernel_helper(context, &program[2], &kernel[2], 1, &int_mul4_kernel_code, "test_int_mul4");
-    if (err != CL_SUCCESS)
-    {
-        log_error("create_single_kernel_helper failed\n");
-        return -1;
-    }
-
-    err = create_single_kernel_helper(context, &program[3], &kernel[3], 1, &int_mad4_kernel_code, "test_int_mad4");
-    if (err != CL_SUCCESS)
-    {
-        log_error("create_single_kernel_helper failed\n");
-        return -1;
-    }
-
-  err  = clSetKernelArg(kernel[0], 0, sizeof streams[0], &streams[0]);
-  err |= clSetKernelArg(kernel[0], 1, sizeof streams[1], &streams[1]);
-  err |= clSetKernelArg(kernel[0], 2, sizeof streams[3], &streams[3]);
-    if (err != CL_SUCCESS)
-    {
-        log_error("clSetKernelArgs failed\n");
-        return -1;
-    }
-
-  err  = clSetKernelArg(kernel[1], 0, sizeof streams[0], &streams[0]);
-  err |= clSetKernelArg(kernel[1], 1, sizeof streams[1], &streams[1]);
-  err |= clSetKernelArg(kernel[1], 2, sizeof streams[3], &streams[3]);
-    if (err != CL_SUCCESS)
-    {
-        log_error("clSetKernelArgs failed\n");
-        return -1;
-    }
-
-  err  = clSetKernelArg(kernel[2], 0, sizeof streams[0], &streams[0]);
-  err |= clSetKernelArg(kernel[2], 1, sizeof streams[1], &streams[1]);
-  err |= clSetKernelArg(kernel[2], 2, sizeof streams[3], &streams[3]);
-    if (err != CL_SUCCESS)
-    {
-        log_error("clSetKernelArgs failed\n");
-        return -1;
-    }
-
-  err  = clSetKernelArg(kernel[3], 0, sizeof streams[0], &streams[0]);
-  err |= clSetKernelArg(kernel[3], 1, sizeof streams[1], &streams[1]);
-  err |= clSetKernelArg(kernel[3], 2, sizeof streams[2], &streams[2]);
-  err |= clSetKernelArg(kernel[3], 3, sizeof streams[3], &streams[3]);
-    if (err != CL_SUCCESS)
-    {
-        log_error("clSetKernelArgs failed\n");
-        return -1;
-    }
-
-  threads[0] = (unsigned int)num_elements;
-  for (i=0; i<4; i++)
-  {
-    err = clEnqueueNDRangeKernel(queue, kernel[i], 1, NULL, threads, NULL, 0, NULL, NULL);
-    if (err != CL_SUCCESS)
-    {
-      log_error("clEnqueueNDRangeKernel failed\n");
-      return -1;
-    }
-
-    err = clEnqueueReadBuffer(queue, streams[3], CL_TRUE, 0, length, output_ptr, 0, NULL, NULL);
-    if (err != CL_SUCCESS)
-    {
-      log_error("clEnqueueReadBuffer failed\n");
-      return -1;
-    }
-
-    switch (i)
-    {
-      case 0:
-        err = verify_int_add4(input_ptr[0], input_ptr[1], output_ptr, num_elements);
-        break;
-      case 1:
-        err = verify_int_sub4(input_ptr[0], input_ptr[1], output_ptr, num_elements);
-        break;
-      case 2:
-        err = verify_int_mul4(input_ptr[0], input_ptr[1], output_ptr, num_elements);
-        break;
-      case 3:
-        err = verify_int_mad4(input_ptr[0], input_ptr[1], input_ptr[2], output_ptr, num_elements);
-        break;
-    }
-    if (err)
-      break;
-    }
-
-    // cleanup
-    clReleaseMemObject(streams[0]);
-    clReleaseMemObject(streams[1]);
-    clReleaseMemObject(streams[2]);
-    clReleaseMemObject(streams[3]);
-    for (i=0; i<4; i++)
-    {
-        clReleaseKernel(kernel[i]);
-        clReleaseProgram(program[i]);
-    }
-    free(input_ptr[0]);
-    free(input_ptr[1]);
-    free(input_ptr[2]);
-    free(output_ptr);
-    return err;
-}
-
-
diff --git a/test_conformance/basic/test_intmath_long.cpp b/test_conformance/basic/test_intmath_long.cpp
deleted file mode 100644
index f163d3b7..00000000
--- a/test_conformance/basic/test_intmath_long.cpp
+++ /dev/null
@@ -1,342 +0,0 @@
-//
-// Copyright (c) 2017 The Khronos Group Inc.
-// 
-// Licensed under the Apache License, Version 2.0 (the "License");
-// you may not use this file except in compliance with the License.
-// You may obtain a copy of the License at
-//
-//    http://www.apache.org/licenses/LICENSE-2.0
-//
-// Unless required by applicable law or agreed to in writing, software
-// distributed under the License is distributed on an "AS IS" BASIS,
-// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
-// See the License for the specific language governing permissions and
-// limitations under the License.
-//
-#include "harness/compat.h"
-
-#include <stdio.h>
-#include <stdlib.h>
-#include <string.h>
-#include <sys/types.h>
-#include <sys/stat.h>
-
-
-#include "procs.h"
-
-const char *long_add_kernel_code =
-"__kernel void test_long_add(__global long *srcA, __global long *srcB, __global long *dst)\n"
-"{\n"
-"    int  tid = get_global_id(0);\n"
-"\n"
-"    dst[tid] = srcA[tid] + srcB[tid];\n"
-"}\n";
-
-const char *long_sub_kernel_code =
-"__kernel void test_long_sub(__global long *srcA, __global long *srcB, __global long *dst)\n"
-"{\n"
-"    int  tid = get_global_id(0);\n"
-"\n"
-"    dst[tid] = srcA[tid] - srcB[tid];\n"
-"}\n";
-
-const char *long_mul_kernel_code =
-"__kernel void test_long_mul(__global long *srcA, __global long *srcB, __global long *dst)\n"
-"{\n"
-"    int  tid = get_global_id(0);\n"
-"\n"
-"    dst[tid] = srcA[tid] * srcB[tid];\n"
-"}\n";
-
-const char *long_mad_kernel_code =
-"__kernel void test_long_mad(__global long *srcA, __global long *srcB, __global long *srcC, __global long *dst)\n"
-"{\n"
-"    int  tid = get_global_id(0);\n"
-"\n"
-"    dst[tid] = srcA[tid] * srcB[tid] + srcC[tid];\n"
-"}\n";
-
-static const float    MAX_ERR = 1e-5f;
-
-int
-verify_long_add(cl_long *inptrA, cl_long *inptrB, cl_long *outptr, int n)
-{
-    cl_long    r;
-    int            i;
-
-    for (i=0; i<n; i++)
-    {
-        r = inptrA[i] + inptrB[i];
-        if (r != outptr[i])
-        {
-            log_error("LONG_ADD int test failed\n");
-            return -1;
-        }
-    }
-
-    log_info("LONG_ADD int test passed\n");
-    return 0;
-}
-
-int
-verify_long_sub(cl_long *inptrA, cl_long *inptrB, cl_long *outptr, int n)
-{
-    cl_long    r;
-    int            i;
-
-    for (i=0; i<n; i++)
-    {
-        r = inptrA[i] - inptrB[i];
-        if (r != outptr[i])
-        {
-            log_error("LONG_SUB int test failed\n");
-            return -1;
-        }
-    }
-
-    log_info("LONG_SUB int test passed\n");
-    return 0;
-}
-
-int
-verify_long_mul(cl_long *inptrA, cl_long *inptrB, cl_long *outptr, int n)
-{
-    cl_long    r;
-    int            i;
-
-    for (i=0; i<n; i++)
-    {
-        r = inptrA[i] * inptrB[i];
-        if (r != outptr[i])
-        {
-            log_error("LONG_MUL int test failed\n");
-            return -1;
-        }
-    }
-
-    log_info("LONG_MUL int test passed\n");
-    return 0;
-}
-
-int
-verify_long_mad(cl_long *inptrA, cl_long *inptrB, cl_long *inptrC, cl_long *outptr, int n)
-{
-    cl_long    r;
-    int            i;
-
-    for (i=0; i<n; i++)
-    {
-        r = inptrA[i] * inptrB[i] + inptrC[i];
-        if (r != outptr[i])
-        {
-            log_error("LONG_MAD int test failed\n");
-            return -1;
-        }
-    }
-
-    log_info("LONG_MAD int test passed\n");
-    return 0;
-}
-
-int
-test_intmath_long(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
-{
-    cl_mem streams[4];
-    cl_program program[4];
-    cl_kernel kernel[4];
-
-    cl_long *input_ptr[3], *output_ptr, *p;
-    size_t threads[1];
-    int err, i;
-
-    if(! gHasLong )
-    {
-        log_info("64-bit integers are not supported by this device. Skipping test.\n");
-        return CL_SUCCESS;
-    }
-
-    MTdata d = init_genrand( gRandomSeed );
-
-    size_t length = sizeof(cl_long) * num_elements;
-
-    input_ptr[0] = (cl_long*)malloc(length);
-    input_ptr[1] = (cl_long*)malloc(length);
-    input_ptr[2] = (cl_long*)malloc(length);
-    output_ptr   = (cl_long*)malloc(length);
-
-    streams[0] = clCreateBuffer(context, CL_MEM_READ_WRITE, length, NULL, NULL);
-    if (!streams[0])
-    {
-        log_error("clCreateBuffer failed\n");
-        return -1;
-    }
-    streams[1] = clCreateBuffer(context, CL_MEM_READ_WRITE, length, NULL, NULL);
-    if (!streams[1])
-    {
-        log_error("clCreateBuffer failed\n");
-        return -1;
-    }
-    streams[2] = clCreateBuffer(context, CL_MEM_READ_WRITE, length, NULL, NULL);
-    if (!streams[2])
-    {
-        log_error("clCreateBuffer failed\n");
-        return -1;
-    }
-    streams[3] = clCreateBuffer(context, CL_MEM_READ_WRITE, length, NULL, NULL);
-    if (!streams[3])
-    {
-        log_error("clCreateBuffer failed\n");
-        return -1;
-    }
-
-    p = input_ptr[0];
-    for (i=0; i<num_elements; i++)
-        p[i] = (cl_long)genrand_int32(d) | ((cl_long) genrand_int32(d) << 32);
-    p = input_ptr[1];
-    for (i=0; i<num_elements; i++)
-        p[i] = (cl_long)genrand_int32(d) | ((cl_long) genrand_int32(d) << 32);
-    p = input_ptr[2];
-    for (i=0; i<num_elements; i++)
-        p[i] = (cl_long)genrand_int32(d) | ((cl_long) genrand_int32(d) << 32);
-
-    free_mtdata(d);
-    d = NULL;
-
-    err = clEnqueueWriteBuffer(queue, streams[0], CL_TRUE, 0, length, input_ptr[0], 0, NULL, NULL);
-    if (err != CL_SUCCESS)
-    {
-        log_error("clEnqueueWriteBuffer failed\n");
-        return -1;
-    }
-    err = clEnqueueWriteBuffer(queue, streams[1], CL_TRUE, 0, length, input_ptr[1], 0, NULL, NULL);
-    if (err != CL_SUCCESS)
-    {
-        log_error("clEnqueueWriteBuffer failed\n");
-        return -1;
-    }
-    err = clEnqueueWriteBuffer(queue, streams[2], CL_TRUE, 0, length, input_ptr[2], 0, NULL, NULL);
-    if (err != CL_SUCCESS)
-    {
-        log_error("clEnqueueWriteBuffer failed\n");
-        return -1;
-    }
-
-    err = create_single_kernel_helper(context, &program[0], &kernel[0], 1, &long_add_kernel_code, "test_long_add");
-    if (err != CL_SUCCESS)
-    {
-        log_error("create_single_kernel_helper failed\n");
-        return -1;
-    }
-
-    err = create_single_kernel_helper(context, &program[1], &kernel[1], 1, &long_sub_kernel_code, "test_long_sub");
-    if (err != CL_SUCCESS)
-    {
-        log_error("create_single_kernel_helper failed\n");
-        return -1;
-    }
-
-    err = create_single_kernel_helper(context, &program[2], &kernel[2], 1, &long_mul_kernel_code, "test_long_mul");
-    if (err != CL_SUCCESS)
-    {
-        log_error("create_single_kernel_helper failed\n");
-        return -1;
-    }
-
-    err = create_single_kernel_helper(context, &program[3], &kernel[3], 1, &long_mad_kernel_code, "test_long_mad");
-    if (err != CL_SUCCESS)
-    {
-        log_error("create_single_kernel_helper failed\n");
-        return -1;
-    }
-
-  err  = clSetKernelArg(kernel[0], 0, sizeof streams[0], &streams[0]);
-  err |= clSetKernelArg(kernel[0], 1, sizeof streams[1], &streams[1]);
-  err |= clSetKernelArg(kernel[0], 2, sizeof streams[3], &streams[3]);
-    if (err != CL_SUCCESS)
-    {
-        log_error("clSetKernelArgs failed\n");
-        return -1;
-    }
-
-  err  = clSetKernelArg(kernel[1], 0, sizeof streams[0], &streams[0]);
-  err |= clSetKernelArg(kernel[1], 1, sizeof streams[1], &streams[1]);
-  err |= clSetKernelArg(kernel[1], 2, sizeof streams[3], &streams[3]);
-    if (err != CL_SUCCESS)
-    {
-        log_error("clSetKernelArgs failed\n");
-        return -1;
-    }
-
-  err  = clSetKernelArg(kernel[2], 0, sizeof streams[0], &streams[0]);
-  err |= clSetKernelArg(kernel[2], 1, sizeof streams[1], &streams[1]);
-  err |= clSetKernelArg(kernel[2], 2, sizeof streams[3], &streams[3]);
-    if (err != CL_SUCCESS)
-    {
-        log_error("clSetKernelArgs failed\n");
-        return -1;
-    }
-
-  err  = clSetKernelArg(kernel[3], 0, sizeof streams[0], &streams[0]);
-  err |= clSetKernelArg(kernel[3], 1, sizeof streams[1], &streams[1]);
-  err |= clSetKernelArg(kernel[3], 2, sizeof streams[2], &streams[2]);
-  err |= clSetKernelArg(kernel[3], 3, sizeof streams[3], &streams[3]);
-    if (err != CL_SUCCESS)
-    {
-        log_error("clSetKernelArgs failed\n");
-        return -1;
-    }
-
-  threads[0] = (unsigned int)num_elements;
-  for (i=0; i<4; i++)
-  {
-    err = clEnqueueNDRangeKernel(queue, kernel[i], 1, NULL, threads, NULL, 0, NULL, NULL);
-    if (err != CL_SUCCESS)
-    {
-      log_error("clEnqueueNDRangeKernel failed\n");
-      return -1;
-    }
-
-    err = clEnqueueReadBuffer(queue, streams[3], CL_TRUE, 0, length, output_ptr, 0, NULL, NULL);
-    if (err != CL_SUCCESS)
-    {
-      log_error("clEnqueueReadBuffer failed\n");
-      return -1;
-    }
-
-    switch (i)
-    {
-      case 0:
-        err = verify_long_add(input_ptr[0], input_ptr[1], output_ptr, num_elements);
-        break;
-      case 1:
-        err = verify_long_sub(input_ptr[0], input_ptr[1], output_ptr, num_elements);
-        break;
-      case 2:
-        err = verify_long_mul(input_ptr[0], input_ptr[1], output_ptr, num_elements);
-        break;
-      case 3:
-        err = verify_long_mad(input_ptr[0], input_ptr[1], input_ptr[2], output_ptr, num_elements);
-        break;
-    }
-    if (err)
-      break;
-    }
-
-    // cleanup
-    clReleaseMemObject(streams[0]);
-    clReleaseMemObject(streams[1]);
-    clReleaseMemObject(streams[2]);
-    clReleaseMemObject(streams[3]);
-    for (i=0; i<4; i++)
-    {
-        clReleaseKernel(kernel[i]);
-        clReleaseProgram(program[i]);
-    }
-    free(input_ptr[0]);
-    free(input_ptr[1]);
-    free(input_ptr[2]);
-    free(output_ptr);
-    return err;
-}
-
-
diff --git a/test_conformance/basic/test_intmath_long2.cpp b/test_conformance/basic/test_intmath_long2.cpp
deleted file mode 100644
index 3736450e..00000000
--- a/test_conformance/basic/test_intmath_long2.cpp
+++ /dev/null
@@ -1,340 +0,0 @@
-//
-// Copyright (c) 2017 The Khronos Group Inc.
-// 
-// Licensed under the Apache License, Version 2.0 (the "License");
-// you may not use this file except in compliance with the License.
-// You may obtain a copy of the License at
-//
-//    http://www.apache.org/licenses/LICENSE-2.0
-//
-// Unless required by applicable law or agreed to in writing, software
-// distributed under the License is distributed on an "AS IS" BASIS,
-// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
-// See the License for the specific language governing permissions and
-// limitations under the License.
-//
-#include "harness/compat.h"
-
-#include <stdio.h>
-#include <stdlib.h>
-#include <string.h>
-#include <sys/types.h>
-#include <sys/stat.h>
-
-
-#include "procs.h"
-
-const char *long_add2_kernel_code =
-"__kernel void test_long_add2(__global long2 *srcA, __global long2 *srcB, __global long2 *dst)\n"
-"{\n"
-"    int  tid = get_global_id(0);\n"
-"\n"
-"    dst[tid] = srcA[tid] + srcB[tid];\n"
-"}\n";
-
-const char *long_sub2_kernel_code =
-"__kernel void test_long_sub2(__global long2 *srcA, __global long2 *srcB, __global long2 *dst)\n"
-"{\n"
-"    int  tid = get_global_id(0);\n"
-"\n"
-"    dst[tid] = srcA[tid] - srcB[tid];\n"
-"}\n";
-
-const char *long_mul2_kernel_code =
-"__kernel void test_long_mul2(__global long2 *srcA, __global long2 *srcB, __global long2 *dst)\n"
-"{\n"
-"    int  tid = get_global_id(0);\n"
-"\n"
-"    dst[tid] = srcA[tid] * srcB[tid];\n"
-"}\n";
-
-const char *long_mad2_kernel_code =
-"__kernel void test_long_mad2(__global long2 *srcA, __global long2 *srcB, __global long2 *srcC, __global long2 *dst)\n"
-"{\n"
-"    int  tid = get_global_id(0);\n"
-"\n"
-"    dst[tid] = srcA[tid] * srcB[tid] + srcC[tid];\n"
-"}\n";
-
-int
-verify_long_add2(cl_long *inptrA, cl_long *inptrB, cl_long *outptr, int n)
-{
-    cl_long    r;
-    int         i;
-
-    for (i=0; i<n; i++)
-    {
-        r = inptrA[i] + inptrB[i];
-        if (r != outptr[i])
-        {
-            log_error("LONG_ADD long2 test failed\n");
-            return -1;
-        }
-    }
-
-    log_info("LONG_ADD long2 test passed\n");
-    return 0;
-}
-
-int
-verify_long_sub2(cl_long *inptrA, cl_long *inptrB, cl_long *outptr, int n)
-{
-    cl_long    r;
-    int         i;
-
-    for (i=0; i<n; i++)
-    {
-        r = inptrA[i] - inptrB[i];
-        if (r != outptr[i])
-        {
-            log_error("LONG_SUB long2 test failed\n");
-            return -1;
-        }
-    }
-
-    log_info("LONG_SUB long2 test passed\n");
-    return 0;
-}
-
-int
-verify_long_mul2(cl_long *inptrA, cl_long *inptrB, cl_long *outptr, int n)
-{
-    cl_long    r;
-    int         i;
-
-    for (i=0; i<n; i++)
-    {
-        r = inptrA[i] * inptrB[i];
-        if (r != outptr[i])
-        {
-            log_error("LONG_MUL long2 test failed\n");
-            return -1;
-        }
-    }
-
-    log_info("LONG_MUL long2 test passed\n");
-    return 0;
-}
-
-int
-verify_long_mad2(cl_long *inptrA, cl_long *inptrB, cl_long *inptrC, cl_long *outptr, int n)
-{
-    cl_long    r;
-    int         i;
-
-    for (i=0; i<n; i++)
-    {
-        r = inptrA[i] * inptrB[i] + inptrC[i];
-        if (r != outptr[i])
-        {
-            log_error("LONG_MAD long2 test failed\n");
-            return -1;
-        }
-    }
-
-    log_info("LONG_MAD long2 test passed\n");
-    return 0;
-}
-
-int
-test_intmath_long2(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
-{
-    cl_mem streams[4];
-    cl_program program[4];
-    cl_kernel kernel[4];
-
-    cl_long *input_ptr[3], *output_ptr, *p;
-    size_t threads[1];
-    int err, i;
-
-    if(! gHasLong)
-    {
-        log_info("64-bit integers are not supported in this device. Skipping test.\n");
-        return 0;
-    }
-
-    MTdata d = init_genrand( gRandomSeed );
-
-    size_t length = sizeof(cl_long) * 2* num_elements;
-
-    input_ptr[0] = (cl_long*)malloc(length);
-    input_ptr[1] = (cl_long*)malloc(length);
-    input_ptr[2] = (cl_long*)malloc(length);
-    output_ptr   = (cl_long*)malloc(length);
-
-    streams[0] = clCreateBuffer(context, CL_MEM_READ_WRITE, length, NULL, NULL);
-    if (!streams[0])
-    {
-        log_error("clCreateBuffer failed\n");
-        return -1;
-    }
-    streams[1] = clCreateBuffer(context, CL_MEM_READ_WRITE, length, NULL, NULL);
-    if (!streams[1])
-    {
-        log_error("clCreateBuffer failed\n");
-        return -1;
-    }
-    streams[2] = clCreateBuffer(context, CL_MEM_READ_WRITE, length, NULL, NULL);
-    if (!streams[2])
-    {
-        log_error("clCreateBuffer failed\n");
-        return -1;
-    }
-    streams[3] = clCreateBuffer(context, CL_MEM_READ_WRITE, length, NULL, NULL);
-    if (!streams[3])
-    {
-        log_error("clCreateBuffer failed\n");
-        return -1;
-    }
-
-    p = input_ptr[0];
-    for (i=0; i<num_elements * 2; i++)
-        p[i] = (cl_long)genrand_int32(d) | ((cl_long) genrand_int32(d) << 32);
-    p = input_ptr[1];
-    for (i=0; i<num_elements * 2; i++)
-        p[i] = (cl_long)genrand_int32(d) | ((cl_long) genrand_int32(d) << 32);
-    p = input_ptr[2];
-    for (i=0; i<num_elements * 2; i++)
-        p[i] = (cl_long)genrand_int32(d) | ((cl_long) genrand_int32(d) << 32);
-
-    free_mtdata(d);
-    d = NULL;
-
-    err = clEnqueueWriteBuffer(queue, streams[0], CL_TRUE, 0, length, input_ptr[0], 0, NULL, NULL);
-    if (err != CL_SUCCESS)
-    {
-        log_error("clEnqueueWriteBuffer failed\n");
-        return -1;
-    }
-    err = clEnqueueWriteBuffer(queue, streams[1], CL_TRUE, 0, length, input_ptr[1], 0, NULL, NULL);
-    if (err != CL_SUCCESS)
-    {
-        log_error("clEnqueueWriteBuffer failed\n");
-        return -1;
-    }
-    err = clEnqueueWriteBuffer(queue, streams[2], CL_TRUE, 0, length, input_ptr[2], 0, NULL, NULL);
-    if (err != CL_SUCCESS)
-    {
-        log_error("clEnqueueWriteBuffer failed\n");
-        return -1;
-    }
-
-    err = create_single_kernel_helper(context, &program[0], &kernel[0], 1, &long_add2_kernel_code, "test_long_add2");
-    if (err != CL_SUCCESS)
-    {
-        log_error("create_single_kernel_helper failed\n");
-        return -1;
-    }
-
-    err = create_single_kernel_helper(context, &program[1], &kernel[1], 1, &long_sub2_kernel_code, "test_long_sub2");
-    if (err != CL_SUCCESS)
-    {
-        log_error("create_single_kernel_helper failed\n");
-        return -1;
-    }
-
-    err = create_single_kernel_helper(context, &program[2], &kernel[2], 1, &long_mul2_kernel_code, "test_long_mul2");
-    if (err != CL_SUCCESS)
-    {
-        log_error("create_single_kernel_helper failed\n");
-        return -1;
-    }
-
-    err = create_single_kernel_helper(context, &program[3], &kernel[3], 1, &long_mad2_kernel_code, "test_long_mad2");
-    if (err != CL_SUCCESS)
-    {
-        log_error("create_single_kernel_helper failed\n");
-        return -1;
-    }
-
-  err  = clSetKernelArg(kernel[0], 0, sizeof streams[0], &streams[0]);
-  err |= clSetKernelArg(kernel[0], 1, sizeof streams[1], &streams[1]);
-  err |= clSetKernelArg(kernel[0], 2, sizeof streams[3], &streams[3]);
-    if (err != CL_SUCCESS)
-    {
-        log_error("clSetKernelArgs failed\n");
-        return -1;
-    }
-
-  err  = clSetKernelArg(kernel[1], 0, sizeof streams[0], &streams[0]);
-  err |= clSetKernelArg(kernel[1], 1, sizeof streams[1], &streams[1]);
-  err |= clSetKernelArg(kernel[1], 2, sizeof streams[3], &streams[3]);
-    if (err != CL_SUCCESS)
-    {
-        log_error("clSetKernelArgs failed\n");
-        return -1;
-    }
-
-  err  = clSetKernelArg(kernel[2], 0, sizeof streams[0], &streams[0]);
-  err |= clSetKernelArg(kernel[2], 1, sizeof streams[1], &streams[1]);
-  err |= clSetKernelArg(kernel[2], 2, sizeof streams[3], &streams[3]);
-    if (err != CL_SUCCESS)
-    {
-        log_error("clSetKernelArgs failed\n");
-        return -1;
-    }
-
-  err  = clSetKernelArg(kernel[3], 0, sizeof streams[0], &streams[0]);
-  err |= clSetKernelArg(kernel[3], 1, sizeof streams[1], &streams[1]);
-  err |= clSetKernelArg(kernel[3], 2, sizeof streams[2], &streams[2]);
-  err |= clSetKernelArg(kernel[3], 3, sizeof streams[3], &streams[3]);
-    if (err != CL_SUCCESS)
-    {
-        log_error("clSetKernelArgs failed\n");
-        return -1;
-    }
-
-  threads[0] = (unsigned int)num_elements;
-  for (i=0; i<4; i++)
-  {
-    err = clEnqueueNDRangeKernel(queue, kernel[i], 1, NULL, threads, NULL, 0, NULL, NULL);
-    if (err != CL_SUCCESS)
-    {
-      log_error("clEnqueueNDRangeKernel failed\n");
-      return -1;
-    }
-
-    err = clEnqueueReadBuffer(queue, streams[3], CL_TRUE, 0, length, output_ptr, 0, NULL, NULL);
-    if (err != CL_SUCCESS)
-    {
-      log_error("clEnqueueReadBuffer failed\n");
-      return -1;
-    }
-
-    switch (i)
-    {
-      case 0:
-        err = verify_long_add2(input_ptr[0], input_ptr[1], output_ptr, num_elements);
-        break;
-      case 1:
-        err = verify_long_sub2(input_ptr[0], input_ptr[1], output_ptr, num_elements);
-        break;
-      case 2:
-        err = verify_long_mul2(input_ptr[0], input_ptr[1], output_ptr, num_elements);
-        break;
-      case 3:
-        err = verify_long_mad2(input_ptr[0], input_ptr[1], input_ptr[2], output_ptr, num_elements);
-        break;
-    }
-    if (err)
-      break;
-    }
-
-    // cleanup
-    clReleaseMemObject(streams[0]);
-    clReleaseMemObject(streams[1]);
-    clReleaseMemObject(streams[2]);
-    clReleaseMemObject(streams[3]);
-    for (i=0; i<4; i++)
-    {
-        clReleaseKernel(kernel[i]);
-        clReleaseProgram(program[i]);
-    }
-    free(input_ptr[0]);
-    free(input_ptr[1]);
-    free(input_ptr[2]);
-    free(output_ptr);
-    return err;
-}
-
-
diff --git a/test_conformance/basic/test_intmath_long4.cpp b/test_conformance/basic/test_intmath_long4.cpp
deleted file mode 100644
index 1a5faae6..00000000
--- a/test_conformance/basic/test_intmath_long4.cpp
+++ /dev/null
@@ -1,340 +0,0 @@
-//
-// Copyright (c) 2017 The Khronos Group Inc.
-// 
-// Licensed under the Apache License, Version 2.0 (the "License");
-// you may not use this file except in compliance with the License.
-// You may obtain a copy of the License at
-//
-//    http://www.apache.org/licenses/LICENSE-2.0
-//
-// Unless required by applicable law or agreed to in writing, software
-// distributed under the License is distributed on an "AS IS" BASIS,
-// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
-// See the License for the specific language governing permissions and
-// limitations under the License.
-//
-#include "harness/compat.h"
-
-#include <stdio.h>
-#include <stdlib.h>
-#include <string.h>
-#include <sys/types.h>
-#include <sys/stat.h>
-
-
-#include "procs.h"
-
-const char *long_add4_kernel_code =
-"__kernel void test_long_add4(__global long4 *srcA, __global long4 *srcB, __global long4 *dst)\n"
-"{\n"
-"    int  tid = get_global_id(0);\n"
-"\n"
-"    dst[tid] = srcA[tid] + srcB[tid];\n"
-"}\n";
-
-const char *long_sub4_kernel_code =
-"__kernel void test_long_sub4(__global long4 *srcA, __global long4 *srcB, __global long4 *dst)\n"
-"{\n"
-"    int  tid = get_global_id(0);\n"
-"\n"
-"    dst[tid] = srcA[tid] - srcB[tid];\n"
-"}\n";
-
-const char *long_mul4_kernel_code =
-"__kernel void test_long_mul4(__global long4 *srcA, __global long4 *srcB, __global long4 *dst)\n"
-"{\n"
-"    int  tid = get_global_id(0);\n"
-"\n"
-"    dst[tid] = srcA[tid] * srcB[tid];\n"
-"}\n";
-
-const char *long_mad4_kernel_code =
-"__kernel void test_long_mad4(__global long4 *srcA, __global long4 *srcB, __global long4 *srcC, __global long4 *dst)\n"
-"{\n"
-"    int  tid = get_global_id(0);\n"
-"\n"
-"    dst[tid] = srcA[tid] * srcB[tid] + srcC[tid];\n"
-"}\n";
-
-int
-verify_long_add4(cl_long *inptrA, cl_long *inptrB, cl_long *outptr, int n)
-{
-    cl_long    r;
-    int         i;
-
-    for (i=0; i<n; i++)
-    {
-        r = inptrA[i] + inptrB[i];
-        if (r != outptr[i])
-        {
-            log_error("LONG_ADD long4 test failed\n");
-            return -1;
-        }
-    }
-
-    log_info("LONG_ADD long4 test passed\n");
-    return 0;
-}
-
-int
-verify_long_sub4(cl_long *inptrA, cl_long *inptrB, cl_long *outptr, int n)
-{
-    cl_long    r;
-    int         i;
-
-    for (i=0; i<n; i++)
-    {
-        r = inptrA[i] - inptrB[i];
-        if (r != outptr[i])
-        {
-            log_error("LONG_SUB long4 test failed\n");
-            return -1;
-        }
-    }
-
-    log_info("LONG_SUB long4 test passed\n");
-    return 0;
-}
-
-int
-verify_long_mul4(cl_long *inptrA, cl_long *inptrB, cl_long *outptr, int n)
-{
-    cl_long    r;
-    int         i;
-
-    for (i=0; i<n; i++)
-    {
-        r = inptrA[i] * inptrB[i];
-        if (r != outptr[i])
-        {
-            log_error("LONG_MUL long4 test failed\n");
-            return -1;
-        }
-    }
-
-    log_info("LONG_MUL long4 test passed\n");
-    return 0;
-}
-
-int
-verify_long_mad4(cl_long *inptrA, cl_long *inptrB, cl_long *inptrC, cl_long *outptr, int n)
-{
-    cl_long    r;
-    int         i;
-
-    for (i=0; i<n; i++)
-    {
-        r = inptrA[i] * inptrB[i] + inptrC[i];
-        if (r != outptr[i])
-        {
-            log_error("LONG_MAD long4 test failed\n");
-            return -1;
-        }
-    }
-
-    log_info("LONG_MAD long4 test passed\n");
-    return 0;
-}
-
-int
-test_intmath_long4(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
-{
-    cl_mem streams[4];
-    cl_program program[4];
-    cl_kernel kernel[4];
-
-    cl_long *input_ptr[3], *output_ptr, *p;
-    size_t threads[1];
-    int err, i;
-
-    if(! gHasLong )
-    {
-        log_info("64-bit integers are not supported by this device. Skipping test.\n");
-        return 0;
-    }
-
-    MTdata d = init_genrand( gRandomSeed );
-
-    size_t length = sizeof(cl_long) * 4 * num_elements;
-
-    input_ptr[0] = (cl_long*)malloc(length);
-    input_ptr[1] = (cl_long*)malloc(length);
-    input_ptr[2] = (cl_long*)malloc(length);
-    output_ptr   = (cl_long*)malloc(length);
-
-    streams[0] = clCreateBuffer(context, CL_MEM_READ_WRITE, length, NULL, NULL);
-    if (!streams[0])
-    {
-        log_error("clCreateBuffer failed\n");
-        return -1;
-    }
-    streams[1] = clCreateBuffer(context, CL_MEM_READ_WRITE, length, NULL, NULL);
-    if (!streams[1])
-    {
-        log_error("clCreateBuffer failed\n");
-        return -1;
-    }
-    streams[2] = clCreateBuffer(context, CL_MEM_READ_WRITE, length, NULL, NULL);
-    if (!streams[2])
-    {
-        log_error("clCreateBuffer failed\n");
-        return -1;
-    }
-    streams[3] = clCreateBuffer(context, CL_MEM_READ_WRITE, length, NULL, NULL);
-    if (!streams[3])
-    {
-        log_error("clCreateBuffer failed\n");
-        return -1;
-    }
-
-    p = input_ptr[0];
-    for (i=0; i<num_elements * 4; i++)
-        p[i] = (cl_long)genrand_int32(d) | ((cl_long) genrand_int32(d) << 32);
-    p = input_ptr[1];
-    for (i=0; i<num_elements * 4; i++)
-        p[i] = (cl_long)genrand_int32(d) | ((cl_long) genrand_int32(d) << 32);
-    p = input_ptr[2];
-    for (i=0; i<num_elements * 4; i++)
-        p[i] = (cl_long)genrand_int32(d) | ((cl_long) genrand_int32(d) << 32);
-
-    free_mtdata(d);
-    d = NULL;
-
-    err = clEnqueueWriteBuffer(queue, streams[0], CL_TRUE, 0, length, input_ptr[0], 0, NULL, NULL);
-    if (err != CL_SUCCESS)
-    {
-        log_error("clEnqueueWriteBuffer failed\n");
-        return -1;
-    }
-    err = clEnqueueWriteBuffer(queue, streams[1], CL_TRUE, 0, length, input_ptr[1], 0, NULL, NULL);
-    if (err != CL_SUCCESS)
-    {
-        log_error("clEnqueueWriteBuffer failed\n");
-        return -1;
-    }
-    err = clEnqueueWriteBuffer(queue, streams[2], CL_TRUE, 0, length, input_ptr[2], 0, NULL, NULL);
-    if (err != CL_SUCCESS)
-    {
-        log_error("clEnqueueWriteBuffer failed\n");
-        return -1;
-    }
-
-    err = create_single_kernel_helper(context, &program[0], &kernel[0], 1, &long_add4_kernel_code, "test_long_add4");
-    if (err != CL_SUCCESS)
-    {
-        log_error("create_single_kernel_helper failed\n");
-        return -1;
-    }
-
-    err = create_single_kernel_helper(context, &program[1], &kernel[1], 1, &long_sub4_kernel_code, "test_long_sub4");
-    if (err != CL_SUCCESS)
-    {
-        log_error("create_single_kernel_helper failed\n");
-        return -1;
-    }
-
-    err = create_single_kernel_helper(context, &program[2], &kernel[2], 1, &long_mul4_kernel_code, "test_long_mul4");
-    if (err != CL_SUCCESS)
-    {
-        log_error("create_single_kernel_helper failed\n");
-        return -1;
-    }
-
-    err = create_single_kernel_helper(context, &program[3], &kernel[3], 1, &long_mad4_kernel_code, "test_long_mad4");
-    if (err != CL_SUCCESS)
-    {
-        log_error("create_single_kernel_helper failed\n");
-        return -1;
-    }
-
-  err  = clSetKernelArg(kernel[0], 0, sizeof streams[0], &streams[0]);
-  err |= clSetKernelArg(kernel[0], 1, sizeof streams[1], &streams[1]);
-  err |= clSetKernelArg(kernel[0], 2, sizeof streams[3], &streams[3]);
-    if (err != CL_SUCCESS)
-    {
-        log_error("clSetKernelArgs failed\n");
-        return -1;
-    }
-
-  err  = clSetKernelArg(kernel[1], 0, sizeof streams[0], &streams[0]);
-  err |= clSetKernelArg(kernel[1], 1, sizeof streams[1], &streams[1]);
-  err |= clSetKernelArg(kernel[1], 2, sizeof streams[3], &streams[3]);
-    if (err != CL_SUCCESS)
-    {
-        log_error("clSetKernelArgs failed\n");
-        return -1;
-    }
-
-  err  = clSetKernelArg(kernel[2], 0, sizeof streams[0], &streams[0]);
-  err |= clSetKernelArg(kernel[2], 1, sizeof streams[1], &streams[1]);
-  err |= clSetKernelArg(kernel[2], 2, sizeof streams[3], &streams[3]);
-    if (err != CL_SUCCESS)
-    {
-        log_error("clSetKernelArgs failed\n");
-        return -1;
-    }
-
-  err  = clSetKernelArg(kernel[3], 0, sizeof streams[0], &streams[0]);
-  err |= clSetKernelArg(kernel[3], 1, sizeof streams[1], &streams[1]);
-  err |= clSetKernelArg(kernel[3], 2, sizeof streams[2], &streams[2]);
-  err |= clSetKernelArg(kernel[3], 3, sizeof streams[3], &streams[3]);
-    if (err != CL_SUCCESS)
-    {
-        log_error("clSetKernelArgs failed\n");
-        return -1;
-    }
-
-  threads[0] = (unsigned int)num_elements;
-  for (i=0; i<4; i++)
-  {
-    err = clEnqueueNDRangeKernel(queue, kernel[i], 1, NULL, threads, NULL, 0, NULL, NULL);
-    if (err != CL_SUCCESS)
-    {
-      log_error("clEnqueueNDRangeKernel failed\n");
-      return -1;
-    }
-
-    err = clEnqueueReadBuffer(queue, streams[3], CL_TRUE, 0, length, output_ptr, 0, NULL, NULL);
-    if (err != CL_SUCCESS)
-    {
-      log_error("clEnqueueReadBuffer failed\n");
-      return -1;
-    }
-
-    switch (i)
-    {
-      case 0:
-        err = verify_long_add4(input_ptr[0], input_ptr[1], output_ptr, num_elements);
-        break;
-      case 1:
-        err = verify_long_sub4(input_ptr[0], input_ptr[1], output_ptr, num_elements);
-        break;
-      case 2:
-        err = verify_long_mul4(input_ptr[0], input_ptr[1], output_ptr, num_elements);
-        break;
-      case 3:
-        err = verify_long_mad4(input_ptr[0], input_ptr[1], input_ptr[2], output_ptr, num_elements);
-        break;
-    }
-    if (err)
-      break;
-    }
-
-    // cleanup
-    clReleaseMemObject(streams[0]);
-    clReleaseMemObject(streams[1]);
-    clReleaseMemObject(streams[2]);
-    clReleaseMemObject(streams[3]);
-    for (i=0; i<4; i++)
-    {
-        clReleaseKernel(kernel[i]);
-        clReleaseProgram(program[i]);
-    }
-    free(input_ptr[0]);
-    free(input_ptr[1]);
-    free(input_ptr[2]);
-    free(output_ptr);
-    return err;
-}
-
-