Организация данных в ядрах Cuda

Question

Я новичок в Cuda и читаю учебные пособия и другой открытый исходный код, чтобы попытаться понять вещи.Я знаю общую концепцию иерархий потоков.

TL; DR, все прочитанные мною учебные пособия предполагали, что данные, отправляемые в ядро, также организованы в этой иерархии, без явного указания перед запуском ядра.Разве данные, передаваемые ядру, не должны быть переупорядочены в иерархии grid> block> thread перед передачей в ядро?Ниже приведены два фрагмента, которые смутили меня в этом отношении.

Я следовал этому x_plus_y уроку здесь .В этом уроке следующий фрагмент:

_global__
void add(int n, float *x, float *y)
{
  int index = blockIdx.x * blockDim.x + threadIdx.x;
  int stride = blockDim.x * gridDim.x;
  for (int i = index; i < n; i += stride)
    y[i] = x[i] + y[i];
}

В приведенном выше фрагменте мы хотим добавить соответствующие элементы в x и y, но как мне узнать, что x и y помещены в графический процессор так, что индекс i (который вычисляется с использованием blockIdx, blockDim и т. д.) фактически указывает на соответствующие элементы x и y.Если x и y помещаются один за другим в память, не должен ли индекс, используемый для y, учитывать длину x?Мне не хватает некоторого ключевого интуитивного понимания здесь.Кроме того, как мне узнать, где в графическом процессоре был отображен какой-то случайный элемент массива, скажем, x [1011]?Или мне не нужно заботиться о точном позиционировании моих данных из-за какой-то абстракции?

Я также расскажу о другом фрагменте из репозитория с открытым исходным кодом.Это ядро ​​для вычисления метрики расстояния между двумя наборами облаков точек.Каждое облако представляет собой матрицу Nx3 (имеет N 3-D точек).

b - размер пакета (таким образом, b количество облаков передается ядру)

n - количество точек в каждом облаке первого набора

m - количество точек в каждом облаке второго сета.

Например, первый набор облаков может быть (16,1024,3), а второй набор (16,512,3):

__global__ void NmDistanceKernel(int b,int n,const float * xyz,int m,const float * xyz2,float * result,int * result_i){
    const int batch=512;
    __shared__ float buf[batch*3];
    for (int i=blockIdx.x;i<b;i+=gridDim.x){
        for (int k2=0;k2<m;k2+=batch){
            int end_k=min(m,k2+batch)-k2;
            for (int j=threadIdx.x;j<end_k*3;j+=blockDim.x){
                buf[j]=xyz2[(i*m+k2)*3+j];
            }


        for (int j=threadIdx.x+blockIdx.y*blockDim.x;j<n;j+=blockDim.x*gridDim.y){
                float x1=xyz[(i*n+j)*3+0];
                float y1=xyz[(i*n+j)*3+1];
                float z1=xyz[(i*n+j)*3+2];
            }
    }
}

Вышеприведенное ядро ​​запускается следующим образом:

NmDistanceKernel<<<dim3(32,16,1),512>>>(batch_size, n, xyz1.data<float>(), m, xyz2.data<float>(), dist1.data<float>(), idx1.data<int>());

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

Answer 1

Перед вызовом ядра вы должны поместить данные в графический процессор.

Данные в основном передаются в массивах данных, поэтому структура этих массивов на графическом процессоре такая же, как и вкод вашего хоста.

В первом примере массивы x и y передаются отдельно, поэтому индексы для x и y оба начинаются с 0.Вы можете передать их в одном большом массиве, и тогда потребуется скорректировать индексирование.

Это было сделано в другом примере.Массив xyz состоит из значений xy и z всех точек.Заказ идет как x1 y1 z1 x2 y2 z2 x3 y3 z3 ....Вот почему при доступе к значениям вы видите x = [...]+0; y = [...]+1; z = [...]+2;.Для следующего пункта все индексы увеличиваются на 3.

Чтобы затем получить доступ к вашим данным в ядрах, вам нужно обратиться к идентификаторам, которые дает вам CUDA.Вы используете положение потока внутри вашей сетки и блоков.

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

int index = blockIdx.x * blockDim.x + threadIdx.x;

threadIdx.xговорит нам, где находится поток в блоке, поэтому было бы достаточно, если бы мы запускали только один блок.Но тогда разные потоки в разных блоках будут иметь одинаковый индекс.Мы должны отделить их, получив их blockIdx.x.Длина блока составляет blockDim.x, и первый поток во втором блоке должен продолжаться после последнего потока в блоке 1. Таким образом, формируется приведенная выше формула для index.

Затем каждый поток переходит вперед, так чтозатем самый первый поток читает первые данные после данных, которые только что прочитал последний поток, и т. д.

Чем больше измерений использует ваша запущенная сетка, тем сложнее эти вычисления.Попробуйте начать с простых сеток и увеличьте сложность, если вам удобно.