Понимание основных концепций CUDA посредством сложения векторов

Я открываю эту тему, потому что заметил странное поведение в выводе моего кода, пытаясь получить представление о некоторых основных понятиях CUDA, таких как скорость и количество блоков/потоков и т. д. Любая помощь будет оценена по достоинству!

Прежде всего, вот некоторые характеристики моей видеокарты:
Название: GeForce 8600M GT
Количество мультипроцессоров: 4
Максимальное количество потоков на блок: 512
Максимальные размеры сетки: (65535, 65535, 1 )

Я играл со следующим простым кодом. Он заполняет 3 массива длины N единицами и вычисляет сумму. Сумма, очевидно, предсказуема и равна 3N.

#include <iostream>
#include "ArgumentParser.h"

//using namespace std;

__global__ void addVector(int *a, int *b, int *c, int *d, int *N){
    int tid = threadIdx.x + blockIdx.x * blockDim.x;

    if (tid<*N) {
        d[tid] = a[tid] + b[tid] + c[tid];
    }
}

int main(int argc, char *argv[]) {
    //Handy way to pass command-line arguments.
    ArgumentParser parser(argc, argv);
    int nblocks = parser("-nblocks").asInt(1);
    int nthreads = parser("-nthreads").asInt(1);

    //Defining arrays on host.
    int N = 100000;
    int a[N];
    int b[N];
    int c[N];
    int d[N];

    //Pointers to the arrays that will go to the device.
    int *dev_a;
    int *dev_b;
    int *dev_c;
    int *dev_d;
    int *dev_N;

    //Filling up a, b, and c.
    for (int i=0; i<N; i++){
        a[i] = 1;
        b[i] = 1;
        c[i] = 1;
    } 

    //Modifying the memory adress of dev_x so that dev_x is on the device and //
    //the proper memory size is reserved for it.  
    cudaMalloc((void**)&dev_a, N * sizeof(int));
    cudaMalloc((void**)&dev_b, N * sizeof(int));
    cudaMalloc((void**)&dev_c, N * sizeof(int));
    cudaMalloc((void**)&dev_d, N * sizeof(int));
    cudaMalloc((void**)&dev_N, sizeof(int));

    //Copying the content of a/b/c and N to from the host to the device.
    cudaMemcpy(dev_a, a, N * sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(dev_b, b, N * sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(dev_c, c, N * sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(dev_N, &N, sizeof(int), cudaMemcpyHostToDevice);

    //Initializing the cuda timers.
    cudaEvent_t start, stop;
    cudaEventCreate(&start); 
    cudaEventCreate(&stop);
    cudaEventRecord (start, 0);

    //Executing the kernel.
    addVector<<<nblocks, nthreads>>>(dev_a, dev_b, dev_c, dev_d, dev_N);

    cudaEventRecord(stop, 0);
    cudaEventSynchronize(stop);
    float time;
    cudaEventElapsedTime(&time, start, stop);
    printf ("CUDA time: %3.5f s\n", time/1000);

    //Copying the result from device to host.
    cudaMemcpy(d, dev_d, N * sizeof(int), cudaMemcpyDeviceToHost);

    //Freeing the memory allocated on the GPU
    cudaFree(dev_a);
    cudaFree(dev_b);
    cudaFree(dev_c);
    cudaFree(dev_d);
    cudaFree(dev_N);

    //Checking the predictable result.
    int sum=0;
    for (int i=0; i<N; i++){
        sum += d[i];
    }
    printf("Result of the sum: %d. It should be: %d.\n", sum, 3*N);
}

Вопрос 1.
Когда я компилирую код и набираю:

./addArrayCuda -nblocks 1 -nthreads 1

получаю в ответ:

Result of the sum: -642264408. It should be: 300000.

Это кажется разумным. Я использую один блок с одним потоком. Будет добавлен только первый элемент каждого массива. Остальные элементы, являющиеся случайными значениями, в сумме дают нечто непредсказуемое. Должно быть так, что nblocks * nthreads >= N. Итак, попробуем:

./addArrayCuda -nblocks 3125 -nthreads 32

Результат:

Result of the sum: 300000. It should be: 300000.

Это имеет смысл. 3125*32=100000=N. Пока здесь все нормально. Однако, если я повторно запускаю предыдущую команду (с nblocks = nthreads = 1) без перекомпиляции, я получаю:

./addArrayCuda -nblocks 1 -nthreads 1
Result of the sum: 300000. It should be: 300000.

Что случилось?

Вопрос 2. Этот вопрос касается взаимосвязи между nblocks/nthreads и скоростью выполнения. Я понимаю, что этот вопрос может не иметь особого смысла, если проблема в коде объясняет вопрос 1, но позвольте мне все же задать его. Я посмотрел на время выполнения кода (в среднем за 5 прогонов) с различным количеством блоков/потоков, но убедился, что nblocks * nthreads > N. Вот что у меня есть (у меня есть хороший сюжет, но не достаточно репутации, чтобы опубликовать это...):

(nblocks, nthreads) время выполнения [s] коэффициент увеличения
(196, 512) 5.0e-4 -
(391, 256) 4.8e-4 1.0
(782, 128) 4.8e-4 1,0
(1563, 64) 4,9e-4 1,0
(3125, 32) 5,0e-4 1,0
(6250, 16) 5,2e-4 1,0
(12500, 8) 9,0e -4 1,7
(25000, 4) 1,3e-3 1,4
(50000, 2) 2,3e-3 1,8

Моя интерпретация: графический процессор разделен на блоки, и каждый блок разделен на потоки. Каждый такт GPU отправляет ядро ​​на 4 блока (счетчик многопроцессорных систем) и внутри каждого из этих блоков на деформацию (группа из 32 потоков). Это означает, что использование числа потоков, не кратного 32, является пустой тратой ресурсов. Таким образом, мы можем понять общую связь между (nblocks, nthreads) и временем выполнения. От (196, 512) до (3125, 32) количество тактов, используемых GPU, примерно одинаково и приблизительно пропорционально (nblocks/4) * (nthreads/32). Однако мы примерно ожидаем удвоения времени выполнения между (3125, 32) и (6250, 16), (6250, 16) и (12500, 8) и так далее.

Почему это не так? Точнее, почему нет существенной разницы во времени выполнения между (3125, 32) и (6250, 16)?

Я благодарю вас за то, что нашли время, чтобы прочитать до сих пор ;-)