程序啓動時堆棧溢出異常（CUDA Monte Carlo Pi）

Question

我的問題是，當程序第一次進入主程序時，程序啓動時出現堆棧溢出異常。我的程序是使用CUDA的並行Monte Carlo Pi計算器。當我嘗試在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 

這使得你一個大概130MB的靜態聲明數組。我懷疑編譯器運行時庫可以處理如此大的靜態分配，這就是爲什麼你會得到即時堆棧溢出。將其替換爲動態分配，並且堆棧溢出問題將消失。但仔細閱讀Pavan's post，因爲一旦你修復了堆棧溢出問題，CUDA代碼本身在工作之前也需要重新設計。

Answer 2

您正在聲明shared memory = size的大小;像這裏

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__ 

我個人偏好方法兩個用於這些種核的，因爲共享存儲器是正比於線程的數目，但線程的數量是已知的恆定。它也稍微快一點。

編輯

除了我懷疑，這可能會導致問題太forementioned問題。

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

因爲大小本身就是這樣。

size = NUM_BLOCKS * NUM_THREADS * sizeof(float); 

可能是你想這樣做呢？

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

至於實際的堆棧溢出問題，你可能要看看talonmies post。