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