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

Question

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

Я пробовал пример кода 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), работает нормально.

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

Answer 1

Связь между вашим вопросом о том, почему "результаты в основном нулевые" и 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 назначает одно адресное пространство для памяти CPU и GPU, так что копирование памяти между хостом и глобальной памятью нескольких устройств становится легко доступным за счет использованияcudaMemcpy с добрым cudaMemcpyDefault.

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

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