Таблица поиска с использованием Cuda-C

Я придумал решение для этого поста с использованием алгоритмического подхода. Мне было любопытно попробовать подход к таблице поиска, предложенный в одном из комментариев к посту. Я новичок в CUDA C и пытался найти примеры/информацию о том, как это можно сделать. У меня есть значения, хранящиеся в таблице ниже. Я знаю, что мне нужно связать каждый поток, чтобы вытащить каждое из 4 значений. Значения соответствуют индексам SubBlkIdxA, SubBlkIdxB, BlkIdxA и BlkIdxB соответственно для каждого потока. Как только они считываются из таблицы, они передаются функции для вычисления чего-либо.

Я знаю, что если я скажу m_aIdx[3][0], он пойдет {3,0,0,1,}, запись в таблице и прочитает первую запись «3». Чтобы прочитать каждую запись в этом месте по упомянутым выше индексам, я думаю так:

Моя таблица выглядит так:

static __constant__ int16 m_aIdx[64][4] =
{
    {0,1,0,0,},
    {2,3,0,0,},
    {1,0,0,1,},
    {3,0,0,1,},
    {1,2,0,1,},
    {3,2,0,1,},
    and so on ... upto 64 entries
}

Функционируйте следующим образом:

static __device__ void func()
{
    SubBlkIdxA = m_aIdx[3][0];
    SubBlkIdxB = m_aIdx[3][1];
    BlkIdxA = m_aIdx[3][2];
    BlkIdxB = m_aIdx[3][3];

    func1(SubBlkIdxA, SubBlkIdxB, BlkIdxA, BlkIdxB);
}

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


person searocks    schedule 16.07.2013    source источник
comment
Из другого поста кажется, что у вас есть код для генерации индексов в соответствии с алгоритмическим подходом. Выше у вас есть черновик кода (похоже, вам просто нужно обратиться к матрице m_aIdx по индексам потоков/блоков) для неалгоритмического подхода. Почему бы вам не сравнить относительную скорость двух решений?   -  person Vitality    schedule 17.07.2013


Ответы (1)


Либо должно быть довольно быстро. В вашем «алгоритмическом» подходе вы вычисляете индексы на основе данных, хранящихся в регистрах, что будет довольно быстро. В этом подходе вы делаете относительно хорошо объединенную память, доступ к 512 байтам постоянной памяти, что также довольно быстро. (Даже если он был плохо объединен, он будет кэшироваться довольно быстро).

Что меня беспокоит, так это то, как вы используете эти индексы в func1. Если операторы по этим индексам могут привести к некоторому плохому расхождению, а доступ к памяти с этими индексами может привести к некоторым плохо объединенным передачам.

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

P.S. Я не совсем уверен, как устроены ваши подблоки, поскольку я не улавливаю шаблон ваших индексов, и я не понимаю, почему вы создаете подблоки внутри блока вместо использования более мелких блоков.

person Zero880    schedule 16.07.2013
comment
Возможно, для архитектуры Kepler он мог бы использовать кэш текстур только для чтения вместо постоянной памяти? - person Vitality; 17.07.2013
comment
Спасибо вам всем. Я сравнил скорости, и оба метода, по-видимому, дают почти одинаковые значения скорости, но ничего существенного не отличаются. - person searocks; 17.07.2013
comment
В шаблоне индексов мне нужно различать границы между краями подблоков, а также между краями блоков, поскольку я вычисляю значение прочности границы вдоль этих краев. Так обрабатываются данные. - person searocks; 17.07.2013