Вопрос: Как и когда следует использовать передаваемый указатель с API-интерфейсом cuda?


Я достаточно хорошо понимаю, как распределять и копировать линейную память с помощью cudaMalloc() а также cudaMemcpy(), Однако, когда я хочу использовать функции CUDA для выделения и копирования 2D или 3D-матриц, меня часто опровергают различные аргументы, особенно касающиеся разнесенных указателей, которые всегда присутствуют при работе с 2D / 3D-массивами. Документация хороша для того, чтобы предоставить пару примеров того, как их использовать, но предполагает, что я знаком с понятием заполнения и подачи, которого я не знаю.

Я обычно в конечном итоге настраиваю различные примеры, которые я нахожу в документации или где-то еще в Интернете, но слепая отладка, которая следует за этим, довольно болезненна, поэтому мой вопрос:

Что такое подача? Как его использовать? Как распределить и скопировать 2D-и 3D-массивы в CUDA?


37


источник


Ответы:


Ниже приведено объяснение, касающееся указателя и отступов в кудах.

Линейная память и запатентованная память

Во-первых, давайте начнем с причины существования нелинейной памяти. При распределении памяти с помощью cudaMalloc результат подобен распределению с помощью malloc, у нас есть смежный кусок памяти указанного размера, и мы можем поместить в него все, что мы хотим. Если мы хотим выделить вектор с плавающей запятой 10000, мы просто делаем:

float* myVector;
cudaMalloc(&myVector,10000*sizeof(float));

а затем доступ к i-му элементу myVector путем классического индексирования:

float element = myVector[i];

и если мы хотим получить доступ к следующему элементу, мы просто делаем:

float next_element = myvector[i+1];

Он работает очень хорошо, потому что доступ к элементу прямо рядом с первым (по причинам, о которых я не знаю, и я не хочу быть на данный момент) дешево.

Когда мы используем нашу память как 2D-массив, ситуация становится немного иной. Допустим, что наш векторный флоп 10000 фактически представляет собой массив 100x100. Мы можем выделить его, используя ту же функцию cudaMalloc, и если мы хотим прочитать i-ю строку, мы делаем:

float* myArray;
cudaMalloc(&myArray,10000*sizeof(float));
int row[100]; // number of columns
for (int j=0; j<100; ++j)
    row[j] = myArray[i*100+j];

Выравнивание слов

Поэтому мы должны читать память от myArray + 100 * i до myArray + 101 * i-1. Количество выполняемых операций доступа к памяти зависит от количества слов памяти, которые занимает эта строка. Количество байтов в слове памяти зависит от реализации. Чтобы свести к минимуму количество обращений к памяти при чтении одной строки, мы должны заверить, что мы начинаем строку в начале слова, поэтому мы должны заполнять память для каждой строки до начала новой.

Банковский конфликт

Еще одна причина для массивов дополнений - это банковская меканизм в cuda, касающаяся доступа к общей памяти. Когда массив находится в общей памяти, он разделяется на несколько банков памяти. Два потока cuda могут получить к нему доступ одновременно, при условии, что они не получают доступ к памяти, принадлежащей одному банку памяти. Поскольку мы обычно хотим рассматривать каждую строку параллельно, мы можем гарантировать, что мы можем получить доступ к ней с помощью имитации, заполнив каждую строку до начала нового банка.

Теперь вместо выделения двумерного массива cudaMalloc мы будем использовать cudaMallocPitched:

size_t pitch;
float* myArray;
cudaMallocPitch(&myArray,&pitch,100*sizeof(float),100);//width in bytes by height

Обратите внимание, что шаг здесь - это возвращаемое значение функции: cudaMallocPitch проверяет, что должно быть в вашей системе, и возвращает соответствующее значение. Что делает cudaMallocPitch:

  1. Выделите первую строку.
  2. Проверьте, правильно ли выровнено количество выделенных байтов ( т.е.  он кратен 128).
  3. Если нет, выделите дополнительные байты, чтобы достигнуть следующего кратного 128. шаг - это количество байтов, выделенных для одной строки, включая дополнительные байты (байты заполнения).
  4. Повторите для каждой строки.

В конце мы обычно выделяли больше памяти, чем необходимо, потому что каждая строка теперь имеет размер шага, а не размер w * sizeof (float).

Но теперь, когда мы хотим получить доступ к следующему элементу в столбце, мы должны сделать следующее:

float next_column_element = myArray[(j+1)*pitch+i];

Смещение в байтах между двумя последовательными столбцами больше не может быть выведено из размера нашего массива, поэтому мы хотим сохранить высоту, возвращаемую cudaMallocPitch. И так как шаг является кратным размеру заполнения (как правило, самый большой размер слова и размер банка), он отлично работает. Ура.

Копирование данных в / из скатной памяти

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

Допустим, мы хотим скопировать наш массив в массив 100x100, выделенный на нашем хосте с помощью malloc:

float* host_memory = (float*)malloc(100*100*sizeof(float));

Если мы используем cudaMemcpy, мы скопируем всю память, выделенную cudaMallocPitch, включая заполненные байты между каждой строкой. То, что мы должны сделать, чтобы избежать заполнения памяти, - это копировать каждую строку по одной. Мы можем сделать это вручную:

for (size_t i=0;i<100;++i) {
cudaMemcpy(host_memory[i*100],myArray[pitch*i],
    100*sizeof(float),cudaMemcpyDeviceToHost);
}

Или мы можем сказать API-интерфейсу cuda, что мы хотим получить только полезную память из памяти, которую мы выделили с байтами заполнения для его  удобство, поэтому, если бы он мог справиться с собственным беспорядком автоматически, было бы очень приятно, спасибо. И здесь входит cudaMemcpy2D:

cudaMemcpy2D(host_memory,100*sizeof(float)/*destination pitch*/,myArray,pitch,
100*sizeof(float)/*width*/,100/*heigth*/,cudaMemcpyDeviceToHost);

Теперь копия будет выполнена автоматически. Он скопирует количество байтов, заданных по ширине (здесь: 100xsizeof (float)), время ожидания (здесь: 100), пропуск подача  байты каждый раз, когда он переходит к следующей строке. Обратите внимание, что мы все равно должны предоставить высоту тона для памяти назначения, так как она также может быть дополнена. Здесь это не так, поэтому тангаж равен высоте неделаемого массива: это размер строки. Также обратите внимание, что параметр width в функции memcpy выражается в байтах, но параметр heigth выражается в количестве элементов. Это происходит из-за способа копирования, каким-то образом, как я написал копию вручную выше: ширина - это размер каждой копии вдоль строки (элементы, которые смежны в памяти), а высота - это количество раз, которое эта операция должна выполняются. (Эти несоответствия в единицах, как физик, меня очень раздражает).

Работа с 3D-массивами

3D-массивы ничем не отличаются от того, что 2D-массивы на самом деле, нет дополнительного дополнения. 3D-массив - это всего лишь 2D классический  массив заполненных строк. Вот почему при распределении 3D-массива вы получаете только один шаг, который представляет собой разницу между количеством байтов между последовательными точками вдоль строки. Если вы хотите получить доступ к последовательным точкам по размеру глубины, вы можете безопасно умножить высоту тона на количество столбцов, что даст вам slicePitch.

Cuda api для доступа к 3D-памяти слабее, чем у 2D-памяти, но идея такая же:

  • При использовании cudaMalloc3D вы получаете значение тангажа, которое вы должны тщательно сохранить для последующего доступа к памяти.
  • При копировании фрагмента 3D-памяти вы не можете использовать cudaMemcpy, если вы не копируете одну строку. Вы должны использовать любой другой тип utlity, предоставляемый утилитой cuda, который учитывает поле.
  • Когда вы копируете свои данные в / из линейной памяти, вы должны указать шаг указателю, даже если это не имеет значения: этот шаг - это размер строки, выраженный в байтах.
  • Параметры размера выражаются в байтах для размера строки, а также в количестве элементов для измерения столбца и глубины.

70



В ответе Гефеста

Если мы используем cudaMemcpy, мы скопируем всю память, выделенную cudaMallocPitch, включая заполненные байты между каждой строкой. То, что мы должны сделать, чтобы избежать заполнения памяти, - это копировать каждую строку по одной. Мы можем сделать это вручную:

for (size_t i=0;i<100;++i) {
cudaMemcpy(host_memory[i*100],myArray[pitch*100],
    100*sizeof(float),cudaMemcpyDeviceToHost);
}

Здесь адрес «Исходная память» должен быть myArray[i*pitch] скорее, чем myArray[pitch*100],


2



В ответе Эрнеста Галбруна

float next_column_element = myArray[(j+1)*pitch+i];

должны быть

float next_column_element = *((float*)((char*)myArray + (j+1) * pitch) + i);

как в http://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__MEMORY.html#group__CUDART__MEMORY_1g32bd7a39135594788a542ae72217775c

И как указал @RobertCrovella,

float next_column_element = myArray[(j+1)*pitch/sizeof(float)+i];

это тоже не правильный путь.


0