diff --git a/test_conformance/basic/test_if.cpp b/test_conformance/basic/test_if.cpp index c92ec3221..f2a8fa829 100644 --- a/test_conformance/basic/test_if.cpp +++ b/test_conformance/basic/test_if.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 @@ -21,146 +21,119 @@ #include #include +#include +#include #include "procs.h" -const char *conditional_kernel_code = -"__kernel void test_if(__global int *src, __global int *dst)\n" -"{\n" -" int tid = get_global_id(0);\n" -"\n" -" if (src[tid] == 0)\n" -" dst[tid] = 0x12345678;\n" -" else if (src[tid] == 1)\n" -" dst[tid] = 0x23456781;\n" -" else if (src[tid] == 2)\n" -" dst[tid] = 0x34567812;\n" -" else if (src[tid] == 3)\n" -" dst[tid] = 0x45678123;\n" -" else if (src[tid] == 4)\n" -" dst[tid] = 0x56781234;\n" -" else if (src[tid] == 5)\n" -" dst[tid] = 0x67812345;\n" -" else if (src[tid] == 6)\n" -" dst[tid] = 0x78123456;\n" -" else if (src[tid] == 7)\n" -" dst[tid] = 0x81234567;\n" -" else\n" -" dst[tid] = 0x7FFFFFFF;\n" -"\n" -"}\n"; - -const int results[] = { - 0x12345678, - 0x23456781, - 0x34567812, - 0x45678123, - 0x56781234, - 0x67812345, - 0x78123456, - 0x81234567, -}; - -int -verify_if(int *inptr, int *outptr, int n) +namespace { +const char *conditional_kernel_code = R"( +__kernel void test_if(__global int *src, __global int *dst) { - int r, i; + int tid = get_global_id(0); + + if (src[tid] == 0) + dst[tid] = 0x12345678; + else if (src[tid] == 1) + dst[tid] = 0x23456781; + else if (src[tid] == 2) + dst[tid] = 0x34567812; + else if (src[tid] == 3) + dst[tid] = 0x45678123; + else if (src[tid] == 4) + dst[tid] = 0x56781234; + else if (src[tid] == 5) + dst[tid] = 0x67812345; + else if (src[tid] == 6) + dst[tid] = 0x78123456; + else if (src[tid] == 7) + dst[tid] = 0x81234567; + else + dst[tid] = 0x7FFFFFFF; +} +)"; - for (i=0; i input, std::vector output) +{ + const cl_int results[] = { + 0x12345678, 0x23456781, 0x34567812, 0x45678123, + 0x56781234, 0x67812345, 0x78123456, 0x81234567, + }; + + auto predicate = [&results](cl_int a, cl_int b) { + if (a <= 7) + return b == results[a]; else - r = 0x7FFFFFFF; + return b == 0x7FFFFFFF; + }; - if (r != outptr[i]) - { - log_error("IF test failed\n"); - return -1; - } + if (!std::equal(input.begin(), input.end(), output.begin(), predicate)) + { + log_error("IF test failed\n"); + return -1; } log_info("IF test passed\n"); return 0; } -int test_if(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements) +void generate_random_inputs(std::vector &v) { - cl_mem streams[2]; - cl_int *input_ptr, *output_ptr; - cl_program program; - cl_kernel kernel; - size_t threads[1]; - int err, i; - MTdata d = init_genrand( gRandomSeed ); + RandomSeed seed(gRandomSeed); + + auto random_generator = [&seed]() { + return static_cast(get_random_float(0, 32, seed)); + }; + + std::generate(v.begin(), v.end(), random_generator); +} +} +int test_if(cl_device_id device, cl_context context, cl_command_queue queue, + int num_elements) +{ + clMemWrapper streams[2]; + clProgramWrapper program; + clKernelWrapper kernel; + + int err; size_t length = sizeof(cl_int) * num_elements; - input_ptr = (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; - } + std::vector input(num_elements); + std::vector output(num_elements); - for (i=0; i