В ядре CUDA, как мне хранить массив в "локальной памяти потока"?

Ответ 1

Массивы, локальная память и регистры

Здесь есть неправильное представление об определении "локальной памяти". "Локальная память" в CUDA на самом деле является глобальной памятью (и на самом деле ее следует называть "глобальной локальной памятью потока") с чередованной адресацией (что делает итерацию по массиву параллельно немного быстрее, чем блокировка данных каждого потока). Если вы хотите, чтобы все было действительно быстро, вы хотите использовать либо разделяемую память, либо, что еще лучше, регистры (особенно на последних устройствах, где вы получаете до 255 регистров на поток). Объяснение всей иерархии памяти CUDA выходит за рамки этого поста. Вместо этого давайте сконцентрируемся на быстром вычислении небольших массивов.

Небольшие массивы, как и переменные, могут храниться целиком в регистрах. Однако на современном оборудовании NVIDIA размещение массивов в регистрах затруднительно. Зачем? Потому что регистры требуют очень бережного отношения. Если вы сделаете это не совсем правильно, ваши данные окажутся в локальной памяти (которая, опять же, действительно является глобальной памятью, которая является самой медленной из имеющихся у вас). Руководство по программированию CUDA, раздел 5.3.2, сообщает вам, когда используется локальная память:

Локальная память

Доступ к локальной памяти происходит только для некоторых автоматических переменных, как указано в Спецификаторах типов переменных. Автоматические переменные, которые компилятор может поместить в локальную память:

1. Массивы, для которых он не может определить, что они проиндексированы с постоянными величинами,
2. Большие структуры или массивы, которые будут занимать слишком много места в регистре,
3. Любая переменная, если ядро использует больше регистров, чем доступно (это также называется проливом регистров).

Как работает распределение регистра?

Обратите внимание, что распределение регистров - чрезвычайно сложный процесс, поэтому вы не можете (и не должны) вмешиваться в него. Вместо этого компилятор преобразует код CUDA в код PTX (своего рода байт-код), который предполагает наличие машины с бесконечным числом регистров. Вы можете написать встроенный PTX, но это не сделает слишком много, чтобы зарегистрировать распределение. PTX-код - это независимый от устройства код, и это только первый этап. На втором этапе PTX будет скомпилирован в код сборки устройства, называемый SASS. Код SASS имеет фактическое распределение регистров. Компилятор SASS и его оптимизатор также будут в конечном итоге определять, будет ли переменная находиться в регистрах или в локальной памяти. Все, что вы можете сделать, это попытаться понять, что делает компилятор SASS в определенных случаях, и использовать это в своих интересах. В этом может помочь просмотр корреляции кода в Nsight (см. Ниже). Однако, поскольку компилятор и оптимизатор продолжают изменяться, нет никаких гарантий относительно того, что будет или не будет в регистрах.

Недостаточно регистров

Приложение G, раздел 1, сообщает, сколько регистров может иметь поток. Ищите "Максимальное количество 32-битных регистров на поток". Чтобы интерпретировать эту таблицу, вы должны знать свои вычислительные возможности (см. Ниже). Не забывайте, что регистры используются для всех видов вещей, и не просто соотносятся с отдельными переменными. Регистры на всех устройствах до CC 3.5 являются 32-битными каждый. Если компилятор достаточно умен (и компилятор CUDA продолжает изменяться), он может, например, упаковать несколько байтов в один и тот же регистр. Представление корреляции кода Nsight (см. "Анализ доступа к памяти" ниже) также показывает это.

Постоянная и динамическая индексация

Несмотря на то, что ограничение пространства является очевидным препятствием для массивов в реестре, легко отслеживается тот факт, что на текущем оборудовании (Compute Capability 3.x и ниже) компилятор помещает любой массив в локальную память, доступ к которому осуществляется с помощью динамическая индексация. Динамический индекс - это индекс, который компилятор не может понять. Массивы, доступ к которым осуществляется с помощью динамических индексов, не могут быть помещены в регистры, поскольку регистры должны определяться компилятором, и, следовательно, фактический используемый регистр не должен зависеть от значения, определенного во время выполнения. Например, для массива arr arr[k] является индексированием констант тогда и только тогда, когда k является константой или зависит только от констант. Если k каким-либо образом зависит от некоторого непостоянного значения, компилятор не может вычислить значение k и вы получили динамическое индексирование. В циклах, где k начинается и заканчивается с (маленькими) постоянными числами, компилятор (наиболее вероятно) может развернуть ваш цикл и все еще может достичь постоянной индексации.

пример

Например, сортировка небольшого массива может быть выполнена в регистрах, но вы должны использовать сортировку сетей или аналогичные "аппаратные" подходы, и не можете просто использовать стандартный алгоритм, потому что большинство алгоритмов используют динамическую индексацию.

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

__global__
void
testSortingNetwork4(const char * aInput, char * aResult)
{
    const int NBytes = 4;

    char aBytes[NBytes];

    // copy input to local array
    for (int i = 0; i < NBytes; ++i)
    {
        aBytes[i] = aInput[i];
    }

    // sort using sorting network
    CompareAndSwap(aBytes, 0, 2); CompareAndSwap(aBytes, 1, 3); 
    CompareAndSwap(aBytes, 0, 1); CompareAndSwap(aBytes, 2, 3); 
    CompareAndSwap(aBytes, 1, 2); 


    // copy back to result array
    for (int i = 0; i < NBytes; ++i)
    {
        aResult[i] = aBytes[i];
    }
}

Анализ доступа к памяти

Как только вы закончите, вы, как правило, хотите проверить, действительно ли данные хранятся в регистрах или они поступили в локальную память. Первое, что вы можете сделать, это указать вашему компилятору предоставить вам статистику памяти, используя --ptxas-options=-v. Более подробный способ анализа обращений к памяти - использование Nsight.

Nsight имеет много интересных функций. Nsight для Visual Studio имеет встроенный профилировщик и представление корреляции кода CUDA <-> SASS. Функция объясняется здесь. Обратите внимание, что версии Nsight для разных IDE, вероятно, разрабатываются независимо, и, следовательно, их функции могут различаться в разных реализациях.

Если вы будете следовать инструкциям в приведенной выше ссылке (не забудьте добавить соответствующие флаги при компиляции!), Вы можете найти кнопку "Операции с памятью CUDA" в самом низу нижнего меню. В этом представлении вы хотите обнаружить, что нет транзакции памяти, исходящей из строк, которые работают только с соответствующим массивом (например, строки CompareAndSwap в моем примере кода). Потому что, если он не сообщает о доступе к памяти для этих строк, вы (очень вероятно) смогли бы сохранить все вычисления в регистрах и могли бы просто получить ускорение в тысячи, если не в десятки тысяч процентов (вы также можете захотеть проверьте фактическое увеличение скорости, вы выходите из этого!).

Вычисление вычислительных возможностей

Чтобы выяснить, сколько регистров у вас есть, вам нужно знать возможности вашего устройства для вычислений. Стандартный способ получения такой информации об устройстве - запуск образца запроса устройства. Для CUDA 5.5 в 64-разрядной версии Windows, которая по умолчанию находится в каталоге C:\ProgramData\NVIDIA Corporation\CUDA Samples\v5.5\Bin\win64\Release\deviceQuery.exe (В Windows окно консоли будет закрыто немедленно, вы можете захотеть сначала открыть cmd и запустить его оттуда). Он имеет аналогичное расположение в Linux и MAC.

Если у вас есть Nsight для Visual Studio, просто перейдите в Nsight → Windows → Информация о системе.

Не оптимизировать рано

Я делюсь этим сегодня, потому что совсем недавно столкнулся с этой проблемой. Однако, как уже упоминалось в этой теме, принудительное включение данных в регистры - это далеко не первый шаг, который вы хотите предпринять. Сначала убедитесь, что вы действительно понимаете, что происходит, затем шаг за шагом подойдите к проблеме. Глядя на ассемблерный код, безусловно, хороший шаг, но обычно он не должен быть вашим первым. Если вы новичок в CUDA, руководство по рекомендациям CUDA поможет вам разобраться в некоторых из этих шагов.

Ответ 2

Все, что вам нужно, это:

__global__ my_kernel(int a)
{
    unsigned char p[50];
    ........
}

Компилятор автоматически прольет это, чтобы поменять локальную память, если это необходимо. Но имейте в виду, что локальная память хранится в SDRAM с графического процессора, и такая же медленная, как и глобальная память. Поэтому, если вы надеетесь, что это приведет к улучшению производительности, возможно, вы разочарованы.....

Ответ 3

~ Для кого-то, кто сталкивается с этим в будущем ~

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

Вот пример из блога nvidia Максима Милакова в 2015 году:

// Should be multiple of 32
#define THREADBLOCK_SIZE 64 
// Could be any number, but the whole array should fit into shared memory 
#define ARRAY_SIZE 32 

__device__ __forceinline__ int no_bank_conflict_index(int thread_id, int logical_index)
{
    return logical_index * THREADBLOCK_SIZE + thread_id;
}

__global__ void kernel5(float * buf, int * index_buf)
{
    // Declare shared memory array A which will hold virtual 
    // private arrays of size ARRAY_SIZE elements for all 
    // THREADBLOCK_SIZE threads of a threadblock
    __shared__ float A[ARRAY_SIZE * THREADBLOCK_SIZE]; 
    ...
    int index = index_buf[threadIdx.x + blockIdx.x * blockDim.x];

    // Here we assume thread block is 1D so threadIdx.x 
    // enumerates all threads in the thread block
    float val = A[no_bank_conflict_index(threadIdx.x, index)];
    ...
}

Ответ 4

Вы смешиваете локальные и регистрируете пространство памяти.

Отдельные переменные и массивы с постоянным размером автоматически сохраняются в пространстве регистров на чипе, практически без затрат на чтение и запись.

Если вы превысите количество регистров на мультипроцессор, они будут сохранены в локальной памяти.

Локальная память находится в глобальной памяти и имеет ту же самую медленную полосу пропускания для операций чтения и записи.

#DEFINE P_SIZE = 50

__global__ void kernel()
{
    unsigned char p[P_SIZE];
}

Ответ 5

Ключевое слово, которое вы ищете, это __shared__. Большие массивы не помещаются в пространство общей памяти, но компилятор должен использовать разделяемую память для небольшого массива фиксированного размера, как в этом случае. Вы можете использовать ключевое слово __shared__, чтобы это произошло. Вы увидите ошибку времени компиляции, если вы превысите максимальный объем разделяемой памяти для блока.