diff --git a/test_conformance/basic/CMakeLists.txt b/test_conformance/basic/CMakeLists.txt
index fe070a6b..df778a68 100644
--- a/test_conformance/basic/CMakeLists.txt
+++ b/test_conformance/basic/CMakeLists.txt
@@ -2,7 +2,7 @@ set(MODULE_NAME BASIC)
 
 set(${MODULE_NAME}_SOURCES
     main.cpp
-    test_fpmath_float.cpp test_fpmath_float2.cpp test_fpmath_float4.cpp
+    test_fpmath_float.cpp
     test_intmath.cpp
     test_hiloeo.cpp test_local.cpp test_pointercast.cpp
     test_if.cpp test_loop.cpp
diff --git a/test_conformance/basic/test_fpmath_float.cpp b/test_conformance/basic/test_fpmath_float.cpp
index 60d509b0..fced0f4e 100644
--- a/test_conformance/basic/test_fpmath_float.cpp
+++ b/test_conformance/basic/test_fpmath_float.cpp
@@ -1,6 +1,6 @@
 //
 // Copyright (c) 2017 The Khronos Group Inc.
-// 
+//
 // Licensed under the Apache License, Version 2.0 (the "License");
 // you may not use this file except in compliance with the License.
 // You may obtain a copy of the License at
@@ -22,245 +22,175 @@
 #include <sys/stat.h>
 #include "harness/rounding_mode.h"
 
+#include <algorithm>
+#include <functional>
+#include <string>
+#include <vector>
+
 #include "procs.h"
 
-static const char *fpadd_kernel_code =
-"__kernel void test_fpadd(__global float *srcA, __global float *srcB, __global float *dst)\n"
-"{\n"
-"    int  tid = get_global_id(0);\n"
-"\n"
-"    dst[tid] = srcA[tid] + srcB[tid];\n"
-"}\n";
-
-static const char *fpsub_kernel_code =
-"__kernel void test_fpsub(__global float *srcA, __global float *srcB, __global float *dst)\n"
-"{\n"
-"    int  tid = get_global_id(0);\n"
-"\n"
-"    dst[tid] = srcA[tid] - srcB[tid];\n"
-"}\n";
-
-static const char *fpmul_kernel_code =
-"__kernel void test_fpmul(__global float *srcA, __global float *srcB, __global float *dst)\n"
-"{\n"
-"    int  tid = get_global_id(0);\n"
-"\n"
-"    dst[tid] = srcA[tid] * srcB[tid];\n"
-"}\n";
-
-
-static int
-verify_fpadd(float *inptrA, float *inptrB, float *outptr, int n)
+struct TestDef
 {
-    float       r;
-    int         i;
+    const char op;
+    std::function<float(float, float)> ref;
+};
 
-    for (i=0; i<n; i++)
+static const char *fp_kernel_code = R"(
+__kernel void test_fp(__global TYPE *srcA, __global TYPE *srcB, __global TYPE *dst)
+{
+    int  tid = get_global_id(0);
+
+    dst[tid] = srcA[tid] OP srcB[tid];
+})";
+
+static int verify_fp(std::vector<float> (&input)[2], std::vector<float> &output,
+                     const TestDef &test)
+{
+
+    auto &inA = input[0];
+    auto &inB = input[1];
+    for (int i = 0; i < output.size(); i++)
     {
-        r = inptrA[i] + inptrB[i];
-        if (r != outptr[i])
+        float r = test.ref(inA[i], inB[i]);
+        if (r != output[i])
         {
-            log_error("FP_ADD float test failed\n");
+            log_error("FP '%c' float test failed\n", test.op);
             return -1;
         }
     }
 
-    log_info("FP_ADD float test passed\n");
+    log_info("FP '%c' float test passed\n", test.op);
     return 0;
 }
 
-static int
-verify_fpsub(float *inptrA, float *inptrB, float *outptr, int n)
-{
-    float       r;
-    int         i;
 
-    for (i=0; i<n; i++)
+void generate_random_inputs(std::vector<cl_float> (&input)[2])
+{
+    RandomSeed seed(gRandomSeed);
+
+    auto random_generator = [&seed]() {
+        return get_random_float(-MAKE_HEX_FLOAT(0x1.0p31f, 0x1, 31),
+                                MAKE_HEX_FLOAT(0x1.0p31f, 0x1, 31), seed);
+    };
+
+    for (auto &v : input)
     {
-        r = inptrA[i] - inptrB[i];
-        if (r != outptr[i])
-        {
-            log_error("FP_SUB float test failed\n");
-            return -1;
-        }
+        std::generate(v.begin(), v.end(), random_generator);
     }
-
-    log_info("FP_SUB float test passed\n");
-    return 0;
 }
 
-static int
-verify_fpmul(float *inptrA, float *inptrB, float *outptr, int n)
+template <size_t N>
+int test_fpmath(cl_device_id device, cl_context context, cl_command_queue queue,
+                int num_elements, const std::string type_str,
+                const TestDef &test)
 {
-    float       r;
-    int         i;
+    clMemWrapper streams[3];
+    clProgramWrapper program;
+    clKernelWrapper kernel;
 
-    for (i=0; i<n; i++)
-    {
-        r = inptrA[i] * inptrB[i];
-        if (r != outptr[i])
-        {
-            log_error("FP_MUL float test failed\n");
-            return -1;
-        }
-    }
+    int err;
 
-    log_info("FP_MUL float test passed\n");
-    return 0;
-}
+    size_t length = sizeof(cl_float) * num_elements * N;
 
-
-int
-test_fpmath_float(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
-{
-    cl_mem streams[4];
-    cl_program program[3];
-    cl_kernel kernel[3];
-
-    float *input_ptr[3], *output_ptr, *p;
-    size_t threads[1];
-    int err, i;
-    MTdata d = init_genrand( gRandomSeed );
-    size_t length = sizeof(cl_float) * num_elements;
     int isRTZ = 0;
     RoundingMode oldMode = kDefaultRoundingMode;
 
-    // check for floating point capabilities
-    cl_device_fp_config single_config = 0;
-    err = clGetDeviceInfo( device, CL_DEVICE_SINGLE_FP_CONFIG, sizeof( single_config ), &single_config, NULL );
-    if (err) {
-      log_error("clGetDeviceInfo for CL_DEVICE_SINGLE_FP_CONFIG failed: %d", err);
-      return -1;
-    }
-    //If we only support rtz mode
-    if( CL_FP_ROUND_TO_ZERO == ( single_config & (CL_FP_ROUND_TO_ZERO|CL_FP_ROUND_TO_NEAREST) ) )
+    // If we only support rtz mode
+    if (CL_FP_ROUND_TO_ZERO == get_default_rounding_mode(device))
     {
-        //Check to make sure we are an embedded device
-        char profile[32];
-        err = clGetDeviceInfo( device, CL_DEVICE_PROFILE, sizeof(profile), profile, NULL);
-        if( err )
-        {
-            log_error("clGetDeviceInfo for CL_DEVICE_PROFILE failed: %d", err);
-              return -1;
-        }
-        if( 0 != strcmp( profile, "EMBEDDED_PROFILE"))
-        {
-            log_error( "FAILURE:  Device doesn't support CL_FP_ROUND_TO_NEAREST and isn't EMBEDDED_PROFILE\n" );
-            return -1;
-        }
-
         isRTZ = 1;
         oldMode = get_round();
     }
 
 
-    input_ptr[0] = (cl_float*)malloc(length);
-    input_ptr[1] = (cl_float*)malloc(length);
-    input_ptr[2] = (cl_float*)malloc(length);
-    output_ptr   = (cl_float*)malloc(length);
+    std::vector<cl_float> inputs[]{ std::vector<cl_float>(N * num_elements),
+                                    std::vector<cl_float>(N * num_elements) };
+    std::vector<cl_float> output = std::vector<cl_float>(N * num_elements);
 
-    streams[0] = clCreateBuffer(context, CL_MEM_READ_WRITE, length, NULL, &err);
-    test_error( err, "clCreateBuffer failed.");
-    streams[1] = clCreateBuffer(context, CL_MEM_READ_WRITE, length, NULL, &err);
-    test_error( err, "clCreateBuffer failed.");
-    streams[2] = clCreateBuffer(context, CL_MEM_READ_WRITE, length, NULL, &err);
-    test_error( err, "clCreateBuffer failed.");
-    streams[3] = clCreateBuffer(context, CL_MEM_READ_WRITE, length, NULL, &err);
-    test_error( err, "clCreateBuffer failed.");
+    generate_random_inputs(inputs);
 
-    p = input_ptr[0];
-    for (i=0; i<num_elements; i++)
-        p[i] = get_random_float(-MAKE_HEX_FLOAT(0x1.0p31f, 0x1, 31), MAKE_HEX_FLOAT(0x1.0p31f, 0x1, 31), d);
-    p = input_ptr[1];
-    for (i=0; i<num_elements; i++)
-        p[i] = get_random_float(-MAKE_HEX_FLOAT(0x1.0p31f, 0x1, 31), MAKE_HEX_FLOAT(0x1.0p31f, 0x1, 31), d);
-    p = input_ptr[2];
-    for (i=0; i<num_elements; i++)
-        p[i] = get_random_float(-MAKE_HEX_FLOAT(0x1.0p31f, 0x1, 31), MAKE_HEX_FLOAT(0x1.0p31f, 0x1, 31), d);
-
-    err = clEnqueueWriteBuffer(queue, streams[0], CL_TRUE, 0, length, input_ptr[0], 0, NULL, NULL);
-    test_error( err, "clEnqueueWriteBuffer failed.");
-
-    err = clEnqueueWriteBuffer(queue, streams[1], CL_TRUE, 0, length, input_ptr[1], 0, NULL, NULL);
-    test_error( err, "clEnqueueWriteBuffer failed.");
-
-    err = clEnqueueWriteBuffer(queue, streams[2], CL_TRUE, 0, length, input_ptr[2], 0, NULL, NULL);
-    test_error( err, "clEnqueueWriteBuffer failed.");
-
-    err = create_single_kernel_helper(context, &program[0], &kernel[0], 1, &fpadd_kernel_code, "test_fpadd");
-    test_error( err, "create_single_kernel_helper failed");
-
-    err = create_single_kernel_helper(context, &program[1], &kernel[1], 1, &fpsub_kernel_code, "test_fpsub");
-    test_error( err, "create_single_kernel_helper failed");
-
-    err = create_single_kernel_helper(context, &program[2], &kernel[2], 1, &fpmul_kernel_code, "test_fpmul");
-    test_error( err, "create_single_kernel_helper failed");
-
-
-    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]);
-    test_error( err, "clSetKernelArgs failed.");
-
-    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]);
-    test_error( err, "clSetKernelArgs failed.");
-
-    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]);
-    test_error( err, "clSetKernelArgs failed.");
-
-    threads[0] = (unsigned int)num_elements;
-    for (i=0; i<3; i++)
+    for (int i = 0; i < ARRAY_SIZE(streams); i++)
     {
-        err = clEnqueueNDRangeKernel(queue, kernel[i], 1, NULL, threads, NULL, 0, NULL, NULL);
-        test_error( err, "clEnqueueNDRangeKernel failed.");
-
-        err = clEnqueueReadBuffer(queue, streams[3], CL_TRUE, 0, length, output_ptr, 0, NULL, NULL);
-        test_error( err, "clEnqueueReadBuffer failed.");
-
-        if( isRTZ )
-            set_round( kRoundTowardZero, kfloat );
-
-        switch (i)
-        {
-            case 0:
-                err = verify_fpadd(input_ptr[0], input_ptr[1], output_ptr, num_elements);
-                break;
-            case 1:
-                err = verify_fpsub(input_ptr[0], input_ptr[1], output_ptr, num_elements);
-                break;
-            case 2:
-                err = verify_fpmul(input_ptr[0], input_ptr[1], output_ptr, num_elements);
-                break;
-        }
-
-        if( isRTZ )
-            set_round( oldMode, kfloat );
-
-        if (err)
-            break;
+        streams[i] =
+            clCreateBuffer(context, CL_MEM_READ_WRITE, length, NULL, &err);
+        test_error(err, "clCreateBuffer failed.");
+    }
+    for (int i = 0; i < ARRAY_SIZE(inputs); i++)
+    {
+        err = clEnqueueWriteBuffer(queue, streams[i], CL_TRUE, 0, length,
+                                   inputs[i].data(), 0, NULL, NULL);
+        test_error(err, "clEnqueueWriteBuffer failed.");
     }
 
-    // cleanup
-    clReleaseMemObject(streams[0]);
-    clReleaseMemObject(streams[1]);
-    clReleaseMemObject(streams[2]);
-    clReleaseMemObject(streams[3]);
-    for (i=0; i<3; i++)
+    std::string build_options = "-DTYPE=";
+    build_options.append(type_str).append(" -DOP=").append(1, test.op);
+
+    err = create_single_kernel_helper(context, &program, &kernel, 1,
+                                      &fp_kernel_code, "test_fp",
+                                      build_options.c_str());
+
+    test_error(err, "create_single_kernel_helper failed");
+
+    for (int i = 0; i < ARRAY_SIZE(streams); i++)
     {
-        clReleaseKernel(kernel[i]);
-        clReleaseProgram(program[i]);
+        err = clSetKernelArg(kernel, i, sizeof(streams[i]), &streams[i]);
+        test_error(err, "clSetKernelArgs failed.");
     }
-    free(input_ptr[0]);
-    free(input_ptr[1]);
-    free(input_ptr[2]);
-    free(output_ptr);
-    free_mtdata( d );
+
+    size_t threads[] = { static_cast<size_t>(num_elements) };
+    err = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, threads, NULL, 0, NULL,
+                                 NULL);
+    test_error(err, "clEnqueueNDRangeKernel failed.");
+
+    err = clEnqueueReadBuffer(queue, streams[2], CL_TRUE, 0, length,
+                              output.data(), 0, NULL, NULL);
+    test_error(err, "clEnqueueReadBuffer failed.");
+
+    if (isRTZ) set_round(kRoundTowardZero, kfloat);
+
+    err = verify_fp(inputs, output, test);
+
+    if (isRTZ) set_round(oldMode, kfloat);
 
     return err;
 }
 
 
+template <size_t N>
+int test_fpmath_common(cl_device_id device, cl_context context,
+                       cl_command_queue queue, int num_elements,
+                       const std::string type_str)
+{
+    TestDef tests[] = { { '+', std::plus<float>() },
+                        { '-', std::minus<float>() },
+                        { '*', std::multiplies<float>() } };
+    int err = TEST_PASS;
+
+    for (const auto &test : tests)
+    {
+        err |= test_fpmath<N>(device, context, queue, num_elements, type_str,
+                              test);
+    }
+
+    return err;
+}
+
+int test_fpmath_float(cl_device_id device, cl_context context,
+                      cl_command_queue queue, int num_elements)
+{
+    return test_fpmath_common<1>(device, context, queue, num_elements, "float");
+}
+
+int test_fpmath_float2(cl_device_id device, cl_context context,
+                       cl_command_queue queue, int num_elements)
+{
+    return test_fpmath_common<2>(device, context, queue, num_elements,
+                                 "float2");
+}
+
+int test_fpmath_float4(cl_device_id device, cl_context context,
+                       cl_command_queue queue, int num_elements)
+{
+    return test_fpmath_common<4>(device, context, queue, num_elements,
+                                 "float4");
+}
diff --git a/test_conformance/basic/test_fpmath_float2.cpp b/test_conformance/basic/test_fpmath_float2.cpp
deleted file mode 100644
index 1881b4be..00000000
--- a/test_conformance/basic/test_fpmath_float2.cpp
+++ /dev/null
@@ -1,266 +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 "harness/rounding_mode.h"
-
-
-#include "procs.h"
-
-const char *fpadd2_kernel_code =
-"__kernel void test_fpadd2(__global float2 *srcA, __global float2 *srcB, __global float2 *dst)\n"
-"{\n"
-"    int  tid = get_global_id(0);\n"
-"\n"
-"    dst[tid] = srcA[tid] + srcB[tid];\n"
-"}\n";
-
-const char *fpsub2_kernel_code =
-"__kernel void test_fpsub2(__global float2 *srcA, __global float2 *srcB, __global float2 *dst)\n"
-"{\n"
-"    int  tid = get_global_id(0);\n"
-"\n"
-"    dst[tid] = srcA[tid] - srcB[tid];\n"
-"}\n";
-
-const char *fpmul2_kernel_code =
-"__kernel void test_fpmul2(__global float2 *srcA, __global float2 *srcB, __global float2 *dst)\n"
-"{\n"
-"    int  tid = get_global_id(0);\n"
-"\n"
-"    dst[tid] = srcA[tid] * srcB[tid];\n"
-"}\n";
-
-
-int
-verify_fpadd2(float *inptrA, float *inptrB, float *outptr, int n)
-{
-    float       r;
-    int         i;
-
-    for (i=0; i<n; i++)
-    {
-        r = inptrA[i] + inptrB[i];
-        if (r != outptr[i])
-        {
-            log_error("FP_ADD float2 test failed\n");
-            return -1;
-        }
-    }
-
-    log_info("FP_ADD float2 test passed\n");
-    return 0;
-}
-
-int
-verify_fpsub2(float *inptrA, float *inptrB, float *outptr, int n)
-{
-    float       r;
-    int         i;
-
-    for (i=0; i<n; i++)
-    {
-        r = inptrA[i] - inptrB[i];
-        if (r != outptr[i])
-        {
-            log_error("FP_SUB float2 test failed\n");
-            return -1;
-        }
-    }
-
-    log_info("FP_SUB float2 test passed\n");
-    return 0;
-}
-
-int
-verify_fpmul2(float *inptrA, float *inptrB, float *outptr, int n)
-{
-    float       r;
-    int         i;
-
-    for (i=0; i<n; i++)
-    {
-        r = inptrA[i] * inptrB[i];
-        if (r != outptr[i])
-        {
-            log_error("FP_MUL float2 test failed\n");
-            return -1;
-        }
-    }
-
-    log_info("FP_MUL float2 test passed\n");
-    return 0;
-}
-
-
-int
-test_fpmath_float2(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
-{
-    cl_mem streams[4];
-    cl_program program[3];
-    cl_kernel kernel[3];
-
-    cl_float *input_ptr[3], *output_ptr, *p;
-    size_t threads[1];
-    int err, i;
-    MTdata d = init_genrand( gRandomSeed );
-
-    size_t length = sizeof(cl_float) * 2 * num_elements;
-    int isRTZ = 0;
-    RoundingMode oldMode = kDefaultRoundingMode;
-
-    // check for floating point capabilities
-    cl_device_fp_config single_config = 0;
-    err = clGetDeviceInfo( device, CL_DEVICE_SINGLE_FP_CONFIG, sizeof( single_config ), &single_config, NULL );
-    if (err) {
-      log_error("clGetDeviceInfo for CL_DEVICE_SINGLE_FP_CONFIG failed: %d", err);
-      return -1;
-    }
-    //If we only support rtz mode
-    if( CL_FP_ROUND_TO_ZERO == ( single_config & (CL_FP_ROUND_TO_ZERO|CL_FP_ROUND_TO_NEAREST) ) )
-    {
-        //Check to make sure we are an embedded device
-        char profile[32];
-        err = clGetDeviceInfo( device, CL_DEVICE_PROFILE, sizeof(profile), profile, NULL);
-        if( err )
-        {
-            log_error("clGetDeviceInfo for CL_DEVICE_PROFILE failed: %d", err);
-              return -1;
-        }
-        if( 0 != strcmp( profile, "EMBEDDED_PROFILE"))
-        {
-            log_error( "FAILURE:  Device doesn't support CL_FP_ROUND_TO_NEAREST and isn't EMBEDDED_PROFILE\n" );
-            return -1;
-        }
-
-        isRTZ = 1;
-        oldMode = get_round();
-    }
-
-    input_ptr[0] = (cl_float*)malloc(length);
-    input_ptr[1] = (cl_float*)malloc(length);
-    input_ptr[2] = (cl_float*)malloc(length);
-    output_ptr   = (cl_float*)malloc(length);
-
-    streams[0] = clCreateBuffer(context, CL_MEM_READ_WRITE, length, NULL, &err);
-    test_error( err, "clCreateBuffer failed.");
-    streams[1] = clCreateBuffer(context, CL_MEM_READ_WRITE, length, NULL, &err);
-    test_error( err, "clCreateBuffer failed.");
-    streams[2] = clCreateBuffer(context, CL_MEM_READ_WRITE, length, NULL, &err);
-    test_error( err, "clCreateBuffer failed.");
-    streams[3] = clCreateBuffer(context, CL_MEM_READ_WRITE, length, NULL, &err);
-    test_error( err, "clCreateBuffer failed.");
-
-    p = input_ptr[0];
-    for (i=0; i<num_elements*2; i++)
-        p[i] = get_random_float(-MAKE_HEX_FLOAT( 0x1.0p31f, 0x1, 31), MAKE_HEX_FLOAT( 0x1.0p31f, 0x1, 31), d);
-    p = input_ptr[1];
-    for (i=0; i<num_elements*2; i++)
-        p[i] = get_random_float(-MAKE_HEX_FLOAT( 0x1.0p31f, 0x1, 31), MAKE_HEX_FLOAT( 0x1.0p31f, 0x1, 31), d);
-    p = input_ptr[2];
-    for (i=0; i<num_elements*2; i++)
-        p[i] = get_random_float(-MAKE_HEX_FLOAT( 0x1.0p31f, 0x1, 31), MAKE_HEX_FLOAT( 0x1.0p31f, 0x1, 31), d);
-
-    err = clEnqueueWriteBuffer(queue, streams[0], CL_TRUE, 0, length, input_ptr[0], 0, NULL, NULL);
-    test_error(err, "clEnqueueWriteBuffer failed");
-    err = clEnqueueWriteBuffer(queue, streams[1], CL_TRUE, 0, length, input_ptr[1], 0, NULL, NULL);
-    test_error(err, "clEnqueueWriteBuffer failed");
-    err = clEnqueueWriteBuffer(queue, streams[2], CL_TRUE, 0, length, input_ptr[2], 0, NULL, NULL);
-    test_error(err, "clEnqueueWriteBuffer failed");
-
-    err = create_single_kernel_helper(context, &program[0], &kernel[0], 1, &fpadd2_kernel_code, "test_fpadd2");
-    test_error( err, "create_single_kernel_helper failed");
-
-    err = create_single_kernel_helper(context, &program[1], &kernel[1], 1, &fpsub2_kernel_code, "test_fpsub2");
-    test_error( err, "create_single_kernel_helper failed");
-
-    err = create_single_kernel_helper(context, &program[2], &kernel[2], 1, &fpmul2_kernel_code, "test_fpmul2");
-    test_error( err, "create_single_kernel_helper failed");
-
-
-    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]);
-    test_error( err, "clSetKernelArgs failed.");
-
-    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]);
-    test_error( err, "clSetKernelArgs failed.");
-
-    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]);
-    test_error( err, "clSetKernelArgs failed.");
-    free_mtdata(d);
-    d = NULL;
-
-    threads[0] = (unsigned int)num_elements;
-    for (i=0; i<3; i++)
-    {
-        err = clEnqueueNDRangeKernel(queue, kernel[i], 1, NULL, threads, NULL, 0, NULL, NULL);
-      test_error( err, "clEnqueueNDRangeKernel failed.");
-
-        err = clEnqueueReadBuffer(queue, streams[3], CL_TRUE, 0, length, output_ptr, 0, NULL, NULL);
-      test_error( err, "clEnqueueReadBuffer failed.");
-
-        if( isRTZ )
-            set_round( kRoundTowardZero, kfloat );
-
-        switch (i)
-        {
-            case 0:
-                err = verify_fpadd2(input_ptr[0], input_ptr[1], output_ptr, num_elements*2);
-                break;
-            case 1:
-                err = verify_fpsub2(input_ptr[0], input_ptr[1], output_ptr, num_elements*2);
-                break;
-            case 2:
-                err = verify_fpmul2(input_ptr[0], input_ptr[1], output_ptr, num_elements*2);
-                break;
-        }
-
-        if( isRTZ )
-            set_round( oldMode, kfloat );
-
-        if (err)
-            break;
-    }
-
-
-    // cleanup
-    clReleaseMemObject(streams[0]);
-    clReleaseMemObject(streams[1]);
-    clReleaseMemObject(streams[2]);
-    clReleaseMemObject(streams[3]);
-    for (i=0; i<3; 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_fpmath_float4.cpp b/test_conformance/basic/test_fpmath_float4.cpp
deleted file mode 100644
index 999c8ec5..00000000
--- a/test_conformance/basic/test_fpmath_float4.cpp
+++ /dev/null
@@ -1,267 +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"
-#include "harness/rounding_mode.h"
-
-const char *fpadd4_kernel_code =
-"__kernel void test_fpadd4(__global float4 *srcA, __global float4 *srcB, __global float4 *dst)\n"
-"{\n"
-"    int  tid = get_global_id(0);\n"
-"\n"
-"    dst[tid] = srcA[tid] + srcB[tid];\n"
-"}\n";
-
-const char *fpsub4_kernel_code =
-"__kernel void test_fpsub4(__global float4 *srcA, __global float4 *srcB, __global float4 *dst)\n"
-"{\n"
-"    int  tid = get_global_id(0);\n"
-"\n"
-"    dst[tid] = srcA[tid] - srcB[tid];\n"
-"}\n";
-
-const char *fpmul4_kernel_code =
-"__kernel void test_fpmul4(__global float4 *srcA, __global float4 *srcB, __global float4 *dst)\n"
-"{\n"
-"    int  tid = get_global_id(0);\n"
-"\n"
-"    dst[tid] = srcA[tid] * srcB[tid];\n"
-"}\n";
-
-
-int
-verify_fpadd4(float *inptrA, float *inptrB, float *outptr, int n)
-{
-    float       r;
-    int         i;
-
-    for (i=0; i<n; i++)
-    {
-        r = inptrA[i] + inptrB[i];
-        if (r != outptr[i])
-        {
-            log_error("FP_ADD float4 test failed\n");
-            return -1;
-        }
-    }
-
-    log_info("FP_ADD float4 test passed\n");
-    return 0;
-}
-
-int
-verify_fpsub4(float *inptrA, float *inptrB, float *outptr, int n)
-{
-    float       r;
-    int         i;
-
-    for (i=0; i<n; i++)
-    {
-        r = inptrA[i] - inptrB[i];
-        if (r != outptr[i])
-        {
-            log_error("FP_SUB float4 test failed\n");
-            return -1;
-        }
-    }
-
-    log_info("FP_SUB float4 test passed\n");
-    return 0;
-}
-
-int
-verify_fpmul4(float *inptrA, float *inptrB, float *outptr, int n)
-{
-  float       r;
-  int         i;
-
-  for (i=0; i<n; i++)
-  {
-    r = inptrA[i] * inptrB[i];
-    if (r != outptr[i])
-    {
-      log_error("FP_MUL float4 test failed\n");
-      return -1;
-    }
-  }
-
-  log_info("FP_MUL float4 test passed\n");
-  return 0;
-}
-
-
-int
-test_fpmath_float4(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
-{
-    cl_mem streams[4];
-    cl_program program[3];
-    cl_kernel kernel[3];
-
-    cl_float *input_ptr[3], *output_ptr, *p;
-    size_t threads[1];
-    int err, i;
-    MTdata d = init_genrand( gRandomSeed );
-
-    size_t length = sizeof(cl_float) * 4 * num_elements;
-    int isRTZ = 0;
-    RoundingMode oldMode = kDefaultRoundingMode;
-
-    // check for floating point capabilities
-    cl_device_fp_config single_config = 0;
-    err = clGetDeviceInfo( device, CL_DEVICE_SINGLE_FP_CONFIG, sizeof( single_config ), &single_config, NULL );
-    if (err) {
-      log_error("clGetDeviceInfo for CL_DEVICE_SINGLE_FP_CONFIG failed: %d", err);
-      return -1;
-    }
-    //If we only support rtz mode
-    if( CL_FP_ROUND_TO_ZERO == ( single_config & (CL_FP_ROUND_TO_ZERO|CL_FP_ROUND_TO_NEAREST) ) )
-    {
-        //Check to make sure we are an embedded device
-        char profile[32];
-        err = clGetDeviceInfo( device, CL_DEVICE_PROFILE, sizeof(profile), profile, NULL);
-        if( err )
-        {
-            log_error("clGetDeviceInfo for CL_DEVICE_PROFILE failed: %d", err);
-              return -1;
-        }
-        if( 0 != strcmp( profile, "EMBEDDED_PROFILE"))
-        {
-            log_error( "FAILURE:  Device doesn't support CL_FP_ROUND_TO_NEAREST and isn't EMBEDDED_PROFILE\n" );
-            return -1;
-        }
-
-        isRTZ = 1;
-        oldMode = get_round();
-    }
-
-    input_ptr[0] = (cl_float*)malloc(length);
-    input_ptr[1] = (cl_float*)malloc(length);
-    input_ptr[2] = (cl_float*)malloc(length);
-    output_ptr   = (cl_float*)malloc(length);
-
-    streams[0] = clCreateBuffer(context, CL_MEM_READ_WRITE, length, NULL, &err);
-    test_error( err, "clCreateBuffer failed.");
-    streams[1] = clCreateBuffer(context, CL_MEM_READ_WRITE, length, NULL, &err);
-    test_error( err, "clCreateBuffer failed.");
-    streams[2] = clCreateBuffer(context, CL_MEM_READ_WRITE, length, NULL, &err);
-    test_error( err, "clCreateBuffer failed.");
-    streams[3] = clCreateBuffer(context, CL_MEM_READ_WRITE, length, NULL, &err);
-    test_error( err, "clCreateBuffer failed.");
-
-    p = input_ptr[0];
-    for (i=0; i<num_elements*4; i++)
-        p[i] = get_random_float(-MAKE_HEX_FLOAT(0x1.0p31f, 0x1, 31), MAKE_HEX_FLOAT(0x1.0p31f, 0x1, 31), d);
-    p = input_ptr[1];
-    for (i=0; i<num_elements*4; i++)
-        p[i] = get_random_float(-MAKE_HEX_FLOAT(0x1.0p31f, 0x1, 31), MAKE_HEX_FLOAT(0x1.0p31f, 0x1, 31), d);
-    p = input_ptr[2];
-    for (i=0; i<num_elements*4; i++)
-        p[i] = get_random_float(-MAKE_HEX_FLOAT(0x1.0p31f, 0x1, 31), MAKE_HEX_FLOAT(0x1.0p31f, 0x1, 31), d);
-
-    free_mtdata(d);
-
-    err = clEnqueueWriteBuffer(queue, streams[0], CL_TRUE, 0, length, input_ptr[0], 0, NULL, NULL);
-    test_error(err, "clEnqueueWriteBuffer failed");
-    err = clEnqueueWriteBuffer(queue, streams[1], CL_TRUE, 0, length, input_ptr[1], 0, NULL, NULL);
-    test_error(err, "clEnqueueWriteBuffer failed");
-    err = clEnqueueWriteBuffer(queue, streams[2], CL_TRUE, 0, length, input_ptr[2], 0, NULL, NULL);
-    test_error(err, "clEnqueueWriteBuffer failed");
-
-    err = create_single_kernel_helper(context, &program[0], &kernel[0], 1, &fpadd4_kernel_code, "test_fpadd4");
-    test_error( err, "create_single_kernel_helper failed");
-
-    err = create_single_kernel_helper(context, &program[1], &kernel[1], 1, &fpsub4_kernel_code, "test_fpsub4");
-    test_error( err, "create_single_kernel_helper failed");
-
-    err = create_single_kernel_helper(context, &program[2], &kernel[2], 1, &fpmul4_kernel_code, "test_fpmul4");
-    test_error( err, "create_single_kernel_helper failed");
-
-
-    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]);
-    test_error( err, "clSetKernelArgs failed.");
-
-    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]);
-    test_error( err, "clSetKernelArgs failed.");
-
-    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]);
-    test_error( err, "clSetKernelArgs failed.");
-
-
-  threads[0] = (unsigned int)num_elements;
-  for (i=0; i<3; i++)
-  {
-    err = clEnqueueNDRangeKernel(queue, kernel[i], 1, NULL, threads, NULL, 0, NULL, NULL);
-    test_error( err, "clEnqueueNDRangeKernel failed.");
-
-    err = clEnqueueReadBuffer(queue, streams[3], CL_TRUE, 0, length, output_ptr, 0, NULL, NULL);
-    test_error( err, "clEnqueueReadBuffer failed.");
-
-    if( isRTZ )
-        set_round( kRoundTowardZero, kfloat );
-
-    switch (i)
-    {
-      case 0:
-        err = verify_fpadd4(input_ptr[0], input_ptr[1], output_ptr, num_elements*4);
-        break;
-      case 1:
-        err = verify_fpsub4(input_ptr[0], input_ptr[1], output_ptr, num_elements*4);
-        break;
-      case 2:
-        err = verify_fpmul4(input_ptr[0], input_ptr[1], output_ptr, num_elements*4);
-        break;
-    }
-
-    if( isRTZ )
-        set_round( oldMode, kfloat );
-
-    if (err)
-      break;
-    }
-
-
-    // cleanup
-    clReleaseMemObject(streams[0]);
-    clReleaseMemObject(streams[1]);
-    clReleaseMemObject(streams[2]);
-    clReleaseMemObject(streams[3]);
-    for (i=0; i<3; i++)
-    {
-        clReleaseKernel(kernel[i]);
-        clReleaseProgram(program[i]);
-    }
-    free(input_ptr[0]);
-    free(input_ptr[1]);
-    free(input_ptr[2]);
-    free(output_ptr);
-    return err;
-}
-
-