Объяснение для изменения вывода кода CUDA путем изменения количества потоков

Question

Я хотел определить, сколько чисел формы x^2 + 1 являются простыми, для 1 < = x < = 10^7. Я просто хотел распараллелить его с CUDA и проверить разницу, поэтому я использовал тривиальную проверку прочности, и я не заинтересован в улучшении ее алгоритма.

Я устроил сетку и уложил ее на свой промежуток времени, записал результаты в общей памяти каждого блока, выполнил уменьшение на gpu над каждым блоком и, наконец, выполнил уменьшение CPU, чтобы получить окончательный результат.

Моя проблема заключается в том, что результат вывода изменяется, когда я изменяю количество блоков и количество потоков в каждом блоке. Еще одна вещь, которую я не могу объяснить, заключается в том, что для конфигурации из 8 блоков и 2048 потоков на блок код работает менее чем за 100 мс, но когда я уменьшаю количество потоков до 1024 и удваиваю количество блоков, код вызывает таймаут в memcpy от устройства до хоста !! Как я могу объяснить это поведение и где правильность попадает в проблему?

Я использую GTX 480 nvidia gpu.

Мой код:

#include <stdio.h> 
static void HandleError(cudaError_t err, const char *file, int line) 
{ 
    if (err != cudaSuccess) { 
     printf("%s in %s at line %d\n", cudaGetErrorString(err), file, line); 
     exit(EXIT_FAILURE); 
    } 
} 

#define HANDLE_ERROR(err) (HandleError(err, __FILE__, __LINE__)) 
#define N 10000000 
#define BLOCKS 8 
#define THREADS 2048 

__device__ int isprime(int x) 
{ 
    long long n = (long long)x*x + 1; 
    for(int p=3; p<=x+1; p+=2) 
     if (n % p == 0) return 0; 
    return 1; 
} 

__global__ void solve(int n, int* result) 
{ 
    __shared__ int ipc[THREADS]; 

    int tid = threadIdx.x; 
    int x = blockIdx.x*blockDim.x + threadIdx.x + 2; 

    // sliding grid window over interval of to-be-computed data 
    int acc = 0; 
    while(x <= n) 
    { 
     if (isprime(x)) acc++; 
     x += blockDim.x*gridDim.x; 
    } 
    ipc[tid] = acc; 
    __syncthreads(); 


    // reduction over each block in parallel 
    for(int s=blockDim.x/2; s>0; s>>=1) 
    { 
     if (tid < s) 
     { 
      ipc[tid] += ipc[tid+s]; 
     } 
     __syncthreads(); 
    } 

    if (tid == 0) result[blockIdx.x] = ipc[0]; 
} 

int main() 
{ 
    int *dev; 
    int res[BLOCKS]; 

    int ans = 0; 

    HANDLE_ERROR(cudaMalloc((void**)&dev, BLOCKS * sizeof(int))); 

    solve<<<BLOCKS, THREADS>>>(N, dev); 

    HANDLE_ERROR(cudaMemcpy(res, dev, BLOCKS*sizeof(int), cudaMemcpyDeviceToHost)); 

    // final reduction over results for each block 
    for(int j=0; j<BLOCKS; j++) 
     ans += res[j]; 

    printf("ans = %d\n", ans); 

    HANDLE_ERROR(cudaFree(dev)); 
    return 0; 
} 

Answer 1

Вы не можете запустить 2048 потоков на блок на любой текущий GPU:

#define THREADS 2048 
... 
solve<<<BLOCKS, THREADS>>>(N, dev); 
       ^
        | 
       2048 is illegal here 

Вы не делаете надлежащего cuda error checking на вызов ядра, так ваш код не скажет вам, что эта ошибка происходит.

Таким образом, в случае 2048 нитей на блок, ядро ​​даже не выполняет (и ваши результаты должны быть поддельными.)

В случае, когда вы режете нити пополам, тайм-аут, вероятно, из-за ядро занимает слишком много времени, чтобы выполнить, и windows TDR mechanism ног в.

Я попытался запустить свой код с BLOCKS = 16 и THREADS = 1024

с N = 100000, общее время выполнения было около 1,5 секунды на моем M2050 GPU. При N = 1000000 время выполнения составляло около 75 секунд. При N = 10000000, что у вас есть, время исполнения очень длинное.