Параллельная запись в битовый набор (массив массива) на GPU

Question

Вот вопрос об алгоритмах графического процессора, относящийся к проблеме, которую я пытаюсь ускорить:

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

bool is_a_foo[131072][512];

bool в этой структуре представляют, являются ли данные в другом месте (которые имеют схожие измерения ... но это не имеет значения) как foo.Для простоты, давайте предположим, что я просто работаю на одном блоке графического процессора, с каждым разрывающимся потоком (на шаге блокировки через __syncwarp() ... но, пожалуйста, не позволяйте этому быть слишком отвлекающим, как на практике я делаючто-то более чувственное) локации 0 -> 131071.Другими словами, код каждого потока выглядит примерно так:

// assume is_a_foo is initialized earlier to 0's by some sort of memset call
// assume that the values for is_a_foo can go from false->true but never from true->false
for (int i = 0; i < 131072; ++i) {
    if (something_kind_of_expensive_but_not_the_bottleneck()) {
        is_a_foo[ i ][thread] = true;
    }
}

Если каждый bool представлен как 8 битов, данные не теряются.Тем не менее, давайте предположим, что я хотел бы уменьшить объем занимаемой памяти / кеша и потребление пропускной способности.Вместо этого мы могли бы представить вышеупомянутую структуру данных как:

unsigned int is_a_foo[131072][512 / (sizeof(unsigned int) * 8)];

И мы можем выполнить битовую арифметику, чтобы установить интересующий бит равным 1.

Проблема в том, что без какой-либо специальной обработки,записи в is_a_foo будут разбивать друг друга, и не каждый бит, который должен быть установлен в 1, обязательно будет установлен в 1.

В случае, если мы хотим сделать что-то особенное, мыможно использовать atomicCAS, чтобы гарантировать, что никакие записи не будут потеряны.К сожалению, это кажется довольно дорогим.Действительно, в моем приложении, где запуск ядра занимает около 30 миллисекунд, время выполнения ядра увеличивается на ~ 33%.В настоящее время неясно, связано ли дополнительное время с атомной операцией или дополнительными инструкциями, но я подозреваю, что это атомная операция.

Одна вещь, которая смогла бы смягчить ущерб, была бы, если бы я мог работать на unsigned charс вместо unsigned int с.К сожалению, CUDA не предоставляет такого интерфейса.И, когда я работаю на unsigned short s, я получаю ошибку компилятора о том, что функция недоступна для unsigned short s (подробности доступны по запросу).

Все это спрашивать, Существуют ли алгоритмы / структуры данных, которые хорошо подходят для этого типа операций на графическом процессоре ?

Answer 1

Мне не известно ни о каком графическом процессоре с поддержкой CUDA с размером деформации 512, поэтому я предполагаю, что вы хотели записать размер блока и __syncthreads() вместо размера деформации и __syncwarp() (размер деформации составляет 32 на каждая существующая архитектура CUDA существует). Я также могу обратить ваше внимание на то, что существует функция atomicOr().

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

Answer 2

Рассматривали ли вы упаковку своих бит по-другому?Если последовательные биты в int принадлежат компоненту first вашего 2D-массива, а не второму, вы выиграете от меньшего объема памяти, избегая ложного совместного использования.

Рассмотрите структуру:

static constexpr bits = sizeof(unsigned int) * 8;

class IsAFoo {
  private:
    static constexpr size = 131072/bits;
    unsigned int data[size][512];
  public:
    __host__ __device__ void set(int i, int thread, bool value) {
      unsigned int bit = 1u << (i%bits);
      if (value)
        data[i/bits][thread] |= bit;
      else
        data[i/bits][thread] &= ~(bit);
    }
    __host__ __device__ bool get(int i, int thread) {
      return bool(data[i/bits][thread] & (1u << (i%bits));
    }
}

__device__ IsAFoo is_a_foo;

... и тогда весь остальной алгоритм будет работать как прежде - вам просто нужно использовать вышеупомянутые функции set и get.Это, очевидно, предполагает, что где-либо еще в вашей программе вы не пытаетесь изменить массив, используя другой шаблон, например set(threadIdx.x, commonValue).

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

unsigned int tmpFlags = 0;
for (int i = 0; i < 131072; ++i) {
    if (something_kind_of_expensive_but_not_the_bottleneck()) {
        tmpFlags |= 1u << (i % bits)
    }
    if (i % bits == bits - 1) {
        is_a_foo.setBulk(i, threadIdx.x, tmpFlags)
        tmpFlags = 0;
    }
}

(при условии, что setBulk указан в классе IsAFoo).Это сократит общее количество глобальных операций с памятью в 32 раза за счет одного дополнительного действующего регистра и нескольких арифметических операций.