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 isDeviceSane(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 int 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 egpuCompatible;
 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 egpuIncompatible;
 234     }
 235
 236     /* Only AMD, Intel, and NVIDIA GPUs are supported for now */
 237     switch (deviceInfo->deviceVendor)
 238     {
 239         case DeviceVendor::Nvidia: return egpuCompatible;
 240         case DeviceVendor::Amd:
 241             return runningOnCompatibleOSForAmd() ? egpuCompatible : egpuIncompatible;
 242         case DeviceVendor::Intel:
 243             return GMX_OPENCL_NB_CLUSTER_SIZE == 4 ? egpuCompatible : egpuIncompatibleClusterSize;
 244         default: return egpuIncompatible;
 245     }
 246 }
 247
 248
 249 /*! \brief Check whether the \c ocl_gpu_device is suitable for use by mdrun
 250  *
 251  * Runs sanity checks: checking that the runtime can compile a dummy kernel
 252  * and this can be executed;
 253  * Runs compatibility checks verifying the device OpenCL version requirement
 254  * and vendor/OS support.
 255  *
 256  * \param[in]  deviceId      The runtime-reported numeric ID of the device.
 257  * \param[in]  deviceInfo    The device info pointer.
 258  * \returns  An e_gpu_detect_res_t to indicate how the GPU coped with
 259  *           the sanity and compatibility check.
 260  */
 261 static int checkGpu(size_t deviceId, const DeviceInformation* deviceInfo)
 262 {
 263
 264     int supportStatus = isDeviceSupported(deviceInfo);
 265     if (supportStatus != egpuCompatible)
 266     {
 267         return supportStatus;
 268     }
 269
 270     std::string errorMessage;
 271     if (!isDeviceSane(deviceInfo, &errorMessage))
 272     {
 273         gmx_warning("While sanity checking device #%zu, %s", deviceId, errorMessage.c_str());
 274         return egpuInsane;
 275     }
 276
 277     return egpuCompatible;
 278 }
 279
 280 } // namespace gmx
 281
 282 /*! \brief Returns an DeviceVendor value corresponding to the input OpenCL vendor name.
 283  *
 284  *  \param[in] vendorName  String with OpenCL vendor name.
 285  *  \returns               DeviceVendor value for the input vendor name
 286  */
 287 static DeviceVendor getDeviceVendor(const char* vendorName)
 288 {
 289     if (vendorName)
 290     {
 291         if (strstr(vendorName, "NVIDIA"))
 292         {
 293             return DeviceVendor::Nvidia;
 294         }
 295         else if (strstr(vendorName, "AMD") || strstr(vendorName, "Advanced Micro Devices"))
 296         {
 297             return DeviceVendor::Amd;
 298         }
 299         else if (strstr(vendorName, "Intel"))
 300         {
 301             return DeviceVendor::Intel;
 302         }
 303     }
 304     return DeviceVendor::Unknown;
 305 }
 306
 307 bool isGpuDetectionFunctional(std::string* errorMessage)
 308 {
 309     cl_uint numPlatforms;
 310     cl_int  status = clGetPlatformIDs(0, nullptr, &numPlatforms);
 311     GMX_ASSERT(status != CL_INVALID_VALUE, "Incorrect call of clGetPlatformIDs detected");
 312 #ifdef cl_khr_icd
 313     if (status == CL_PLATFORM_NOT_FOUND_KHR)
 314     {
 315         // No valid ICDs found
 316         if (errorMessage != nullptr)
 317         {
 318             errorMessage->assign("No valid OpenCL driver found");
 319         }
 320         return false;
 321     }
 322 #endif
 323     GMX_RELEASE_ASSERT(
 324             status == CL_SUCCESS,
 325             gmx::formatString("An unexpected value was returned from clGetPlatformIDs %d: %s",
 326                               status, ocl_get_error_string(status).c_str())
 327                     .c_str());
 328     bool foundPlatform = (numPlatforms > 0);
 329     if (!foundPlatform && errorMessage != nullptr)
 330     {
 331         errorMessage->assign("No OpenCL platforms found even though the driver was valid");
 332     }
 333     return foundPlatform;
 334 }
 335
 336 void findGpus(gmx_gpu_info_t* gpu_info)
 337 {
 338     cl_uint         ocl_platform_count;
 339     cl_platform_id* ocl_platform_ids;
 340     cl_device_type  req_dev_type = CL_DEVICE_TYPE_GPU;
 341
 342     ocl_platform_ids = nullptr;
 343
 344     if (getenv("GMX_OCL_FORCE_CPU") != nullptr)
 345     {
 346         req_dev_type = CL_DEVICE_TYPE_CPU;
 347     }
 348
 349     while (true)
 350     {
 351         cl_int status = clGetPlatformIDs(0, nullptr, &ocl_platform_count);
 352         if (CL_SUCCESS != status)
 353         {
 354             GMX_THROW(gmx::InternalError(
 355                     gmx::formatString("An unexpected value %d was returned from clGetPlatformIDs: ", status)
 356                     + ocl_get_error_string(status)));
 357         }
 358
 359         if (1 > ocl_platform_count)
 360         {
 361             // TODO this should have a descriptive error message that we only support one OpenCL platform
 362             break;
 363         }
 364
 365         snew(ocl_platform_ids, ocl_platform_count);
 366
 367         status = clGetPlatformIDs(ocl_platform_count, ocl_platform_ids, nullptr);
 368         if (CL_SUCCESS != status)
 369         {
 370             GMX_THROW(gmx::InternalError(
 371                     gmx::formatString("An unexpected value %d was returned from clGetPlatformIDs: ", status)
 372                     + ocl_get_error_string(status)));
 373         }
 374
 375         for (unsigned int i = 0; i < ocl_platform_count; i++)
 376         {
 377             cl_uint ocl_device_count;
 378
 379             /* If requesting req_dev_type devices fails, just go to the next platform */
 380             if (CL_SUCCESS != clGetDeviceIDs(ocl_platform_ids[i], req_dev_type, 0, nullptr, &ocl_device_count))
 381             {
 382                 continue;
 383             }
 384
 385             if (1 <= ocl_device_count)
 386             {
 387                 gpu_info->n_dev += ocl_device_count;
 388             }
 389         }
 390
 391         if (1 > gpu_info->n_dev)
 392         {
 393             break;
 394         }
 395
 396         snew(gpu_info->deviceInfo, gpu_info->n_dev);
 397
 398         {
 399             int           device_index;
 400             cl_device_id* ocl_device_ids;
 401
 402             snew(ocl_device_ids, gpu_info->n_dev);
 403             device_index = 0;
 404
 405             for (unsigned int i = 0; i < ocl_platform_count; i++)
 406             {
 407                 cl_uint ocl_device_count;
 408
 409                 /* If requesting req_dev_type devices fails, just go to the next platform */
 410                 if (CL_SUCCESS
 411                     != clGetDeviceIDs(ocl_platform_ids[i], req_dev_type, gpu_info->n_dev,
 412                                       ocl_device_ids, &ocl_device_count))
 413                 {
 414                     continue;
 415                 }
 416
 417                 if (1 > ocl_device_count)
 418                 {
 419                     break;
 420                 }
 421
 422                 for (unsigned int j = 0; j < ocl_device_count; j++)
 423                 {
 424                     gpu_info->deviceInfo[device_index].oclPlatformId = ocl_platform_ids[i];
 425                     gpu_info->deviceInfo[device_index].oclDeviceId   = ocl_device_ids[j];
 426
 427                     gpu_info->deviceInfo[device_index].device_name[0] = 0;
 428                     clGetDeviceInfo(ocl_device_ids[j], CL_DEVICE_NAME,
 429                                     sizeof(gpu_info->deviceInfo[device_index].device_name),
 430                                     gpu_info->deviceInfo[device_index].device_name, nullptr);
 431
 432                     gpu_info->deviceInfo[device_index].device_version[0] = 0;
 433                     clGetDeviceInfo(ocl_device_ids[j], CL_DEVICE_VERSION,
 434                                     sizeof(gpu_info->deviceInfo[device_index].device_version),
 435                                     gpu_info->deviceInfo[device_index].device_version, nullptr);
 436
 437                     gpu_info->deviceInfo[device_index].vendorName[0] = 0;
 438                     clGetDeviceInfo(ocl_device_ids[j], CL_DEVICE_VENDOR,
 439                                     sizeof(gpu_info->deviceInfo[device_index].vendorName),
 440                                     gpu_info->deviceInfo[device_index].vendorName, nullptr);
 441
 442                     gpu_info->deviceInfo[device_index].compute_units = 0;
 443                     clGetDeviceInfo(ocl_device_ids[j], CL_DEVICE_MAX_COMPUTE_UNITS,
 444                                     sizeof(gpu_info->deviceInfo[device_index].compute_units),
 445                                     &(gpu_info->deviceInfo[device_index].compute_units), nullptr);
 446
 447                     gpu_info->deviceInfo[device_index].adress_bits = 0;
 448                     clGetDeviceInfo(ocl_device_ids[j], CL_DEVICE_ADDRESS_BITS,
 449                                     sizeof(gpu_info->deviceInfo[device_index].adress_bits),
 450                                     &(gpu_info->deviceInfo[device_index].adress_bits), nullptr);
 451
 452                     gpu_info->deviceInfo[device_index].deviceVendor =
 453                             getDeviceVendor(gpu_info->deviceInfo[device_index].vendorName);
 454
 455                     clGetDeviceInfo(ocl_device_ids[j], CL_DEVICE_MAX_WORK_ITEM_SIZES, 3 * sizeof(size_t),
 456                                     &gpu_info->deviceInfo[device_index].maxWorkItemSizes, nullptr);
 457
 458                     clGetDeviceInfo(ocl_device_ids[j], CL_DEVICE_MAX_WORK_GROUP_SIZE, sizeof(size_t),
 459                                     &gpu_info->deviceInfo[device_index].maxWorkGroupSize, nullptr);
 460
 461                     gpu_info->deviceInfo[device_index].stat =
 462                             gmx::checkGpu(device_index, gpu_info->deviceInfo + device_index);
 463
 464                     if (egpuCompatible == gpu_info->deviceInfo[device_index].stat)
 465                     {
 466                         gpu_info->n_dev_compatible++;
 467                     }
 468
 469                     device_index++;
 470                 }
 471             }
 472
 473             gpu_info->n_dev = device_index;
 474
 475             /* Dummy sort of devices -  AMD first, then NVIDIA, then Intel */
 476             // TODO: Sort devices based on performance.
 477             if (0 < gpu_info->n_dev)
 478             {
 479                 int last = -1;
 480                 for (int i = 0; i < gpu_info->n_dev; i++)
 481                 {
 482                     if (gpu_info->deviceInfo[i].deviceVendor == DeviceVendor::Amd)
 483                     {
 484                         last++;
 485
 486                         if (last < i)
 487                         {
 488                             std::swap(gpu_info->deviceInfo[i], gpu_info->deviceInfo[last]);
 489                         }
 490                     }
 491                 }
 492
 493                 /* if more than 1 device left to be sorted */
 494                 if ((gpu_info->n_dev - 1 - last) > 1)
 495                 {
 496                     for (int i = 0; i < gpu_info->n_dev; i++)
 497                     {
 498                         if (gpu_info->deviceInfo[i].deviceVendor == DeviceVendor::Nvidia)
 499                         {
 500                             last++;
 501
 502                             if (last < i)
 503                             {
 504                                 std::swap(gpu_info->deviceInfo[i], gpu_info->deviceInfo[last]);
 505                             }
 506                         }
 507                     }
 508                 }
 509             }
 510
 511             sfree(ocl_device_ids);
 512         }
 513
 514         break;
 515     }
 516
 517     sfree(ocl_platform_ids);
 518 }
 519
 520 void get_gpu_device_info_string(char* s, const gmx_gpu_info_t& gpu_info, int index)
 521 {
 522     assert(s);
 523
 524     if (index < 0 && index >= gpu_info.n_dev)
 525     {
 526         return;
 527     }
 528
 529     DeviceInformation* dinfo = &gpu_info.deviceInfo[index];
 530
 531     bool bGpuExists = (dinfo->stat != egpuNonexistent && dinfo->stat != egpuInsane);
 532
 533     if (!bGpuExists)
 534     {
 535         sprintf(s, "#%d: %s, stat: %s", index, "N/A", gpu_detect_res_str[dinfo->stat]);
 536     }
 537     else
 538     {
 539         sprintf(s, "#%d: name: %s, vendor: %s, device version: %s, stat: %s", index, dinfo->device_name,
 540                 dinfo->vendorName, dinfo->device_version, gpu_detect_res_str[dinfo->stat]);
 541     }
 542 }
 543
 544
 545 void init_gpu(const DeviceInformation* deviceInfo)
 546 {
 547     assert(deviceInfo);
 548
 549     // If the device is NVIDIA, for safety reasons we disable the JIT
 550     // caching as this is known to be broken at least until driver 364.19;
 551     // the cache does not always get regenerated when the source code changes,
 552     // e.g. if the path to the kernel sources remains the same
 553
 554     if (deviceInfo->deviceVendor == DeviceVendor::Nvidia)
 555     {
 556         // Ignore return values, failing to set the variable does not mean
 557         // that something will go wrong later.
 558 #ifdef _MSC_VER
 559         _putenv("CUDA_CACHE_DISABLE=1");
 560 #else
 561         // Don't override, maybe a dev is testing.
 562         setenv("CUDA_CACHE_DISABLE", "1", 0);
 563 #endif
 564     }
 565 }
 566
 567 DeviceInformation* getDeviceInfo(const gmx_gpu_info_t& gpu_info, int deviceId)
 568 {
 569     if (deviceId < 0 || deviceId >= gpu_info.n_dev)
 570     {
 571         gmx_incons("Invalid GPU deviceId requested");
 572     }
 573     return &gpu_info.deviceInfo[deviceId];
 574 }
 575
 576 size_t sizeof_gpu_dev_info()
 577 {
 578     return sizeof(DeviceInformation);
 579 }
 580
 581 int gpu_info_get_stat(const gmx_gpu_info_t& info, int index)
 582 {
 583     return info.deviceInfo[index].stat;
 584 }