opencl/source/opencl_device.cxx

   1 /* -*- Mode: C++; tab-width: 4; indent-tabs-mode: nil; c-basic-offset: 4 -*- */
   2 /*
   3  * This file is part of the LibreOffice project.
   4  *
   5  * This Source Code Form is subject to the terms of the Mozilla Public
   6  * License, v. 2.0. If a copy of the MPL was not distributed with this
   7  * file, You can obtain one at http://mozilla.org/MPL/2.0/.
   8  */
   9
  10 #include <float.h>
  11 #include <iostream>
  12 #include <memory>
  13 #include <string_view>
  14 #include <vector>
  15 #include <algorithm>
  16
  17 #include <comphelper/random.hxx>
  18 #include <o3tl/safeint.hxx>
  19 #include <opencl/openclconfig.hxx>
  20 #include <opencl/platforminfo.hxx>
  21 #include <sal/log.hxx>
  22 #include <rtl/math.hxx>
  23 #include <tools/time.hxx>
  24
  25 #include <opencl/OpenCLZone.hxx>
  26
  27 #include <opencl_device.hxx>
  28 #include <opencl_device_selection.h>
  29
  30 #define INPUTSIZE  15360
  31 #define OUTPUTSIZE 15360
  32
  33 #define STRINGIFY(...) #__VA_ARGS__"\n"
  34
  35 namespace {
  36
  37 void DS_CHECK_STATUS(cl_int status, char const * name) {
  38     if (CL_SUCCESS != status)
  39     {
  40     SAL_INFO("opencl.device", "Error code is " << status << " at " << name);
  41     }
  42 }
  43
  44 bool bIsDeviceSelected = false;
  45 ds_device selectedDevice;
  46
  47 struct LibreOfficeDeviceEvaluationIO
  48 {
  49     std::vector<double> input0;
  50     std::vector<double> input1;
  51     std::vector<double> input2;
  52     std::vector<double> input3;
  53     std::vector<double> output;
  54     tools::ULong inputSize;
  55     tools::ULong outputSize;
  56 };
  57
  58 const char* source = STRINGIFY(
  59 \n#if defined(KHR_DP_EXTENSION)
  60 \n#pragma OPENCL EXTENSION cl_khr_fp64 : enable
  61 \n#elif defined(AMD_DP_EXTENSION)
  62 \n#pragma OPENCL EXTENSION cl_amd_fp64 : enable
  63 \n#endif
  64     \n
  65     int isNan(fp_t a) { return a != a; }
  66     fp_t fsum(fp_t a, fp_t b) { return a + b; }
  67
  68     fp_t fAverage(__global fp_t* input)
  69 {
  70     fp_t sum = 0;
  71     int count = 0;
  72     for (int i = 0; i < INPUTSIZE; i++)
  73     {
  74         if (!isNan(input[i]))
  75         {
  76             sum = fsum(input[i], sum);
  77             count += 1;
  78         }
  79     }
  80     return sum / (fp_t)count;
  81 }
  82     fp_t fMin(__global fp_t* input)
  83 {
  84     fp_t min = MAXFLOAT;
  85     for (int i = 0; i < INPUTSIZE; i++)
  86     {
  87         if (!isNan(input[i]))
  88         {
  89             min = fmin(input[i], min);
  90         }
  91     }
  92     return min;
  93 }
  94     fp_t fSoP(__global fp_t* input0, __global fp_t* input1)
  95 {
  96     fp_t sop = 0.0;
  97     for (int i = 0; i < INPUTSIZE; i++)
  98     {
  99         sop += (isNan(input0[i]) ? 0 : input0[i]) * (isNan(input1[i]) ? 0 : input1[i]);
 100     }
 101     return sop;
 102 }
 103     __kernel void DynamicKernel(
 104         __global fp_t* result, __global fp_t* input0, __global fp_t* input1, __global fp_t* input2, __global fp_t* input3)
 105 {
 106     int gid0 = get_global_id(0);
 107     fp_t tmp0 = fAverage(input0);
 108     fp_t tmp1 = fMin(input1) * fSoP(input2, input3);
 109     result[gid0] = fsum(tmp0, tmp1);
 110 }
 111     );
 112
 113 size_t sourceSize[] = { strlen(source) };
 114
 115 /* Random number generator */
 116 double random(double min, double max)
 117 {
 118     if (rtl::math::approxEqual(min, max))
 119         return min;
 120     return comphelper::rng::uniform_real_distribution(min, max);
 121 }
 122
 123 /* Populate input */
 124 void populateInput(std::unique_ptr<LibreOfficeDeviceEvaluationIO> const & testData)
 125 {
 126     double* input0 = testData->input0.data();
 127     double* input1 = testData->input1.data();
 128     double* input2 = testData->input2.data();
 129     double* input3 = testData->input3.data();
 130     for (tools::ULong i = 0; i < testData->inputSize; i++)
 131     {
 132         input0[i] = random(0, i);
 133         input1[i] = random(0, i);
 134         input2[i] = random(0, i);
 135         input3[i] = random(0, i);
 136     }
 137 }
 138
 139 /* Evaluate devices */
 140 ds_status evaluateScoreForDevice(ds_device& rDevice, std::unique_ptr<LibreOfficeDeviceEvaluationIO> const & testData)
 141 {
 142     if (rDevice.eType == DeviceType::OpenCLDevice)
 143     {
 144         /* Evaluating an OpenCL device */
 145         SAL_INFO("opencl.device", "Device: \"" << rDevice.sDeviceName << "\" (OpenCL) evaluation...");
 146         cl_int clStatus;
 147
 148         /* Check for 64-bit float extensions */
 149         std::unique_ptr<char[]> aExtInfo;
 150         {
 151             size_t aDevExtInfoSize = 0;
 152
 153             OpenCLZone zone;
 154             clStatus = clGetDeviceInfo(rDevice.aDeviceID, CL_DEVICE_EXTENSIONS, 0, nullptr, &aDevExtInfoSize);
 155             DS_CHECK_STATUS(clStatus, "evaluateScoreForDevice::clGetDeviceInfo");
 156
 157             aExtInfo.reset(new char[aDevExtInfoSize]);
 158             clStatus = clGetDeviceInfo(rDevice.aDeviceID, CL_DEVICE_EXTENSIONS, sizeof(char) * aDevExtInfoSize, aExtInfo.get(), nullptr);
 159             DS_CHECK_STATUS(clStatus, "evaluateScoreForDevice::clGetDeviceInfo");
 160         }
 161
 162         bool bKhrFp64Flag = false;
 163         bool bAmdFp64Flag = false;
 164         const char* buildOption = nullptr;
 165         std::string tmpStr("-Dfp_t=double -Dfp_t4=double4 -Dfp_t16=double16 -DINPUTSIZE=");
 166         std::ostringstream tmpOStrStr;
 167         tmpOStrStr << std::dec << INPUTSIZE;
 168         tmpStr.append(tmpOStrStr.str());
 169
 170         if ((std::string(aExtInfo.get())).find("cl_khr_fp64") != std::string::npos)
 171         {
 172             bKhrFp64Flag = true;
 173             //buildOption = "-D KHR_DP_EXTENSION -Dfp_t=double -Dfp_t4=double4 -Dfp_t16=double16";
 174             tmpStr.append(" -DKHR_DP_EXTENSION");
 175             buildOption = tmpStr.c_str();
 176             SAL_INFO("opencl.device", "... has cl_khr_fp64");
 177         }
 178         else if ((std::string(aExtInfo.get())).find("cl_amd_fp64") != std::string::npos)
 179         {
 180             bAmdFp64Flag = true;
 181             //buildOption = "-D AMD_DP_EXTENSION -Dfp_t=double -Dfp_t4=double4 -Dfp_t16=double16";
 182             tmpStr.append(" -DAMD_DP_EXTENSION");
 183             buildOption = tmpStr.c_str();
 184             SAL_INFO("opencl.device", "... has cl_amd_fp64");
 185         }
 186
 187         if (!bKhrFp64Flag && !bAmdFp64Flag)
 188         {
 189             /* No 64-bit float support */
 190             rDevice.fTime = DBL_MAX;
 191             rDevice.bErrors = false;
 192             SAL_INFO("opencl.device", "... no fp64 support");
 193         }
 194         else
 195         {
 196             /* 64-bit float support present */
 197
 198             OpenCLZone zone;
 199
 200             /* Create context and command queue */
 201             cl_context  clContext = clCreateContext(nullptr, 1, &rDevice.aDeviceID, nullptr, nullptr, &clStatus);
 202             DS_CHECK_STATUS(clStatus, "evaluateScoreForDevice::clCreateContext");
 203             cl_command_queue clQueue = clCreateCommandQueue(clContext, rDevice.aDeviceID, 0, &clStatus);
 204             DS_CHECK_STATUS(clStatus, "evaluateScoreForDevice::clCreateCommandQueue");
 205
 206             /* Build program */
 207             cl_program clProgram = clCreateProgramWithSource(clContext, 1, &source, sourceSize, &clStatus);
 208             DS_CHECK_STATUS(clStatus, "evaluateScoreForDevice::clCreateProgramWithSource");
 209             clStatus = clBuildProgram(clProgram, 1, &rDevice.aDeviceID, buildOption, nullptr, nullptr);
 210             DS_CHECK_STATUS(clStatus, "evaluateScoreForDevice::clBuildProgram");
 211             if (CL_SUCCESS != clStatus)
 212             {
 213                 /* Build program failed */
 214                 size_t length;
 215                 char* buildLog;
 216                 clStatus = clGetProgramBuildInfo(clProgram, rDevice.aDeviceID, CL_PROGRAM_BUILD_LOG, 0, nullptr, &length);
 217                 buildLog = static_cast<char*>(malloc(length));
 218                 clGetProgramBuildInfo(clProgram, rDevice.aDeviceID, CL_PROGRAM_BUILD_LOG, length, buildLog, &length);
 219                 SAL_INFO("opencl.device", "Build Errors:\n" << buildLog);
 220                 free(buildLog);
 221
 222                 rDevice.fTime = DBL_MAX;
 223                 rDevice.bErrors = true;
 224             }
 225             else
 226             {
 227                 /* Build program succeeded */
 228                 sal_uInt64 kernelTime = tools::Time::GetMonotonicTicks();
 229
 230                 /* Run kernel */
 231                 cl_kernel clKernel = clCreateKernel(clProgram, "DynamicKernel", &clStatus);
 232                 DS_CHECK_STATUS(clStatus, "evaluateScoreForDevice::clCreateKernel");
 233                 cl_mem clResult = clCreateBuffer(clContext, CL_MEM_WRITE_ONLY | CL_MEM_USE_HOST_PTR, sizeof(cl_double) * testData->outputSize, testData->output.data(), &clStatus);
 234                 DS_CHECK_STATUS(clStatus, "evaluateScoreForDevice::clCreateBuffer::clResult");
 235                 cl_mem clInput0 = clCreateBuffer(clContext, CL_MEM_READ_ONLY  | CL_MEM_USE_HOST_PTR, sizeof(cl_double) * testData->inputSize,  testData->input0.data(), &clStatus);
 236                 DS_CHECK_STATUS(clStatus, "evaluateScoreForDevice::clCreateBuffer::clInput0");
 237                 cl_mem clInput1 = clCreateBuffer(clContext, CL_MEM_READ_ONLY  | CL_MEM_USE_HOST_PTR, sizeof(cl_double) * testData->inputSize,  testData->input1.data(), &clStatus);
 238                 DS_CHECK_STATUS(clStatus, "evaluateScoreForDevice::clCreateBuffer::clInput1");
 239                 cl_mem clInput2 = clCreateBuffer(clContext, CL_MEM_READ_ONLY  | CL_MEM_USE_HOST_PTR, sizeof(cl_double) * testData->inputSize,  testData->input2.data(), &clStatus);
 240                 DS_CHECK_STATUS(clStatus, "evaluateScoreForDevice::clCreateBuffer::clInput2");
 241                 cl_mem clInput3 = clCreateBuffer(clContext, CL_MEM_READ_ONLY  | CL_MEM_USE_HOST_PTR, sizeof(cl_double) * testData->inputSize,  testData->input3.data(), &clStatus);
 242                 DS_CHECK_STATUS(clStatus, "evaluateScoreForDevice::clCreateBuffer::clInput3");
 243                 clStatus = clSetKernelArg(clKernel, 0, sizeof(cl_mem), static_cast<void*>(&clResult));
 244                 DS_CHECK_STATUS(clStatus, "evaluateScoreForDevice::clSetKernelArg::clResult");
 245                 clStatus = clSetKernelArg(clKernel, 1, sizeof(cl_mem), static_cast<void*>(&clInput0));
 246                 DS_CHECK_STATUS(clStatus, "evaluateScoreForDevice::clSetKernelArg::clInput0");
 247                 clStatus = clSetKernelArg(clKernel, 2, sizeof(cl_mem), static_cast<void*>(&clInput1));
 248                 DS_CHECK_STATUS(clStatus, "evaluateScoreForDevice::clSetKernelArg::clInput1");
 249                 clStatus = clSetKernelArg(clKernel, 3, sizeof(cl_mem), static_cast<void*>(&clInput2));
 250                 DS_CHECK_STATUS(clStatus, "evaluateScoreForDevice::clSetKernelArg::clInput2");
 251                 clStatus = clSetKernelArg(clKernel, 4, sizeof(cl_mem), static_cast<void*>(&clInput3));
 252                 DS_CHECK_STATUS(clStatus, "evaluateScoreForDevice::clSetKernelArg::clInput3");
 253                 size_t globalWS[1] = { testData->outputSize };
 254                 size_t const localSize[1] = { 64 };
 255                 clStatus = clEnqueueNDRangeKernel(clQueue, clKernel, 1, nullptr, globalWS, localSize, 0, nullptr, nullptr);
 256                 DS_CHECK_STATUS(clStatus, "evaluateScoreForDevice::clEnqueueNDRangeKernel");
 257                 clFinish(clQueue);
 258                 clReleaseMemObject(clInput3);
 259                 clReleaseMemObject(clInput2);
 260                 clReleaseMemObject(clInput1);
 261                 clReleaseMemObject(clInput0);
 262                 clReleaseMemObject(clResult);
 263                 clReleaseKernel(clKernel);
 264
 265                 rDevice.fTime = tools::Time::GetMonotonicTicks() - kernelTime;
 266                 rDevice.bErrors = false;
 267             }
 268
 269             clReleaseProgram(clProgram);
 270             clReleaseCommandQueue(clQueue);
 271             clReleaseContext(clContext);
 272         }
 273     }
 274     else
 275     {
 276         /* Evaluating a Native CPU device */
 277         SAL_INFO("opencl.device", "Device: \"CPU\" (Native) evaluation...");
 278         sal_uInt64 kernelTime = tools::Time::GetMonotonicTicks();
 279
 280         tools::ULong j;
 281         for (j = 0; j < testData->outputSize; j++)
 282         {
 283             double fAverage = 0.0f;
 284             double fMin = DBL_MAX;
 285             double fSoP = 0.0f;
 286             for (tools::ULong i = 0; i < testData->inputSize; i++)
 287             {
 288                 fAverage += testData->input0[i];
 289                 fMin = std::min(fMin, testData->input1[i]);
 290                 fSoP += testData->input2[i] * testData->input3[i];
 291             }
 292             fAverage /= testData->inputSize;
 293             testData->output[j] = fAverage + (fMin * fSoP);
 294             // Don't run for much longer than one second
 295             if (j > 0 && j % 100 == 0)
 296             {
 297                 rDevice.fTime = tools::Time::GetMonotonicTicks() - kernelTime;
 298                 if (rDevice.fTime >= 1)
 299                     break;
 300             }
 301         }
 302
 303         rDevice.fTime = tools::Time::GetMonotonicTicks() - kernelTime;
 304
 305         // Scale time to how long it would have taken to go all the way to outputSize
 306         rDevice.fTime /= (static_cast<double>(j) / testData->outputSize);
 307
 308         // InterpretTail - the S/W fallback is nothing like as efficient
 309         // as any good openCL implementation: no SIMD, tons of branching
 310         // in the inner loops etc. Generously characterise it as only 10x
 311         // slower than the above.
 312         rDevice.fTime *= 10.0;
 313         rDevice.bErrors = false;
 314     }
 315     return DS_SUCCESS;
 316 }
 317
 318 ds_status profileDevices(std::unique_ptr<ds_profile> const & pProfile, std::unique_ptr<LibreOfficeDeviceEvaluationIO> const & pTestData)
 319 {
 320     ds_status status = DS_SUCCESS;
 321
 322     if (!pProfile)
 323         return DS_INVALID_PROFILE;
 324
 325     for (ds_device& rDevice : pProfile->devices)
 326     {
 327         ds_status evaluatorStatus = evaluateScoreForDevice(rDevice, pTestData);
 328         if (evaluatorStatus != DS_SUCCESS)
 329         {
 330             status = evaluatorStatus;
 331             return status;
 332         }
 333     }
 334     return status;
 335 }
 336
 337 /* Pick best device */
 338 int pickBestDevice(std::unique_ptr<ds_profile> const & profile)
 339 {
 340     double bestScore = DBL_MAX;
 341
 342     int nBestDeviceIndex = -1;
 343
 344     for (std::vector<ds_device>::size_type d = 0; d < profile->devices.size();
 345          d++)
 346     {
 347         ds_device& device = profile->devices[d];
 348
 349         // Check denylist and allowlist for actual devices
 350         if (device.eType == DeviceType::OpenCLDevice)
 351         {
 352             // There is a silly impedance mismatch here. Why do we
 353             // need two different ways to describe an OpenCL platform
 354             // and an OpenCL device driver?
 355
 356             OpenCLPlatformInfo aPlatform;
 357             OpenCLDeviceInfo aDevice;
 358
 359             // We know that only the below fields are used by checkForKnownBadCompilers()
 360             aPlatform.maVendor = OStringToOUString(device.sPlatformVendor, RTL_TEXTENCODING_UTF8);
 361             aDevice.maName = OStringToOUString(device.sDeviceName, RTL_TEXTENCODING_UTF8);
 362             aDevice.maDriver = OStringToOUString(device.sDriverVersion, RTL_TEXTENCODING_UTF8);
 363
 364             // If denylisted or not allowlisted, ignore it
 365             if (OpenCLConfig::get().checkImplementation(aPlatform, aDevice))
 366             {
 367                 SAL_INFO("opencl.device", "Device[" << d << "] " << device.sDeviceName << " is denylisted or not allowlisted");
 368                 device.fTime = DBL_MAX;
 369                 device.bErrors = false;
 370             }
 371         }
 372
 373         double fScore = DBL_MAX;
 374         if (device.fTime >= 0.0
 375             || rtl::math::approxEqual(device.fTime, DBL_MAX))
 376         {
 377             fScore = device.fTime;
 378         }
 379         else
 380         {
 381             SAL_INFO("opencl.device", "Unusual null score");
 382         }
 383
 384         if (device.eType == DeviceType::OpenCLDevice)
 385         {
 386             SAL_INFO("opencl.device", "Device[" << d << "] " << device.sDeviceName << " (OpenCL) score is " << fScore);
 387         }
 388         else
 389         {
 390             SAL_INFO("opencl.device", "Device[" << d << "] CPU (Native) score is " << fScore);
 391         }
 392         if (fScore < bestScore)
 393         {
 394             bestScore = fScore;
 395             nBestDeviceIndex = d;
 396         }
 397     }
 398     if (nBestDeviceIndex != -1 && profile->devices[nBestDeviceIndex].eType == DeviceType::OpenCLDevice)
 399     {
 400         SAL_INFO("opencl.device", "Selected Device[" << nBestDeviceIndex << "]: " << profile->devices[nBestDeviceIndex].sDeviceName << "(OpenCL).");
 401     }
 402     else
 403     {
 404         SAL_INFO("opencl.device", "Selected Device[" << nBestDeviceIndex << "]: CPU (Native).");
 405     }
 406     return nBestDeviceIndex;
 407 }
 408
 409 /* Return device ID for matching device name */
 410 int matchDevice(std::unique_ptr<ds_profile> const & profile, char* deviceName)
 411 {
 412     int deviceMatch = -1;
 413     for (size_t d = 0; d < profile->devices.size() - 1; d++)
 414     {
 415         if (profile->devices[d].sDeviceName.indexOf(deviceName) != -1)
 416             deviceMatch = d;
 417     }
 418     if (std::string("NATIVE_CPU").find(deviceName) != std::string::npos)
 419         deviceMatch = profile->devices.size() - 1;
 420     return deviceMatch;
 421 }
 422
 423 class LogWriter
 424 {
 425 private:
 426     SvFileStream maStream;
 427 public:
 428     explicit LogWriter(OUString const & aFileName)
 429         : maStream(aFileName, StreamMode::WRITE)
 430     {}
 431
 432     void text(std::string_view rText)
 433     {
 434         maStream.WriteOString(rText);
 435         maStream.WriteChar('\n');
 436     }
 437
 438     void log(std::string_view rKey, std::string_view rValue)
 439     {
 440         maStream.WriteOString(rKey);
 441         maStream.WriteCharPtr(": ");
 442         maStream.WriteOString(rValue);
 443         maStream.WriteChar('\n');
 444     }
 445
 446     void log(std::string_view rKey, int rValue)
 447     {
 448         log(rKey, OString::number(rValue));
 449     }
 450
 451     void log(std::string_view rKey, bool rValue)
 452     {
 453         log(rKey, OString::boolean(rValue));
 454     }
 455 };
 456
 457
 458 void writeDevicesLog(std::unique_ptr<ds_profile> const & rProfile, std::u16string_view sProfilePath, int nSelectedIndex)
 459 {
 460     OUString aCacheFile(OUString::Concat(sProfilePath) + "opencl_devices.log");
 461     LogWriter aWriter(aCacheFile);
 462
 463     int nIndex = 0;
 464
 465     for (const ds_device& rDevice : rProfile->devices)
 466     {
 467         if (rDevice.eType == DeviceType::OpenCLDevice)
 468         {
 469             aWriter.log("Device Index", nIndex);
 470             aWriter.log("  Selected", nIndex == nSelectedIndex);
 471             aWriter.log("  Device Name", rDevice.sDeviceName);
 472             aWriter.log("  Device Vendor", rDevice.sDeviceVendor);
 473             aWriter.log("  Device Version", rDevice.sDeviceVersion);
 474             aWriter.log("  Driver Version", rDevice.sDriverVersion);
 475             aWriter.log("  Device Type", rDevice.sDeviceType);
 476             aWriter.log("  Device Extensions", rDevice.sDeviceExtensions);
 477             aWriter.log("  Device OpenCL C Version", rDevice.sDeviceOpenCLVersion);
 478
 479             aWriter.log("  Device Available", rDevice.bDeviceAvailable);
 480             aWriter.log("  Device Compiler Available", rDevice.bDeviceCompilerAvailable);
 481             aWriter.log("  Device Linker Available", rDevice.bDeviceLinkerAvailable);
 482
 483             aWriter.log("  Platform Name", rDevice.sPlatformName);
 484             aWriter.log("  Platform Vendor", rDevice.sPlatformVendor);
 485             aWriter.log("  Platform Version", rDevice.sPlatformVersion);
 486             aWriter.log("  Platform Profile", rDevice.sPlatformProfile);
 487             aWriter.log("  Platform Extensions", rDevice.sPlatformExtensions);
 488             aWriter.text("");
 489         }
 490         nIndex++;
 491     }
 492 }
 493
 494 } // end anonymous namespace
 495
 496 ds_device const & getDeviceSelection(
 497     std::u16string_view sProfilePath, bool bForceSelection)
 498 {
 499     /* Run only if device is not yet selected */
 500     if (!bIsDeviceSelected || bForceSelection)
 501     {
 502         /* Setup */
 503         std::unique_ptr<ds_profile> aProfile;
 504         ds_status status;
 505         status = initDSProfile(aProfile, "LibreOffice v1");
 506
 507         if (status != DS_SUCCESS)
 508         {
 509             // failed to initialize profile.
 510             selectedDevice.eType = DeviceType::NativeCPU;
 511             return selectedDevice;
 512         }
 513
 514         /* Try reading scores from file */
 515         OUString sFilePath = OUString::Concat(sProfilePath) + "opencl_profile.xml";
 516
 517         if (!bForceSelection)
 518         {
 519             status = readProfile(sFilePath, aProfile);
 520         }
 521         else
 522         {
 523             status = DS_INVALID_PROFILE;
 524             SAL_INFO("opencl.device", "Performing forced profiling.");
 525         }
 526         if (DS_SUCCESS != status)
 527         {
 528             if (!bForceSelection)
 529             {
 530                 SAL_INFO("opencl.device", "Profile file not available (" << sFilePath << "); performing profiling.");
 531             }
 532
 533             /* Populate input data for micro-benchmark */
 534             std::unique_ptr<LibreOfficeDeviceEvaluationIO> testData(new LibreOfficeDeviceEvaluationIO);
 535             testData->inputSize  = INPUTSIZE;
 536             testData->outputSize = OUTPUTSIZE;
 537             testData->input0.resize(testData->inputSize);
 538             testData->input1.resize(testData->inputSize);
 539             testData->input2.resize(testData->inputSize);
 540             testData->input3.resize(testData->inputSize);
 541             testData->output.resize(testData->outputSize);
 542             populateInput(testData);
 543
 544             /* Perform evaluations */
 545             status = profileDevices(aProfile, testData);
 546
 547             if (DS_SUCCESS == status)
 548             {
 549                 /* Write scores to file */
 550                 status = writeProfile(sFilePath, aProfile);
 551                 if (DS_SUCCESS == status)
 552                 {
 553                     SAL_INFO("opencl.device", "Scores written to file (" << sFilePath << ").");
 554                 }
 555                 else
 556                 {
 557                     SAL_INFO("opencl.device", "Error saving scores to file (" << sFilePath << "); scores not written to file.");
 558                 }
 559             }
 560             else
 561             {
 562                 SAL_INFO("opencl.device", "Unable to evaluate performance; scores not written to file.");
 563             }
 564         }
 565         else
 566         {
 567             SAL_INFO("opencl.device", "Profile read from file (" << sFilePath << ").");
 568         }
 569
 570         /* Pick best device */
 571         int bestDeviceIdx = pickBestDevice(aProfile);
 572
 573         /* Override if necessary */
 574         char* overrideDeviceStr = getenv("SC_OPENCL_DEVICE_OVERRIDE");
 575         if (nullptr != overrideDeviceStr)
 576         {
 577             int overrideDeviceIdx = matchDevice(aProfile, overrideDeviceStr);
 578             if (-1 != overrideDeviceIdx)
 579             {
 580                 SAL_INFO("opencl.device", "Overriding Device Selection (SC_OPENCL_DEVICE_OVERRIDE=" << overrideDeviceStr << ").");
 581                 bestDeviceIdx = overrideDeviceIdx;
 582                 if (aProfile->devices[bestDeviceIdx].eType == DeviceType::OpenCLDevice)
 583                 {
 584                     SAL_INFO("opencl.device", "Selected Device[" << bestDeviceIdx << "]: " << aProfile->devices[bestDeviceIdx].sDeviceName << " (OpenCL).");
 585                 }
 586                 else
 587                 {
 588                     SAL_INFO("opencl.device", "Selected Device[" << bestDeviceIdx << "]: CPU (Native).");
 589                 }
 590             }
 591             else
 592             {
 593                 SAL_INFO("opencl.device", "Ignoring invalid SC_OPENCL_DEVICE_OVERRIDE=" << overrideDeviceStr << ").");
 594             }
 595         }
 596
 597         /* Final device selection */
 598         if (bestDeviceIdx >=0 && o3tl::make_unsigned( bestDeviceIdx ) < aProfile->devices.size() )
 599         {
 600             selectedDevice = aProfile->devices[bestDeviceIdx];
 601             bIsDeviceSelected = true;
 602
 603             writeDevicesLog(aProfile, sProfilePath, bestDeviceIdx);
 604         } else {
 605             selectedDevice.eType = DeviceType::NativeCPU;
 606         }
 607     }
 608     return selectedDevice;
 609 }
 610
 611 /* vim:set shiftwidth=4 softtabstop=4 expandtab: */