src/gromacs/hardware/device_management_ocl.cpp

   1 /*
   2  * This file is part of the GROMACS molecular simulation package.
   3  *
   4  * Copyright (c) 2012,2013,2014,2015,2016, by the GROMACS development team.
   5  * Copyright (c) 2017,2018,2019,2020, by the GROMACS development team, led by
   6  * Mark Abraham, David van der Spoel, Berk Hess, and Erik Lindahl,
   7  * and including many others, as listed in the AUTHORS file in the
   8  * top-level source directory and at http://www.gromacs.org.
   9  *
  10  * GROMACS is free software; you can redistribute it and/or
  11  * modify it under the terms of the GNU Lesser General Public License
  12  * as published by the Free Software Foundation; either version 2.1
  13  * of the License, or (at your option) any later version.
  14  *
  15  * GROMACS is distributed in the hope that it will be useful,
  16  * but WITHOUT ANY WARRANTY; without even the implied warranty of
  17  * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
  18  * Lesser General Public License for more details.
  19  *
  20  * You should have received a copy of the GNU Lesser General Public
  21  * License along with GROMACS; if not, see
  22  * http://www.gnu.org/licenses, or write to the Free Software Foundation,
  23  * Inc., 51 Franklin Street, Fifth Floor, Boston, MA  02110-1301  USA.
  24  *
  25  * If you want to redistribute modifications to GROMACS, please
  26  * consider that scientific software is very special. Version
  27  * control is crucial - bugs must be traceable. We will be happy to
  28  * consider code for inclusion in the official distribution, but
  29  * derived work must not be called official GROMACS. Details are found
  30  * in the README & COPYING files - if they are missing, get the
  31  * official version at http://www.gromacs.org.
  32  *
  33  * To help us fund GROMACS development, we humbly ask that you cite
  34  * the research papers on the package. Check out http://www.gromacs.org.
  35  */
  36 /*! \internal \file
  37  *  \brief Defines the OpenCL implementations of the device management.
  38  *
  39  *  \author Anca Hamuraru <anca@streamcomputing.eu>
  40  *  \author Dimitrios Karkoulis <dimitris.karkoulis@gmail.com>
  41  *  \author Teemu Virolainen <teemu@streamcomputing.eu>
  42  *  \author Mark Abraham <mark.j.abraham@gmail.com>
  43  *  \author Szilárd Páll <pall.szilard@gmail.com>
  44  *  \author Artem Zhmurov <zhmurov@gmail.com>
  45  *
  46  * \ingroup module_hardware
  47  */
  48 #include "gmxpre.h"
  49
  50 #include "config.h"
  51
  52 #include "gromacs/gpu_utils/oclraii.h"
  53 #include "gromacs/gpu_utils/oclutils.h"
  54 #include "gromacs/hardware/device_management.h"
  55 #include "gromacs/utility/fatalerror.h"
  56 #include "gromacs/utility/smalloc.h"
  57 #include "gromacs/utility/stringutil.h"
  58
  59 #include "device_information.h"
  60
  61 namespace gmx
  62 {
  63
  64 /*! \brief Return true if executing on compatible OS for AMD OpenCL.
  65  *
  66  * This is assumed to be true for OS X version of at least 10.10.4 and
  67  * all other OS flavors.
  68  *
  69  * \return true if version is 14.4 or later (= OS X version 10.10.4),
  70  *         or OS is not Darwin.
  71  */
  72 static bool runningOnCompatibleOSForAmd()
  73 {
  74 #ifdef __APPLE__
  75     int    mib[2];
  76     char   kernelVersion[256];
  77     size_t len = sizeof(kernelVersion);
  78
  79     mib[0] = CTL_KERN;
  80
  81     int major = strtod(kernelVersion, NULL);
  82     int minor = strtod(strchr(kernelVersion, '.') + 1, NULL);
  83
  84     // Kernel 14.4 corresponds to OS X 10.10.4
  85     return (major > 14 || (major == 14 && minor >= 4));
  86 #else
  87     return true;
  88 #endif
  89 }
  90
  91 /*! \brief Return true if executing on compatible GPU for NVIDIA OpenCL.
  92  *
  93  * There are known issues with OpenCL when running on NVIDIA Volta or newer (CC 7+).
  94  * As a workaround, we recommend using CUDA on such hardware.
  95  *
  96  * This function relies on cl_nv_device_attribute_query. In case it's not functioning properly,
  97  * we trust the user and mark the device as compatible.
  98  *
  99  * \return true if running on Pascal (CC 6.x) or older, or if we can not determine device generation.
 100  */
 101 static bool runningOnCompatibleHWForNvidia(const DeviceInformation& deviceInfo)
 102 {
 103     // The macro is defined in Intel's and AMD's headers, but it's not strictly required to be there.
 104 #ifndef CL_DEVICE_COMPUTE_CAPABILITY_MAJOR_NV
 105     return true;
 106 #else
 107     static const unsigned int ccMajorBad = 7; // Volta and Turing
 108     unsigned int              ccMajor;
 109     cl_device_id              devId = deviceInfo.oclDeviceId;
 110     const cl_int              err   = clGetDeviceInfo(
 111             devId, CL_DEVICE_COMPUTE_CAPABILITY_MAJOR_NV, sizeof(ccMajor), &ccMajor, nullptr);
 112     if (err != CL_SUCCESS)
 113     {
 114         return true; // Err on a side of trusting the user to know what they are doing.
 115     }
 116     return ccMajor < ccMajorBad;
 117 #endif
 118 }
 119
 120 /*!
 121  * \brief Checks that device \c deviceInfo is compatible with GROMACS.
 122  *
 123  *  Vendor and OpenCL version support checks are executed an the result
 124  *  of these returned.
 125  *
 126  * \param[in]  deviceInfo  The device info pointer.
 127  * \returns                The status enumeration value for the checked device:
 128  */
 129 static DeviceStatus isDeviceFunctional(const DeviceInformation& deviceInfo)
 130 {
 131     if (getenv("GMX_GPU_DISABLE_COMPATIBILITY_CHECK") != nullptr)
 132     {
 133         // Assume the device is compatible because checking has been disabled.
 134         return DeviceStatus::Compatible;
 135     }
 136     if (getenv("GMX_OCL_DISABLE_COMPATIBILITY_CHECK") != nullptr)
 137     {
 138         fprintf(stderr,
 139                 "Environment variable GMX_OCL_DISABLE_COMPATIBILITY_CHECK is deprecated and will "
 140                 "be removed in release 2022. Please use GMX_GPU_DISABLE_COMPATIBILITY_CHECK "
 141                 "instead.\n");
 142         return DeviceStatus::Compatible;
 143     }
 144
 145     // OpenCL device version check, ensure >= REQUIRED_OPENCL_MIN_VERSION
 146     constexpr unsigned int minVersionMajor = REQUIRED_OPENCL_MIN_VERSION_MAJOR;
 147     constexpr unsigned int minVersionMinor = REQUIRED_OPENCL_MIN_VERSION_MINOR;
 148
 149     // Based on the OpenCL spec we're checking the version supported by
 150     // the device which has the following format:
 151     //      OpenCL<space><major_version.minor_version><space><vendor-specific information>
 152     unsigned int deviceVersionMinor, deviceVersionMajor;
 153     const int    valuesScanned = std::sscanf(
 154             deviceInfo.device_version, "OpenCL %u.%u", &deviceVersionMajor, &deviceVersionMinor);
 155     const bool versionLargeEnough =
 156             ((valuesScanned == 2)
 157              && ((deviceVersionMajor > minVersionMajor)
 158                  || (deviceVersionMajor == minVersionMajor && deviceVersionMinor >= minVersionMinor)));
 159     if (!versionLargeEnough)
 160     {
 161         return DeviceStatus::Incompatible;
 162     }
 163
 164     /* Only AMD, Intel, and NVIDIA GPUs are supported for now */
 165     switch (deviceInfo.deviceVendor)
 166     {
 167         case DeviceVendor::Nvidia:
 168             return runningOnCompatibleHWForNvidia(deviceInfo) ? DeviceStatus::Compatible
 169                                                               : DeviceStatus::IncompatibleNvidiaVolta;
 170         case DeviceVendor::Amd:
 171             return runningOnCompatibleOSForAmd() ? DeviceStatus::Compatible : DeviceStatus::Incompatible;
 172         case DeviceVendor::Intel:
 173             return GMX_OPENCL_NB_CLUSTER_SIZE == 4 ? DeviceStatus::Compatible
 174                                                    : DeviceStatus::IncompatibleClusterSize;
 175         default: return DeviceStatus::Incompatible;
 176     }
 177 }
 178
 179 /*! \brief Make an error string following an OpenCL API call.
 180  *
 181  *  It is meant to be called with \p status != CL_SUCCESS, but it will
 182  *  work correctly even if it is called with no OpenCL failure.
 183  *
 184  * \todo Make use of this function more.
 185  *
 186  * \param[in]  message  Supplies context, e.g. the name of the API call that returned the error.
 187  * \param[in]  status   OpenCL API status code
 188  * \returns             A string describing the OpenCL error.
 189  */
 190 inline std::string makeOpenClInternalErrorString(const char* message, cl_int status)
 191 {
 192     if (message != nullptr)
 193     {
 194         return gmx::formatString("%s did %ssucceed %d: %s",
 195                                  message,
 196                                  ((status != CL_SUCCESS) ? "not " : ""),
 197                                  status,
 198                                  ocl_get_error_string(status).c_str());
 199     }
 200     else
 201     {
 202         return gmx::formatString("%sOpenCL error encountered %d: %s",
 203                                  ((status != CL_SUCCESS) ? "" : "No "),
 204                                  status,
 205                                  ocl_get_error_string(status).c_str());
 206     }
 207 }
 208
 209 /*!
 210  * \brief Checks that device \c deviceInfo is sane (ie can run a kernel).
 211  *
 212  * Compiles and runs a dummy kernel to determine whether the given
 213  * OpenCL device functions properly.
 214  *
 215  *
 216  * \param[in]  deviceInfo      The device info pointer.
 217  * \param[out] errorMessage    An error message related to a failing OpenCL API call.
 218  * \throws     std::bad_alloc  When out of memory.
 219  * \returns                    Whether the device passed sanity checks
 220  */
 221 static bool isDeviceFunctional(const DeviceInformation& deviceInfo, std::string* errorMessage)
 222 {
 223     cl_context_properties properties[] = {
 224         CL_CONTEXT_PLATFORM, reinterpret_cast<cl_context_properties>(deviceInfo.oclPlatformId), 0
 225     };
 226     // uncrustify spacing
 227
 228     cl_int    status;
 229     auto      deviceId = deviceInfo.oclDeviceId;
 230     ClContext context(clCreateContext(properties, 1, &deviceId, nullptr, nullptr, &status));
 231     if (status != CL_SUCCESS)
 232     {
 233         errorMessage->assign(makeOpenClInternalErrorString("clCreateContext", status));
 234         return false;
 235     }
 236     ClCommandQueue commandQueue(clCreateCommandQueue(context, deviceId, 0, &status));
 237     if (status != CL_SUCCESS)
 238     {
 239         errorMessage->assign(makeOpenClInternalErrorString("clCreateCommandQueue", status));
 240         return false;
 241     }
 242
 243     // Some compilers such as Apple's require kernel functions to have at least one argument
 244     const char* lines[] = { "__kernel void dummyKernel(__global void* input){}" };
 245     ClProgram   program(clCreateProgramWithSource(context, 1, lines, nullptr, &status));
 246     if (status != CL_SUCCESS)
 247     {
 248         errorMessage->assign(makeOpenClInternalErrorString("clCreateProgramWithSource", status));
 249         return false;
 250     }
 251
 252     if ((status = clBuildProgram(program, 0, nullptr, nullptr, nullptr, nullptr)) != CL_SUCCESS)
 253     {
 254         errorMessage->assign(makeOpenClInternalErrorString("clBuildProgram", status));
 255         return false;
 256     }
 257
 258     ClKernel kernel(clCreateKernel(program, "dummyKernel", &status));
 259     if (status != CL_SUCCESS)
 260     {
 261         errorMessage->assign(makeOpenClInternalErrorString("clCreateKernel", status));
 262         return false;
 263     }
 264
 265     clSetKernelArg(kernel, 0, sizeof(void*), nullptr);
 266
 267     const size_t localWorkSize = 1, globalWorkSize = 1;
 268     if ((status = clEnqueueNDRangeKernel(
 269                  commandQueue, kernel, 1, nullptr, &globalWorkSize, &localWorkSize, 0, nullptr, nullptr))
 270         != CL_SUCCESS)
 271     {
 272         errorMessage->assign(makeOpenClInternalErrorString("clEnqueueNDRangeKernel", status));
 273         return false;
 274     }
 275     return true;
 276 }
 277
 278 /*! \brief Check whether the \c ocl_gpu_device is suitable for use by mdrun
 279  *
 280  * Runs sanity checks: checking that the runtime can compile a dummy kernel
 281  * and this can be executed;
 282  * Runs compatibility checks verifying the device OpenCL version requirement
 283  * and vendor/OS support.
 284  *
 285  * \param[in]  deviceId      The runtime-reported numeric ID of the device.
 286  * \param[in]  deviceInfo    The device info pointer.
 287  * \returns  A DeviceStatus to indicate if the GPU device is supported and if it was able to run
 288  *           basic functionality checks.
 289  */
 290 static DeviceStatus checkGpu(size_t deviceId, const DeviceInformation& deviceInfo)
 291 {
 292
 293     DeviceStatus supportStatus = isDeviceFunctional(deviceInfo);
 294     if (supportStatus != DeviceStatus::Compatible)
 295     {
 296         return supportStatus;
 297     }
 298
 299     std::string errorMessage;
 300     if (!isDeviceFunctional(deviceInfo, &errorMessage))
 301     {
 302         gmx_warning("While sanity checking device #%zu, %s", deviceId, errorMessage.c_str());
 303         return DeviceStatus::NonFunctional;
 304     }
 305
 306     return DeviceStatus::Compatible;
 307 }
 308
 309 } // namespace gmx
 310
 311 bool isDeviceDetectionFunctional(std::string* errorMessage)
 312 {
 313     cl_uint numPlatforms;
 314     cl_int  status = clGetPlatformIDs(0, nullptr, &numPlatforms);
 315     GMX_ASSERT(status != CL_INVALID_VALUE, "Incorrect call of clGetPlatformIDs detected");
 316 #ifdef cl_khr_icd
 317     if (status == CL_PLATFORM_NOT_FOUND_KHR)
 318     {
 319         // No valid ICDs found
 320         if (errorMessage != nullptr)
 321         {
 322             errorMessage->assign("No valid OpenCL driver found");
 323         }
 324         return false;
 325     }
 326 #endif
 327     GMX_RELEASE_ASSERT(
 328             status == CL_SUCCESS,
 329             gmx::formatString("An unexpected value was returned from clGetPlatformIDs %d: %s",
 330                               status,
 331                               ocl_get_error_string(status).c_str())
 332                     .c_str());
 333     bool foundPlatform = (numPlatforms > 0);
 334     if (!foundPlatform && errorMessage != nullptr)
 335     {
 336         errorMessage->assign("No OpenCL platforms found even though the driver was valid");
 337     }
 338     return foundPlatform;
 339 }
 340
 341 std::vector<std::unique_ptr<DeviceInformation>> findDevices()
 342 {
 343     cl_uint         ocl_platform_count;
 344     cl_platform_id* ocl_platform_ids;
 345     cl_device_type  req_dev_type = CL_DEVICE_TYPE_GPU;
 346
 347     ocl_platform_ids = nullptr;
 348
 349     if (getenv("GMX_OCL_FORCE_CPU") != nullptr)
 350     {
 351         req_dev_type = CL_DEVICE_TYPE_CPU;
 352     }
 353
 354     int                                             numDevices = 0;
 355     std::vector<std::unique_ptr<DeviceInformation>> deviceInfoList(0);
 356
 357     while (true)
 358     {
 359         cl_int status = clGetPlatformIDs(0, nullptr, &ocl_platform_count);
 360         if (CL_SUCCESS != status)
 361         {
 362             GMX_THROW(gmx::InternalError(
 363                     gmx::formatString("An unexpected value %d was returned from clGetPlatformIDs: ", status)
 364                     + ocl_get_error_string(status)));
 365         }
 366
 367         if (1 > ocl_platform_count)
 368         {
 369             // TODO this should have a descriptive error message that we only support one OpenCL platform
 370             break;
 371         }
 372
 373         snew(ocl_platform_ids, ocl_platform_count);
 374
 375         status = clGetPlatformIDs(ocl_platform_count, ocl_platform_ids, nullptr);
 376         if (CL_SUCCESS != status)
 377         {
 378             GMX_THROW(gmx::InternalError(
 379                     gmx::formatString("An unexpected value %d was returned from clGetPlatformIDs: ", status)
 380                     + ocl_get_error_string(status)));
 381         }
 382
 383         for (unsigned int i = 0; i < ocl_platform_count; i++)
 384         {
 385             cl_uint ocl_device_count;
 386
 387             /* If requesting req_dev_type devices fails, just go to the next platform */
 388             if (CL_SUCCESS != clGetDeviceIDs(ocl_platform_ids[i], req_dev_type, 0, nullptr, &ocl_device_count))
 389             {
 390                 continue;
 391             }
 392
 393             if (1 <= ocl_device_count)
 394             {
 395                 numDevices += ocl_device_count;
 396             }
 397         }
 398
 399         if (1 > numDevices)
 400         {
 401             break;
 402         }
 403
 404         deviceInfoList.resize(numDevices);
 405
 406         {
 407             int           device_index;
 408             cl_device_id* ocl_device_ids;
 409
 410             snew(ocl_device_ids, numDevices);
 411             device_index = 0;
 412
 413             for (unsigned int i = 0; i < ocl_platform_count; i++)
 414             {
 415                 cl_uint ocl_device_count;
 416
 417                 /* If requesting req_dev_type devices fails, just go to the next platform */
 418                 if (CL_SUCCESS
 419                     != clGetDeviceIDs(ocl_platform_ids[i], req_dev_type, numDevices, ocl_device_ids, &ocl_device_count))
 420                 {
 421                     continue;
 422                 }
 423
 424                 if (1 > ocl_device_count)
 425                 {
 426                     break;
 427                 }
 428
 429                 for (unsigned int j = 0; j < ocl_device_count; j++)
 430                 {
 431                     deviceInfoList[device_index] = std::make_unique<DeviceInformation>();
 432
 433                     deviceInfoList[device_index]->id = device_index;
 434
 435                     deviceInfoList[device_index]->oclPlatformId = ocl_platform_ids[i];
 436                     deviceInfoList[device_index]->oclDeviceId   = ocl_device_ids[j];
 437
 438                     deviceInfoList[device_index]->device_name[0] = 0;
 439                     clGetDeviceInfo(ocl_device_ids[j],
 440                                     CL_DEVICE_NAME,
 441                                     sizeof(deviceInfoList[device_index]->device_name),
 442                                     deviceInfoList[device_index]->device_name,
 443                                     nullptr);
 444
 445                     deviceInfoList[device_index]->device_version[0] = 0;
 446                     clGetDeviceInfo(ocl_device_ids[j],
 447                                     CL_DEVICE_VERSION,
 448                                     sizeof(deviceInfoList[device_index]->device_version),
 449                                     deviceInfoList[device_index]->device_version,
 450                                     nullptr);
 451
 452                     deviceInfoList[device_index]->vendorName[0] = 0;
 453                     clGetDeviceInfo(ocl_device_ids[j],
 454                                     CL_DEVICE_VENDOR,
 455                                     sizeof(deviceInfoList[device_index]->vendorName),
 456                                     deviceInfoList[device_index]->vendorName,
 457                                     nullptr);
 458
 459                     deviceInfoList[device_index]->compute_units = 0;
 460                     clGetDeviceInfo(ocl_device_ids[j],
 461                                     CL_DEVICE_MAX_COMPUTE_UNITS,
 462                                     sizeof(deviceInfoList[device_index]->compute_units),
 463                                     &(deviceInfoList[device_index]->compute_units),
 464                                     nullptr);
 465
 466                     deviceInfoList[device_index]->adress_bits = 0;
 467                     clGetDeviceInfo(ocl_device_ids[j],
 468                                     CL_DEVICE_ADDRESS_BITS,
 469                                     sizeof(deviceInfoList[device_index]->adress_bits),
 470                                     &(deviceInfoList[device_index]->adress_bits),
 471                                     nullptr);
 472
 473                     deviceInfoList[device_index]->deviceVendor =
 474                             getDeviceVendor(deviceInfoList[device_index]->vendorName);
 475
 476                     clGetDeviceInfo(ocl_device_ids[j],
 477                                     CL_DEVICE_MAX_WORK_ITEM_SIZES,
 478                                     3 * sizeof(size_t),
 479                                     &deviceInfoList[device_index]->maxWorkItemSizes,
 480                                     nullptr);
 481
 482                     clGetDeviceInfo(ocl_device_ids[j],
 483                                     CL_DEVICE_MAX_WORK_GROUP_SIZE,
 484                                     sizeof(size_t),
 485                                     &deviceInfoList[device_index]->maxWorkGroupSize,
 486                                     nullptr);
 487
 488                     deviceInfoList[device_index]->status =
 489                             gmx::checkGpu(device_index, *deviceInfoList[device_index]);
 490
 491                     device_index++;
 492                 }
 493             }
 494
 495             numDevices = device_index;
 496
 497             /* Dummy sort of devices -  AMD first, then NVIDIA, then Intel */
 498             // TODO: Sort devices based on performance.
 499             if (0 < numDevices)
 500             {
 501                 int last = -1;
 502                 for (int i = 0; i < numDevices; i++)
 503                 {
 504                     if (deviceInfoList[i]->deviceVendor == DeviceVendor::Amd)
 505                     {
 506                         last++;
 507
 508                         if (last < i)
 509                         {
 510                             std::swap(deviceInfoList[i], deviceInfoList[last]);
 511                         }
 512                     }
 513                 }
 514
 515                 /* if more than 1 device left to be sorted */
 516                 if ((numDevices - 1 - last) > 1)
 517                 {
 518                     for (int i = 0; i < numDevices; i++)
 519                     {
 520                         if (deviceInfoList[i]->deviceVendor == DeviceVendor::Nvidia)
 521                         {
 522                             last++;
 523
 524                             if (last < i)
 525                             {
 526                                 std::swap(deviceInfoList[i], deviceInfoList[last]);
 527                             }
 528                         }
 529                     }
 530                 }
 531             }
 532
 533             sfree(ocl_device_ids);
 534         }
 535
 536         break;
 537     }
 538
 539     sfree(ocl_platform_ids);
 540     return deviceInfoList;
 541 }
 542
 543 void setActiveDevice(const DeviceInformation& deviceInfo)
 544 {
 545     // If the device is NVIDIA, for safety reasons we disable the JIT
 546     // caching as this is known to be broken at least until driver 364.19;
 547     // the cache does not always get regenerated when the source code changes,
 548     // e.g. if the path to the kernel sources remains the same
 549
 550     if (deviceInfo.deviceVendor == DeviceVendor::Nvidia)
 551     {
 552         // Ignore return values, failing to set the variable does not mean
 553         // that something will go wrong later.
 554 #ifdef _MSC_VER
 555         _putenv("CUDA_CACHE_DISABLE=1");
 556 #else
 557         // Don't override, maybe a dev is testing.
 558         setenv("CUDA_CACHE_DISABLE", "1", 0);
 559 #endif
 560     }
 561 }
 562
 563 void releaseDevice(DeviceInformation* /* deviceInfo */) {}
 564
 565 std::string getDeviceInformationString(const DeviceInformation& deviceInfo)
 566 {
 567     bool gpuExists = (deviceInfo.status != DeviceStatus::Nonexistent
 568                       && deviceInfo.status != DeviceStatus::NonFunctional);
 569
 570     if (!gpuExists)
 571     {
 572         return gmx::formatString(
 573                 "#%d: %s, status: %s", deviceInfo.id, "N/A", c_deviceStateString[deviceInfo.status]);
 574     }
 575     else
 576     {
 577         return gmx::formatString("#%d: name: %s, vendor: %s, device version: %s, status: %s",
 578                                  deviceInfo.id,
 579                                  deviceInfo.device_name,
 580                                  deviceInfo.vendorName,
 581                                  deviceInfo.device_version,
 582                                  c_deviceStateString[deviceInfo.status]);
 583     }
 584 }