Доступ к памяти CUDA P2P и __constant__ памяти - PullRequest
Новая
       36

Доступ к памяти CUDA P2P и __constant__ памяти

0 голосов
/ 09 марта 2020

Я нигде не могу найти ответ, и, возможно, я его упустил, но кажется, что нельзя использовать __constant__ память (наряду с cudaMemcpyToSymbol) и одноранговый доступ с помощью UVA.

I Я пробовал пример кода SimpleP2P NVIDIA, который отлично работает на 4 NV100 с nvlink, который у меня есть, но пока я объявляю фактор 2 в ядре как: 

__constant__ float M_; // in global space

float M = 2.0;
cudaMemcpyToSymbol(M_, &M, sizeof(float), 0, cudaMemcpyDefault);

результаты в основном равны нулю. Если я определяю его с помощью препроцессора C (например, #define M_ 2.0), он работает нормально.

Так что мне интересно, правда ли это или я что-то не так делаю? и есть ли другие виды памяти, к которым также нельзя получить доступ таким образом (например, память текстур)?

1 Ответ

1 голос
/ 10 марта 2020

Отношение между вашим вопросом о том, почему "результаты в основном равны нулю" и P2P-доступом с помощью UVA, мне не совсем понятно.

это правда или я что-то не так делаю?

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

__constant__ float M_ выделяет переменную M_ в постоянной памяти из всех видимых устройств CUDA. Чтобы установить значение на нескольких устройствах, вы должны сделать что-то вроде: 

__constant__ float M_; // <= This declares M_ on the constant memory of all CUDA visible devices

__global__ void showMKernel() {
    printf("****** M_ = %f\n", M_);
}

int main()
{

float M = 2.0;

 // Make sure that the return values are properly checked for cudaSuccess ...

int deviceCount = -1;
cudaGetDeviceCount(&deviceCount);

// Set M_ on the constant memory of each device:
for (int i = 0; i < deviceCount; i++) {
  cudaSetDevice(i);
  cudaMemcpyToSymbol(M_, &M, sizeof(float), 0, cudaMemcpyDefault);
}

// Now, run a kernel to show M_:
for (int i = 0; i < deviceCount; i++) 
{
  cudaSetDevice(i);
  printf("Device %g :\n", i);
  showMKernel<<<1,1>>>();
  cudaDeviceSynchronize();
}

}

, который возвращает: 

Device 0 :
****** M = 2.000000
Device 1 :
****** M = 2.000000
// so on for other devices

Теперь, если я заменю 

// Set M_ on the constant memory of each device:
for (int i = 0; i < deviceCount; i++) {
  cudaSetDevice(i);
  cudaMemcpyToSymbol(M_, &M, sizeof(float), 0, cudaMemcpyDefault);
}

с

cudaMemcpyToSymbol(M_, &M, sizeof(float), 0, cudaMemcpyDefault);

это только установит значение M_ на активном устройстве и поэтому возвращает 

Device 0 :
****** M = 2.000000
Device 1 :
****** M = 0.000000 // <= I assume this is what you meant by 'the results are basically zero'
// M = 0 for other devices too

, есть ли другие виды памяти, к которой также нельзя получить доступ таким образом (например, к текстурной памяти)?

Опять же, я не совсем уверен, что такое таким образом . Я думаю, что в общем случае вы не можете получить доступ к постоянной памяти или памяти текстур одного устройства с других устройств, хотя я не уверен на 100%.

UVA назначает одно адресное пространство для памяти ЦП и ГП так, чтобы копирование памяти между хост и глобальная память нескольких устройств становятся легко доступными благодаря использованию cudaMemcpy с типом cudaMemcpyDefault.

Кроме того, P2P-связь между устройствами обеспечивает прямой доступ и передачу данных между глобальной памятью нескольких устройств.

Аналогично приведенному выше примеру __constant__, когда вы объявляете текстуру, такую ​​как texture <float> some_texture, some_texture будет определяться для каждого видимого устройства, однако вам потребуется явно привязать some_texture к вашей текстурной ссылке на каждом устройстве при работе с несколькими устройствами.

...