cuda錯誤：意外啓動失敗

Question

我是cuda的新手，我正嘗試在CUDA上實現與Simpson method的數值集成。

我的代碼出現錯誤「意外啓動失敗」。它看起來像是在gpu內存中的segfault。但這很奇怪，因爲它取決於變量step，它控制迭代次數，而不是任何內存操作。例如，當我運行integrate_with_cuda與step = 0.00001它工作正常，結果是正確的，但如果我更改step在0.000001，我的程序下降。

這是我的代碼：

#include "device_launch_parameters.h" 
#include "cuda_runtime_api.h" 
#include "cuda.h" 
#include "cuda_safe_call.h" 
#include <cmath> 
#include <iostream> 

typedef double(*cuda_func)(double, double); 

struct cuda_expr { 
    cuda_func func; 
    int dest; 
    int op1; 
    int op2; 
}; 

enum cuda_method { 
    cm_Add, 
    cm_Mult 
}; 

__device__ double add_func(double x, double y) { 
    return x + y; 
} 
__device__ cuda_func p_add_func = add_func; 

__device__ double mult_func(double x, double y) { 
    return x*y; 
} 
__device__ cuda_func p_mult_func = mult_func; 

__host__ cuda_func get_cuda_func(cuda_method method) { 
    cuda_func result = NULL; 

    switch (method) { 
    case cm_Add: 
     cudaMemcpyFromSymbol(&result, p_add_func, sizeof(cuda_func)); 
     break; 
    case cm_Mult: 
     cudaMemcpyFromSymbol(&result, p_mult_func, sizeof(cuda_func)); 
     break; 
    } 
    return result; 
} 

__device__ double atomicAdd(double* address, double val) 
{ 
    unsigned long long int* address_as_ull = 
     (unsigned long long int*)address; 
    unsigned long long int old = *address_as_ull, assumed; 
    do { 
     assumed = old; 
     old = atomicCAS(address_as_ull, assumed, 
      __double_as_longlong(val + 
      __longlong_as_double(assumed))); 
    } while (assumed != old); 
    return __longlong_as_double(old); 
} 

__device__ void computate_expr(cuda_expr* expr, int expr_length, double* vars, int vars_count) { 
    for (cuda_expr* step = expr, *end = expr + expr_length; step != end; ++step) { 
     vars[step->dest] = (*step->func)(vars[step->op1], vars[step->op2]); 
    } 
} 

__device__ double simpson_step(cuda_expr* expr, int expr_length, double* vars, int vars_count, double a, double b, double c) { 
    double f_a; 
    double f_b; 
    double f_c; 
    vars[0] = a; 
    computate_expr(expr, expr_length, vars, vars_count); 
    f_a = vars[vars_count - 1]; 
    vars[0] = b; 
    computate_expr(expr, expr_length, vars, vars_count); 
    f_b = vars[vars_count - 1]; 
    vars[0] = c; 
    computate_expr(expr, expr_length, vars, vars_count); 
    f_c = vars[vars_count - 1]; 
    return (c - a)/6 * (f_a + 4 * f_b + f_c); 
} 

__global__ void integrate_kernel(cuda_expr* expr, int expr_length, double* vars, int vars_count, double from, double to, double step, double* res) { 
    int index = blockIdx.x*blockDim.x + threadIdx.x; 
    int threads_count = gridDim.x*blockDim.x; 
    double* my_vars = vars + index * vars_count; 

    double my_from = from + index*(to - from)/threads_count; 
    double my_to = from + (index + 1)*(to - from)/threads_count; 

    double my_res = 0; 

    double a = my_from; 
    double b = my_from + step/2; 
    double c = my_from + step; 

    while (c < (my_to + step/10)) { 
     my_res += simpson_step(expr, expr_length, my_vars, vars_count, a, b, c); 
     a += step; 
     b += step; 
     c += step; 
    } 
    atomicAdd(res, my_res); 
} 

__host__ double integrate_with_cuda(const cuda_expr* expr, int expr_length, double* vars, int vars_count, double from, double to, double step) { 
    const int blockSize = 32; 
    const int gridSize = 2; 
    const int threadsCount = blockSize*gridSize; 

    cuda_expr* d_expr; 
    CudaSafeCall(cudaMalloc((void**)&d_expr, expr_length*sizeof(cuda_expr))); 
    CudaSafeCall(cudaMemcpy(d_expr, expr, expr_length*sizeof(cuda_expr), cudaMemcpyHostToDevice)); 

    double* d_vars; //allocate own vars array for every thread 
    CudaSafeCall(cudaMalloc((void**)&d_vars, threadsCount*vars_count*sizeof(double))); 
    for (int i = 0; i < threadsCount; ++i) { 
     CudaSafeCall(cudaMemcpy(d_vars + i*vars_count, vars, vars_count*sizeof(double), cudaMemcpyHostToDevice)); 
    } 

    double* d_res; 
    double result = 0; 
    CudaSafeCall(cudaMalloc((void**)&d_res, sizeof(double))); 
    CudaSafeCall(cudaMemcpy(d_res, &result, sizeof(double), cudaMemcpyHostToDevice)); 

    integrate_kernel<<<gridSize, blockSize>>>(d_expr, expr_length, d_vars, vars_count, from, to, step, d_res); 

    CudaSafeCall(cudaMemcpy(&result, d_res, sizeof(double), cudaMemcpyDeviceToHost)); 

    CudaSafeCall(cudaFree(d_expr)); 
    CudaSafeCall(cudaFree(d_vars)); 
    CudaSafeCall(cudaFree(d_res)); 
    return result; 
} 

int main() { 
    cuda_expr expr[3] = { 
     { get_cuda_func(cuda_method::cm_Add), 4, 1, 0 }, 
     { get_cuda_func(cuda_method::cm_Add), 3, 0, 2 }, 
     { get_cuda_func(cuda_method::cm_Mult), 5, 3, 4 } 
    }; 
    double vars[6] = {0, 10, 1, 0, 0, 0}; 

    double res = integrate_with_cuda(expr, 3, vars, 6, 0, 10, 0.00001); 

    std::cout << res << std::endl; 
    system("PAUSE"); 
} 

我想，我需要給它是如何工作的一些解釋。函數integrate_with_cuda將cuda_expr的輸入數組和雙精度數組作爲變量。 cuda_expr數組表示數學表達式的語法樹，它在數組中展開。 cuda_expr :: func指向設備函數，該函數將與args vars [cuda_expr :: op1]和vars [cuda_expr :: op2]一起調用，結果將放入vars [cuda_expr :: dest]中。變量數組中的第一個單元格保留爲x變量。
main函數中的測試示例表示表達式(1+x)*(x+10)。計算數組中的第一個cuda_expr從變量中獲得第二個和第一個（它是x）單元，將它們添加並放到變量[4]中，第二個cuda_expr從變量中獲取第一個和第三個單元，將它們添加到變量[5]，最後一個cuda_expr獲取第4和第5個單元格（第一個和第二個cuda_expr將結果放入它們中），將其放大並放到最後一個變量單元格中。 變量的最後一個單元格是計算後的表達式的結果。

我使用MS的Visual Studio 2013（與V110平臺工具包），定期標誌（sm_30弓沒有CUDA調試）：

nvcc.exe -gencode=arch=compute_30,code=\"sm_30,compute_30\" --use-local-env --cl-version 2012 -ccbin "C:\Program Files (x86)\Microsoft Visual Studio 11.0\VC\bin\x86_amd64" -I"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v6.0\include" -I"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v6.0\include"  --keep-dir x64\Tests -maxrregcount=0 --machine 64 --compile -cudart static  -DWIN32 -D_DEBUG -D_UNICODE -DUNICODE -Xcompiler "/EHsc /W3 /nologo /Od /Zi /MDd " -o x64\Tests\integration_on_cuda.cu.obj integration_on_cuda.cu 

感謝。對不起，我的英文:)

Answer 1

當我在linux上運行你的代碼時，我得到983.333是否指定step爲0.00001或0.000001。但是，如果步驟較小，代碼運行時間會更長。

既然你在windows上運行，這可能不過是你打the windows TDR mechanism。在Windows上運行時間超過2秒的內核可能會觸發TDR機制。通常當這種情況發生時，隨着GPU經過窗口觸發的重置，您將看到屏幕閃爍爲黑色，然後重新繪製自己。您也可能會看到系統信息。如果從VS內運行代碼與從命令行運行代碼，則確切的行爲也可能不同。

請參考以上鍊接或在CUDA標籤上搜索如何修改TDR機制。