Недопустимая ошибка символа устройства в зависимости от флагов nvcc

Question

Игрушечная программа:

#include <iostream>
#include <vector>

// Matrix side size (they are square).
const int N = 3;
const int num_mats = 14;

// Rotation matrices.
__constant__ float rot_mats_device[num_mats*N*N];

int main() {
  std::vector<float> rot_mats_host(num_mats*N*N);
  for (int i = 0; i < rot_mats_host.size(); i++)
    rot_mats_host[i] = i;

  auto errMemcpyToSymbol = cudaMemcpyToSymbol(rot_mats_device,
                                              rot_mats_host.data(),
                                              sizeof(rot_mats_device));

  if (errMemcpyToSymbol != cudaSuccess) {
    std::cout << "MemcpyToSymbol error: " <<
      cudaGetErrorString(errMemcpyToSymbol) << std::endl;
  }
}

Составлено с

nvcc -arch=sm_52 -std=c++11 cuda_invalid_symbol_error.cu -o cuda_invalid_symbol_error

не дает никаких ошибок во время выполнения. Однако с

nvcc -gencode arch=compute_52,code=sm_52 -std=c++11 cuda_invalid_symbol_error.cu -o cuda_invalid_symbol_error

это не удастся с сообщением MemcpyToSymbol error: invalid device symbol,

Почему последние инструкции для компиляции дают ошибку времени выполнения?

Спецификации: cuda 8.0, Ubuntu 16.04, GeForce GTX 1060 (я знаю, что копия этой карты - 6.1).

Answer 1

Почему последние инструкции для компиляции дают ошибку времени выполнения?

-arch=sm_xx является сокращением для:

-gencode arch=compute_xx,code=sm_xx -gencode arch=compute_xx,code=compute_xx

В вашем случае, где xx 52, эта команда встраивает код PTC cc 5.2 (вторая gencode экземпляр) и код SASS 5.2 cc (первый gencode пример). Код SASS для cc 5.2 не будет работать на вашем устройстве cc6.1, поэтому среда выполнения JIT компилирует код PTC cc 5.2 для создания объекта, совместимого с вашей архитектурой cc 6.1. Все счастливы и все работает.

Когда вы вместо этого скомпилируете с:

nvcc -gencode arch=compute_52,code=sm_52 ...

вы пропускаете код PTX из скомпилированного объекта. Присутствует только код SAC 5.2 cc. Этот код не будет работать на вашем устройстве cc6.1, и у среды выполнения нет других опций, поэтому "скрытая" ошибка NO_BINARY_FOR_GPU возникает, когда среда выполнения пытается загрузить образ графического процессора для вашей программы. Поскольку изображение не загружается, символ устройства отсутствует / может использоваться. Поскольку его нет / можно использовать, вы получаете invalid device symbol ошибка при попытке обратиться к нему с помощью API времени выполнения CUDA.

Если до этого вы выполняли еще один вызов API среды выполнения CUDA, который вызвал достаточный или эквивалентный уровень инициализации среды выполнения CUDA (и проверил возвращенный код ошибки), вы бы получили ошибку NO_BINARY_FOR_GPU или аналогичную. Конечно, например, если вы попытались запустить ядро ​​графического процессора, вы получите эту ошибку. Могут быть другие вызовы API CUDA времени выполнения, которые вызовут эквивалентный или достаточный уровень отложенной инициализации, но у меня нет их списка.