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