исключение переполнения стека при запуске программы (CUDA Monte Carlo Pi)

Question

Моя проблема в том, что я получаю исключение переполнения стека при запуске программы, когда программа впервые входит в main. Моя программа - это калькулятор Parallel Monte Carlo Pi с использованием CUDA. Когда я пытаюсь отладить программу в Visual Studio, появляется исключение перед любой точкой останова, которую я могу выбрать. Любая помощь приветствуется.

#include <stdio.h>
#include <stdlib.h>
#include <cuda.h>
#include <curand.h>
#include <curand_kernel.h>

#define NUM_THREAD 512
#define NUM_BLOCK 65534

///////////////////////////////////////////////////////////////////////////////////////////
///////////////////////////////////////////////////////////////////////////////////////////
// Function to sum an array
__global__ void reduce0(float *g_odata) {
extern __shared__ int sdata[];

// each thread loads one element from global to shared mem
unsigned int tid = threadIdx.x;
unsigned int i = blockIdx.x*blockDim.x + threadIdx.x;
sdata[tid] = g_odata[i];
__syncthreads();

// do reduction in shared mem
for (unsigned int s=1; s < blockDim.x; s *= 2) { // step = s x 2
    if (tid % (2*s) == 0) { // only threadIDs divisible by the step participate
        sdata[tid] += sdata[tid + s];
    }
    __syncthreads();
}

// write result for this block to global mem
if (tid == 0) g_odata[blockIdx.x] = sdata[0];
}

///////////////////////////////////////////////////////////////////////////////////////////
///////////////////////////////////////////////////////////////////////////////////////////
__global__ void monteCarlo(float *g_odata, int  trials, curandState *states){
    extern __shared__ int sdata[];
//  unsigned int tid = threadIdx.x;
    unsigned int i = blockIdx.x*blockDim.x + threadIdx.x;
    unsigned int k, incircle;
    float x, y, z;
    incircle = 0;

    curand_init(1234, i, 0, &states[i]);

    for(k = 0; k < trials; k++){

    x = curand_uniform(&states[i]);
    y = curand_uniform(&states[i]);
    z = sqrt(x*x + y*y);
    if (z <= 1) incircle++;
    else{}
    }
    __syncthreads();
    g_odata[i] = incircle;
}
///////////////////////////////////////////////////////////////////////////////////////////
///////////////////////////////////////////////////////////////////////////////////////////
int main() {

    float* solution = (float*)calloc(100, sizeof(float));
    float *sumDev, sumHost[NUM_BLOCK*NUM_THREAD];
    int trials, total; 
    curandState *devStates;



    trials = 100;
    total = trials*NUM_THREAD*NUM_BLOCK;

    dim3 dimGrid(NUM_BLOCK,1,1); // Grid dimensions
    dim3 dimBlock(NUM_THREAD,1,1); // Block dimensions
    size_t size = NUM_BLOCK*NUM_THREAD*sizeof(float); //Array memory size
    cudaMalloc((void **) &sumDev, size); // Allocate array on device
    cudaMalloc((void **) &devStates, size*sizeof(curandState));
    // Do calculation on device by calling CUDA kernel
    monteCarlo <<<dimGrid, dimBlock, size>>> (sumDev, trials, devStates);
        // call reduction function to sum
    reduce0 <<<dimGrid, dimBlock, size>>> (sumDev);
    // Retrieve result from device and store it in host array
    cudaMemcpy(sumHost, sumDev, size, cudaMemcpyDeviceToHost);

    *solution = 4*(sumHost[0]/total);
    printf("%.*f\n", 1000, *solution);
    free (solution);
    //*solution = NULL;
    return 0;
}

Answer 1

Я бы предположил, что проблема заключается в следующем:

float *sumDev, sumHost[NUM_BLOCK*NUM_THREAD];

для

#define NUM_THREAD 512
#define NUM_BLOCK 65534

Это оставляет статически объявленный массив размером примерно 130 МБ. Я сомневаюсь, что библиотека времени выполнения компилятора может справиться с таким большим статическим распределением, поэтому вы получаете мгновенное переполнение стека. Замените его динамическим распределением, и проблема переполнения стека исчезнет. Но затем внимательно прочитайте пост Павана , потому что, как только вы исправите переполнение стека, сам код CUDA также нуждается в некоторой переработке, прежде чем он заработает.

Answer 2

Вы объявляете размер разделяемой памяти = размер;как здесь

monteCarlo <<<dimGrid, dimBlock, size>>>

Значение размера = 512 * 65534 * 4 = 2 ^ 9 * 2 ^ 16 * 2 ^ 2 = 2 ^ 27 (больше максимального значения общей памяти на любой карте Iможно придумать).

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

Так что вам нужно либо

1)
это для запуска ваших ядер

monteCarlo <<<dimGrid, dimBlock, (NUM_THREADS * sizeof(int))>>>

2)
Или используйте это для запуска ваших ядер

monteCarlo <<<dimGrid, dimBlock>>> 

И это для объявления вашей общей памяти внутри вашего ядра.

__shared__ int sdata[NUM_THREADS]; // Note: no extern before __shared__

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

РЕДАКТИРОВАТЬ

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

 cudaMalloc((void **) &devStates, size*sizeof(curandState));

BecuaseРазмер сам по себе такой.

size = NUM_BLOCKS * NUM_THREADS * sizeof(float);

Может быть, вы захотите сделать это вместо этого?

cudaMalloc((void **) &devStates, (NUM_BLOCKS *NUM_THREADS)*sizeof(curandState));

Что касается фактической проблемы переполнения стека, вы можете посмотреть на сообщение talonmies.