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