Ответ 1

Вот объяснение о наклонном указателе и заполнении в CUDA.

Линейная память против мягкой памяти

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

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

и затем получить доступ к i-му элементу myVector с помощью классической индексации:

float element = myVector[i];

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

float next_element = myvector[i+1];

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

Все становится немного по-другому, когда мы используем нашу память в качестве двумерного массива. Допустим, наш вектор с плавающей точкой 10000 на самом деле является массивом 100х100. Мы можем выделить его, используя ту же функцию 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 * я в myArray + 101 * i-1. Число операций доступа к памяти зависит от количества слов памяти, которые занимает эта строка. Количество байтов в слове памяти зависит от реализации. Чтобы минимизировать количество обращений к памяти при чтении одной строки, мы должны убедиться, что мы начинаем строку в начале слова, следовательно, мы должны заполнять память для каждой строки до начала новой.

Банковские конфликты

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

Теперь вместо выделения 2D-массива с помощью cudaMalloc, мы будем использовать cudaMallocPitch:

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

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

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

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

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

float* row_start = (float*)((char*)myArray + row * pitch);
float column_element = row_start[column];

Смещение в байтах между двумя последовательными столбцами больше не может быть выведено из размера нашего массива, поэтому мы хотим сохранить высоту, возвращаемую 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)/*no pitch on host*/,
             myArray, pitch/*CUDA pitch*/,
             100*sizeof(float)/*width in bytes*/, 100/*heigth*/, 
             cudaMemcpyDeviceToHost);

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

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

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

API CUDA для доступа к 3D-памяти немного отличается от API для 2D-памяти, но идея та же:

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

Ответ 2

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

Если мы используем 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].

Ответ 3

В сообщении Ernest_Galbrun

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];

не является правильным способом.