src/gromacs/gpu_utils/gpu_utils_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 Define functions for detection and initialization for OpenCL devices.
  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  */
  45
  46 #include "gmxpre.h"
  47
  48 #include "config.h"
  49
  50 #include <assert.h>
  51 #include <stdio.h>
  52 #include <stdlib.h>
  53 #include <string.h>
  54
  55 #include <cstdio>
  56 #ifdef __APPLE__
  57 #    include <sys/sysctl.h>
  58 #endif
  59
  60 #include <memory.h>
  61
  62 #include "gromacs/gpu_utils/gpu_utils.h"
  63 #include "gromacs/gpu_utils/ocl_compiler.h"
  64 #include "gromacs/gpu_utils/oclraii.h"
  65 #include "gromacs/gpu_utils/oclutils.h"
  66 #include "gromacs/hardware/hw_info.h"
  67 #include "gromacs/utility/cstringutil.h"
  68 #include "gromacs/utility/exceptions.h"
  69 #include "gromacs/utility/fatalerror.h"
  70 #include "gromacs/utility/smalloc.h"
  71 #include "gromacs/utility/stringutil.h"
  72
  73 /*! \brief Return true if executing on compatible OS for AMD OpenCL.
  74  *
  75  * This is assumed to be true for OS X version of at least 10.10.4 and
  76  * all other OS flavors.
  77  *
  78  * Uses the BSD sysctl() interfaces to extract the kernel version.
  79  *
  80  * \return true if version is 14.4 or later (= OS X version 10.10.4),
  81  *         or OS is not Darwin.
  82  */
  83 static bool runningOnCompatibleOSForAmd()
  84 {
  85 #ifdef __APPLE__
  86     int    mib[2];
  87     char   kernelVersion[256];
  88     size_t len = sizeof(kernelVersion);
  89
  90     mib[0] = CTL_KERN;
  91     mib[1] = KERN_OSRELEASE;
  92
  93     sysctl(mib, sizeof(mib) / sizeof(mib[0]), kernelVersion, &len, NULL, 0);
  94
  95     int major = strtod(kernelVersion, NULL);
  96     int minor = strtod(strchr(kernelVersion, '.') + 1, NULL);
  97
  98     // Kernel 14.4 corresponds to OS X 10.10.4
  99     return (major > 14 || (major == 14 && minor >= 4));
 100 #else
 101     return true;
 102 #endif
 103 }
 104
 105 namespace gmx
 106 {
 107
 108 /*! \brief Make an error string following an OpenCL API call.
 109  *
 110  *  It is meant to be called with \p status != CL_SUCCESS, but it will
 111  *  work correctly even if it is called with no OpenCL failure.
 112  *
 113  * \param[in]  message  Supplies context, e.g. the name of the API call that returned the error.
 114  * \param[in]  status   OpenCL API status code
 115  * \returns             A string describing the OpenCL error.
 116  */
 117 static std::string makeOpenClInternalErrorString(const char* message, cl_int status)
 118 {
 119     if (message != nullptr)
 120     {
 121         return formatString("%s did %ssucceed %d: %s", message, ((status != CL_SUCCESS) ? "not " : ""),
 122                             status, ocl_get_error_string(status).c_str());
 123     }
 124     else
 125     {
 126         return formatString("%sOpenCL error encountered %d: %s", ((status != CL_SUCCESS) ? "" : "No "),
 127                             status, ocl_get_error_string(status).c_str());
 128     }
 129 }
 130
 131 /*!
 132  * \brief Checks that device \c deviceInfo is sane (ie can run a kernel).
 133  *
 134  * Compiles and runs a dummy kernel to determine whether the given
 135  * OpenCL device functions properly.
 136  *
 137  *
 138  * \param[in]  deviceInfo      The device info pointer.
 139  * \param[out] errorMessage    An error message related to a failing OpenCL API call.
 140  * \throws     std::bad_alloc  When out of memory.
 141  * \returns                    Whether the device passed sanity checks
 142  */
 143 static bool isDeviceFunctional(const DeviceInformation* deviceInfo, std::string* errorMessage)
 144 {
 145     cl_context_properties properties[] = {
 146         CL_CONTEXT_PLATFORM, reinterpret_cast<cl_context_properties>(deviceInfo->oclPlatformId), 0
 147     };
 148     // uncrustify spacing
 149
 150     cl_int    status;
 151     auto      deviceId = deviceInfo->oclDeviceId;
 152     ClContext context(clCreateContext(properties, 1, &deviceId, nullptr, nullptr, &status));
 153     if (status != CL_SUCCESS)
 154     {
 155         errorMessage->assign(makeOpenClInternalErrorString("clCreateContext", status));
 156         return false;
 157     }
 158     ClCommandQueue commandQueue(clCreateCommandQueue(context, deviceId, 0, &status));
 159     if (status != CL_SUCCESS)
 160     {
 161         errorMessage->assign(makeOpenClInternalErrorString("clCreateCommandQueue", status));
 162         return false;
 163     }
 164
 165     // Some compilers such as Apple's require kernel functions to have at least one argument
 166     const char* lines[] = { "__kernel void dummyKernel(__global void* input){}" };
 167     ClProgram   program(clCreateProgramWithSource(context, 1, lines, nullptr, &status));
 168     if (status != CL_SUCCESS)
 169     {
 170         errorMessage->assign(makeOpenClInternalErrorString("clCreateProgramWithSource", status));
 171         return false;
 172     }
 173
 174     if ((status = clBuildProgram(program, 0, nullptr, nullptr, nullptr, nullptr)) != CL_SUCCESS)
 175     {
 176         errorMessage->assign(makeOpenClInternalErrorString("clBuildProgram", status));
 177         return false;
 178     }
 179
 180     ClKernel kernel(clCreateKernel(program, "dummyKernel", &status));
 181     if (status != CL_SUCCESS)
 182     {
 183         errorMessage->assign(makeOpenClInternalErrorString("clCreateKernel", status));
 184         return false;
 185     }
 186
 187     clSetKernelArg(kernel, 0, sizeof(void*), nullptr);
 188
 189     const size_t localWorkSize = 1, globalWorkSize = 1;
 190     if ((status = clEnqueueNDRangeKernel(commandQueue, kernel, 1, nullptr, &globalWorkSize,
 191                                          &localWorkSize, 0, nullptr, nullptr))
 192         != CL_SUCCESS)
 193     {
 194         errorMessage->assign(makeOpenClInternalErrorString("clEnqueueNDRangeKernel", status));
 195         return false;
 196     }
 197     return true;
 198 }
 199
 200 /*!
 201  * \brief Checks that device \c deviceInfo is compatible with GROMACS.
 202  *
 203  *  Vendor and OpenCL version support checks are executed an the result
 204  *  of these returned.
 205  *
 206  * \param[in]  deviceInfo  The device info pointer.
 207  * \returns                The result of the compatibility checks.
 208  */
 209 static DeviceStatus isDeviceSupported(const DeviceInformation* deviceInfo)
 210 {
 211     if (getenv("GMX_OCL_DISABLE_COMPATIBILITY_CHECK") != nullptr)
 212     {
 213         // Assume the device is compatible because checking has been disabled.
 214         return DeviceStatus::Compatible;
 215     }
 216
 217     // OpenCL device version check, ensure >= REQUIRED_OPENCL_MIN_VERSION
 218     constexpr unsigned int minVersionMajor = REQUIRED_OPENCL_MIN_VERSION_MAJOR;
 219     constexpr unsigned int minVersionMinor = REQUIRED_OPENCL_MIN_VERSION_MINOR;
 220
 221     // Based on the OpenCL spec we're checking the version supported by
 222     // the device which has the following format:
 223     //      OpenCL<space><major_version.minor_version><space><vendor-specific information>
 224     unsigned int deviceVersionMinor, deviceVersionMajor;
 225     const int    valuesScanned = std::sscanf(deviceInfo->device_version, "OpenCL %u.%u",
 226                                           &deviceVersionMajor, &deviceVersionMinor);
 227     const bool   versionLargeEnough =
 228             ((valuesScanned == 2)
 229              && ((deviceVersionMajor > minVersionMajor)
 230                  || (deviceVersionMajor == minVersionMajor && deviceVersionMinor >= minVersionMinor)));
 231     if (!versionLargeEnough)
 232     {
 233         return DeviceStatus::Incompatible;
 234     }
 235
 236     /* Only AMD, Intel, and NVIDIA GPUs are supported for now */
 237     switch (deviceInfo->deviceVendor)
 238     {
 239         case DeviceVendor::Nvidia: return DeviceStatus::Compatible;
 240         case DeviceVendor::Amd:
 241             return runningOnCompatibleOSForAmd() ? DeviceStatus::Compatible : DeviceStatus::Incompatible;
 242         case DeviceVendor::Intel:
 243             return GMX_OPENCL_NB_CLUSTER_SIZE == 4 ? DeviceStatus::Compatible
 244                                                    : DeviceStatus::IncompatibleClusterSize;
 245         default: return DeviceStatus::Incompatible;
 246     }
 247 }
 248
 249
 250 /*! \brief Check whether the \c ocl_gpu_device is suitable for use by mdrun
 251  *
 252  * Runs sanity checks: checking that the runtime can compile a dummy kernel
 253  * and this can be executed;
 254  * Runs compatibility checks verifying the device OpenCL version requirement
 255  * and vendor/OS support.
 256  *
 257  * \param[in]  deviceId      The runtime-reported numeric ID of the device.
 258  * \param[in]  deviceInfo    The device info pointer.
 259  * \returns  A DeviceStatus to indicate if the GPU device is supported and if it was able to run
 260  *           basic functionality checks.
 261  */
 262 static DeviceStatus checkGpu(size_t deviceId, const DeviceInformation* deviceInfo)
 263 {
 264
 265     DeviceStatus supportStatus = isDeviceSupported(deviceInfo);
 266     if (supportStatus != DeviceStatus::Compatible)
 267     {
 268         return supportStatus;
 269     }
 270
 271     std::string errorMessage;
 272     if (!isDeviceFunctional(deviceInfo, &errorMessage))
 273     {
 274         gmx_warning("While sanity checking device #%zu, %s", deviceId, errorMessage.c_str());
 275         return DeviceStatus::NonFunctional;
 276     }
 277
 278     return DeviceStatus::Compatible;
 279 }
 280
 281 } // namespace gmx
 282
 283 /*! \brief Returns an DeviceVendor value corresponding to the input OpenCL vendor name.
 284  *
 285  *  \param[in] vendorName  String with OpenCL vendor name.
 286  *  \returns               DeviceVendor value for the input vendor name
 287  */
 288 static DeviceVendor getDeviceVendor(const char* vendorName)
 289 {
 290     if (vendorName)
 291     {
 292         if (strstr(vendorName, "NVIDIA"))
 293         {
 294             return DeviceVendor::Nvidia;
 295         }
 296         else if (strstr(vendorName, "AMD") || strstr(vendorName, "Advanced Micro Devices"))
 297         {
 298             return DeviceVendor::Amd;
 299         }
 300         else if (strstr(vendorName, "Intel"))
 301         {
 302             return DeviceVendor::Intel;
 303         }
 304     }
 305     return DeviceVendor::Unknown;
 306 }
 307
 308 bool isGpuDetectionFunctional(std::string* errorMessage)
 309 {
 310     cl_uint numPlatforms;
 311     cl_int  status = clGetPlatformIDs(0, nullptr, &numPlatforms);
 312     GMX_ASSERT(status != CL_INVALID_VALUE, "Incorrect call of clGetPlatformIDs detected");
 313 #ifdef cl_khr_icd
 314     if (status == CL_PLATFORM_NOT_FOUND_KHR)
 315     {
 316         // No valid ICDs found
 317         if (errorMessage != nullptr)
 318         {
 319             errorMessage->assign("No valid OpenCL driver found");
 320         }
 321         return false;
 322     }
 323 #endif
 324     GMX_RELEASE_ASSERT(
 325             status == CL_SUCCESS,
 326             gmx::formatString("An unexpected value was returned from clGetPlatformIDs %d: %s",
 327                               status, ocl_get_error_string(status).c_str())
 328                     .c_str());
 329     bool foundPlatform = (numPlatforms > 0);
 330     if (!foundPlatform && errorMessage != nullptr)
 331     {
 332         errorMessage->assign("No OpenCL platforms found even though the driver was valid");
 333     }
 334     return foundPlatform;
 335 }
 336
 337 void findGpus(gmx_gpu_info_t* gpu_info)
 338 {
 339     cl_uint         ocl_platform_count;
 340     cl_platform_id* ocl_platform_ids;
 341     cl_device_type  req_dev_type = CL_DEVICE_TYPE_GPU;
 342
 343     ocl_platform_ids = nullptr;
 344
 345     if (getenv("GMX_OCL_FORCE_CPU") != nullptr)
 346     {
 347         req_dev_type = CL_DEVICE_TYPE_CPU;
 348     }
 349
 350     while (true)
 351     {
 352         cl_int status = clGetPlatformIDs(0, nullptr, &ocl_platform_count);
 353         if (CL_SUCCESS != status)
 354         {
 355             GMX_THROW(gmx::InternalError(
 356                     gmx::formatString("An unexpected value %d was returned from clGetPlatformIDs: ", status)
 357                     + ocl_get_error_string(status)));
 358         }
 359
 360         if (1 > ocl_platform_count)
 361         {
 362             // TODO this should have a descriptive error message that we only support one OpenCL platform
 363             break;
 364         }
 365
 366         snew(ocl_platform_ids, ocl_platform_count);
 367
 368         status = clGetPlatformIDs(ocl_platform_count, ocl_platform_ids, nullptr);
 369         if (CL_SUCCESS != status)
 370         {
 371             GMX_THROW(gmx::InternalError(
 372                     gmx::formatString("An unexpected value %d was returned from clGetPlatformIDs: ", status)
 373                     + ocl_get_error_string(status)));
 374         }
 375
 376         for (unsigned int i = 0; i < ocl_platform_count; i++)
 377         {
 378             cl_uint ocl_device_count;
 379
 380             /* If requesting req_dev_type devices fails, just go to the next platform */
 381             if (CL_SUCCESS != clGetDeviceIDs(ocl_platform_ids[i], req_dev_type, 0, nullptr, &ocl_device_count))
 382             {
 383                 continue;
 384             }
 385
 386             if (1 <= ocl_device_count)
 387             {
 388                 gpu_info->n_dev += ocl_device_count;
 389             }
 390         }
 391
 392         if (1 > gpu_info->n_dev)
 393         {
 394             break;
 395         }
 396
 397         snew(gpu_info->deviceInfo, gpu_info->n_dev);
 398
 399         {
 400             int           device_index;
 401             cl_device_id* ocl_device_ids;
 402
 403             snew(ocl_device_ids, gpu_info->n_dev);
 404             device_index = 0;
 405
 406             for (unsigned int i = 0; i < ocl_platform_count; i++)
 407             {
 408                 cl_uint ocl_device_count;
 409
 410                 /* If requesting req_dev_type devices fails, just go to the next platform */
 411                 if (CL_SUCCESS
 412                     != clGetDeviceIDs(ocl_platform_ids[i], req_dev_type, gpu_info->n_dev,
 413                                       ocl_device_ids, &ocl_device_count))
 414                 {
 415                     continue;
 416                 }
 417
 418                 if (1 > ocl_device_count)
 419                 {
 420                     break;
 421                 }
 422
 423                 for (unsigned int j = 0; j < ocl_device_count; j++)
 424                 {
 425                     gpu_info->deviceInfo[device_index].oclPlatformId = ocl_platform_ids[i];
 426                     gpu_info->deviceInfo[device_index].oclDeviceId   = ocl_device_ids[j];
 427
 428                     gpu_info->deviceInfo[device_index].device_name[0] = 0;
 429                     clGetDeviceInfo(ocl_device_ids[j], CL_DEVICE_NAME,
 430                                     sizeof(gpu_info->deviceInfo[device_index].device_name),
 431                                     gpu_info->deviceInfo[device_index].device_name, nullptr);
 432
 433                     gpu_info->deviceInfo[device_index].device_version[0] = 0;
 434                     clGetDeviceInfo(ocl_device_ids[j], CL_DEVICE_VERSION,
 435                                     sizeof(gpu_info->deviceInfo[device_index].device_version),
 436                                     gpu_info->deviceInfo[device_index].device_version, nullptr);
 437
 438                     gpu_info->deviceInfo[device_index].vendorName[0] = 0;
 439                     clGetDeviceInfo(ocl_device_ids[j], CL_DEVICE_VENDOR,
 440                                     sizeof(gpu_info->deviceInfo[device_index].vendorName),
 441                                     gpu_info->deviceInfo[device_index].vendorName, nullptr);
 442
 443                     gpu_info->deviceInfo[device_index].compute_units = 0;
 444                     clGetDeviceInfo(ocl_device_ids[j], CL_DEVICE_MAX_COMPUTE_UNITS,
 445                                     sizeof(gpu_info->deviceInfo[device_index].compute_units),
 446                                     &(gpu_info->deviceInfo[device_index].compute_units), nullptr);
 447
 448                     gpu_info->deviceInfo[device_index].adress_bits = 0;
 449                     clGetDeviceInfo(ocl_device_ids[j], CL_DEVICE_ADDRESS_BITS,
 450                                     sizeof(gpu_info->deviceInfo[device_index].adress_bits),
 451                                     &(gpu_info->deviceInfo[device_index].adress_bits), nullptr);
 452
 453                     gpu_info->deviceInfo[device_index].deviceVendor =
 454                             getDeviceVendor(gpu_info->deviceInfo[device_index].vendorName);
 455
 456                     clGetDeviceInfo(ocl_device_ids[j], CL_DEVICE_MAX_WORK_ITEM_SIZES, 3 * sizeof(size_t),
 457                                     &gpu_info->deviceInfo[device_index].maxWorkItemSizes, nullptr);
 458
 459                     clGetDeviceInfo(ocl_device_ids[j], CL_DEVICE_MAX_WORK_GROUP_SIZE, sizeof(size_t),
 460                                     &gpu_info->deviceInfo[device_index].maxWorkGroupSize, nullptr);
 461
 462                     gpu_info->deviceInfo[device_index].stat =
 463                             gmx::checkGpu(device_index, gpu_info->deviceInfo + device_index);
 464
 465                     if (DeviceStatus::Compatible == gpu_info->deviceInfo[device_index].stat)
 466                     {
 467                         gpu_info->n_dev_compatible++;
 468                     }
 469
 470                     device_index++;
 471                 }
 472             }
 473
 474             gpu_info->n_dev = device_index;
 475
 476             /* Dummy sort of devices -  AMD first, then NVIDIA, then Intel */
 477             // TODO: Sort devices based on performance.
 478             if (0 < gpu_info->n_dev)
 479             {
 480                 int last = -1;
 481                 for (int i = 0; i < gpu_info->n_dev; i++)
 482                 {
 483                     if (gpu_info->deviceInfo[i].deviceVendor == DeviceVendor::Amd)
 484                     {
 485                         last++;
 486
 487                         if (last < i)
 488                         {
 489                             std::swap(gpu_info->deviceInfo[i], gpu_info->deviceInfo[last]);
 490                         }
 491                     }
 492                 }
 493
 494                 /* if more than 1 device left to be sorted */
 495                 if ((gpu_info->n_dev - 1 - last) > 1)
 496                 {
 497                     for (int i = 0; i < gpu_info->n_dev; i++)
 498                     {
 499                         if (gpu_info->deviceInfo[i].deviceVendor == DeviceVendor::Nvidia)
 500                         {
 501                             last++;
 502
 503                             if (last < i)
 504                             {
 505                                 std::swap(gpu_info->deviceInfo[i], gpu_info->deviceInfo[last]);
 506                             }
 507                         }
 508                     }
 509                 }
 510             }
 511
 512             sfree(ocl_device_ids);
 513         }
 514
 515         break;
 516     }
 517
 518     sfree(ocl_platform_ids);
 519 }
 520
 521 void get_gpu_device_info_string(char* s, const gmx_gpu_info_t& gpu_info, int index)
 522 {
 523     assert(s);
 524
 525     if (index < 0 && index >= gpu_info.n_dev)
 526     {
 527         return;
 528     }
 529
 530     DeviceInformation* dinfo = &gpu_info.deviceInfo[index];
 531
 532     bool bGpuExists =
 533             (dinfo->stat != DeviceStatus::Nonexistent && dinfo->stat != DeviceStatus::NonFunctional);
 534
 535     if (!bGpuExists)
 536     {
 537         sprintf(s, "#%d: %s, stat: %s", index, "N/A", c_deviceStateString[dinfo->stat]);
 538     }
 539     else
 540     {
 541         sprintf(s, "#%d: name: %s, vendor: %s, device version: %s, stat: %s", index, dinfo->device_name,
 542                 dinfo->vendorName, dinfo->device_version, c_deviceStateString[dinfo->stat]);
 543     }
 544 }
 545
 546
 547 void init_gpu(const DeviceInformation* deviceInfo)
 548 {
 549     assert(deviceInfo);
 550
 551     // If the device is NVIDIA, for safety reasons we disable the JIT
 552     // caching as this is known to be broken at least until driver 364.19;
 553     // the cache does not always get regenerated when the source code changes,
 554     // e.g. if the path to the kernel sources remains the same
 555
 556     if (deviceInfo->deviceVendor == DeviceVendor::Nvidia)
 557     {
 558         // Ignore return values, failing to set the variable does not mean
 559         // that something will go wrong later.
 560 #ifdef _MSC_VER
 561         _putenv("CUDA_CACHE_DISABLE=1");
 562 #else
 563         // Don't override, maybe a dev is testing.
 564         setenv("CUDA_CACHE_DISABLE", "1", 0);
 565 #endif
 566     }
 567 }
 568
 569 DeviceInformation* getDeviceInfo(const gmx_gpu_info_t& gpu_info, int deviceId)
 570 {
 571     if (deviceId < 0 || deviceId >= gpu_info.n_dev)
 572     {
 573         gmx_incons("Invalid GPU deviceId requested");
 574     }
 575     return &gpu_info.deviceInfo[deviceId];
 576 }
 577
 578 size_t sizeof_gpu_dev_info()
 579 {
 580     return sizeof(DeviceInformation);
 581 }
 582
 583 DeviceStatus gpu_info_get_stat(const gmx_gpu_info_t& info, int index)
 584 {
 585     return info.deviceInfo[index].stat;
 586 }