Код OpenCL N-body с двумя машинами GPU NVIDIA A6000 (с подключением NVLink между собой)

Question

Я бы хотел запустить старое N-тело, использующее OpenCL.

У меня 2 карты NVIDIA A6000 с NVLink, компонент, который связывает с аппаратной (и, возможно, программной?) точки зрения эти 2 карты GPU.

Но при исполнении получаю следующий результат:

ядро отказало

Вот используемый код ядра (я добавил прагму, которая, по моему мнению, полезна для карт NVIDIA):

      #pragma OPENCL EXTENSION cl_khr_fp64 : enable

__kernel
void
nbody_sim(
    __global double4* pos ,
    __global double4* vel,
    int numBodies,
    double deltaTime,
    double epsSqr,
    __local double4* localPos,
    __global double4* newPosition,
    __global double4* newVelocity)
{
    unsigned int tid = get_local_id(0);
    unsigned int gid = get_global_id(0);
    unsigned int localSize = get_local_size(0);

    // Gravitational constant
    double G_constant = 227.17085e-74;

    // Number of tiles we need to iterate
    unsigned int numTiles = numBodies / localSize;

    // position of this work-item
    double4 myPos = pos[gid];
    double4 acc = (double4) (0.0f, 0.0f, 0.0f, 0.0f);

    for(int i = 0; i < numTiles; ++i)
    {
        // load one tile into local memory
        int idx = i * localSize + tid;
        localPos[tid] = pos[idx];

        // Synchronize to make sure data is available for processing
        barrier(CLK_LOCAL_MEM_FENCE);

        // Calculate acceleration effect due to each body
        // a[i->j] = m[j] * r[i->j] / (r^2 + epsSqr)^(3/2)
        for(int j = 0; j < localSize; ++j)
        {
            // Calculate acceleration caused by particle j on particle i
            double4 r = localPos[j] - myPos;
            double distSqr = r.x * r.x  +  r.y * r.y  +  r.z * r.z;
            double invDist = 1.0f / sqrt(distSqr + epsSqr);
            double invDistCube = invDist * invDist * invDist;
            double s = G_constant * localPos[j].w * invDistCube;

            // accumulate effect of all particles
            acc += s * r;
        }

        // Synchronize so that next tile can be loaded
        barrier(CLK_LOCAL_MEM_FENCE);
    }

    double4 oldVel = vel[gid];

    // updated position and velocity
    double4 newPos = myPos + oldVel * deltaTime + acc * 0.5f * deltaTime * deltaTime;
    newPos.w = myPos.w;
    double4 newVel = oldVel + acc * deltaTime;

    // write to global memory
    newPosition[gid] = newPos;
    newVelocity[gid] = newVel;
}

Часть кода, которая устанавливает код ядра, приведена ниже:

      int NBody::setupCL()
{
  cl_int status = CL_SUCCESS;
  cl_event writeEvt1, writeEvt2;

  // The block is to move the declaration of prop closer to its use
  cl_command_queue_properties prop = 0;
  commandQueue = clCreateCommandQueue(
      context,
      devices[current_device],
      prop,
      &status);
  CHECK_OPENCL_ERROR( status, "clCreateCommandQueue failed.");

    ...

// create a CL program using the kernel source
  const char *kernelName = "NBody_Kernels.cl";
  FILE *fp = fopen(kernelName, "r");
  if (!fp) {
    fprintf(stderr, "Failed to load kernel.\n");
    exit(1);
  }
  char *source = (char*)malloc(10000);
  int sourceSize = fread( source, 1, 10000, fp);
  fclose(fp);

  // Create a program from the kernel source
  program = clCreateProgramWithSource(context, 1, (const char **)&source, (const size_t *)&sourceSize, &status);

  // Build the program
  status = clBuildProgram(program, 1, devices, NULL, NULL, NULL);

  // get a kernel object handle for a kernel with the given name
  kernel = clCreateKernel(
      program,
      "nbody_sim",
      &status);
  CHECK_OPENCL_ERROR(status, "clCreateKernel failed.");

  status = waitForEventAndRelease(&writeEvt1);
  CHECK_ERROR(status, NBODY_SUCCESS, "WaitForEventAndRelease(writeEvt1) Failed");

  status = waitForEventAndRelease(&writeEvt2);
  CHECK_ERROR(status, NBODY_SUCCESS, "WaitForEventAndRelease(writeEvt2) Failed");

  return NBODY_SUCCESS;
}

Итак, ошибка возникает при создании кода ядра . Есть ли способ рассмотреть the 2 GPU как уникальный графический процессор с NVLINK component? Я имею в виду с точки зрения программного обеспечения?

Как исправить эту ошибку создания кода ядра ?

Answer 1

Код вашего ядра выглядит хорошо, и реализация мозаики кеша правильная. Только убедитесь, что количество тел кратно локальному размеру, или, в качестве альтернативы, дополнительно ограничьте внутренний цикл for глобальным размером.

OpenCL позволяет использовать несколько устройств параллельно. Вам нужно сделать поток с очередью для каждого устройства отдельно . Вам также необходимо позаботиться о связи и синхронизации между устройствами вручную. Передача данных происходит через PCIe (вы также можете осуществлять удаленный прямой доступ к памяти); но вы не можете использовать NVLink с OpenCL. Однако в вашем случае это не должно быть проблемой, поскольку вам потребуется лишь небольшая передача данных по сравнению с объемом арифметических операций.

Еще несколько замечаний:

• Во многих случаях N-body требует FP64 для суммирования сил и определения положений в очень разных масштабах длины. Однако на A6000 производительность FP64 очень низкая, как и на GeForce Ampere. FP32 будет значительно (~64x) быстрее, но, вероятно, здесь недостаточно с точки зрения точности. Для эффективного FP64 вам понадобится A100 или MI100.
• Вместо 1.0/sqrt используйте rsqrt. Это аппаратно поддерживается и почти так же быстро, как умножение.
• обязательно используйте литералы FP32 float (1.0f) или FP64 double (1.0) последовательно. Использование двойных литералов с float запускает двойную арифметику и приведение результата обратно к float, что намного медленнее.