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

Я пытаюсь разработать небольшую программу с CUDA, но, поскольку она была МЕДЛЕННОЙ, я провел несколько тестов и немного погуглил. Я обнаружил, что хотя отдельные переменные по умолчанию хранятся в локальной памяти потоков, массивы обычно не хранятся. Я полагаю, именно поэтому это занимает так много времени, чтобы выполнить. Теперь я задаюсь вопросом: поскольку объем памяти локального потока должен составлять не менее 16 КБ, а размер моих массивов составляет всего 52 символа, есть ли способ (пожалуйста, синтаксис:)) хранить их в локальной памяти?

Не должно ли это быть что-то вроде:

__global__ my_kernel(int a)
{
  __local__ unsigned char p[50];
}
Источник

4 ответа

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

Здесь есть неправильное представление об определении "локальной памяти". "Локальная память" в 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 for 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 поможет вам разобраться в некоторых из этих шагов.

Источник

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

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

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

Источник

~ Для тех, кто столкнется с этим в будущем ~

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

Вот пример из сообщения Максима Милакова в блоге 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)];
    ...
}
Источник

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

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

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

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

#DEFINE P_SIZE = 50

__global__ void kernel()
{
    unsigned char p[P_SIZE];
}
Источник

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

Источник
Другие вопросы по тегам