cudaMemcpyFromSymbol在__device__變量

Question

我試圖在__device__變量，其中，根據規格，駐留「在全球記憶」應用內核函數

#include <stdio.h> 
#include "sys_data.h" 
#include "my_helper.cuh" 
#include "helper_cuda.h" 
#include <cuda_runtime.h> 


double X[10] = {1,-2,3,-4,5,-6,7,-8,9,-10}; 
double Y[10] = {0}; 
__device__ double DEV_X[10]; 


int main(void) { 
    checkCudaErrors(cudaMemcpyToSymbol(DEV_X, X,10*sizeof(double))); 
    vector_projection<double><<<1,10>>>(DEV_X, 10); 
    getLastCudaError("oops"); 
    checkCudaErrors(cudaMemcpyFromSymbol(Y, DEV_X, 10*sizeof(double))); 
    return 0; 
} 

內核功能vector_projection在my_helper.cuh定義爲如下：

template<typename T> __global__ void vector_projection(T *dx, int n) { 
    int tid; 
    tid = threadIdx.x + blockIdx.x * blockDim.x; 
    if (tid < n) { 
     if (dx[tid] < 0) 
      dx[tid] = (T) 0; 
    } 
} 

正如你所看到的，我用cudaMemcpyToSymbol和cudaMemcpyFromSymbol傳輸數據和從設備。不過，我發現了以下錯誤：

CUDA error at ../src/vectorAdd.cu:19 code=4(cudaErrorLaunchFailure) 
    "cudaMemcpyFromSymbol(Y, DEV_X, 10*sizeof(double))" 

腳註：我當然可以迴避使用__device__變量和去的東西like this的正常工作;我只想看看如何用__device__變量做同樣的事情（如果可能的話）。

更新：cuda-memcheck的輸出可以在http://pastebin.com/AW9vmjFs找到。是錯誤消息我得到如下：

========= Invalid __global__ read of size 8 
=========  at 0x000000c8 in /home/ubuntu/Test0001/Debug/../src/my_helper.cuh:75:void vector_projection<double>(double*, int) 
=========  by thread (9,0,0) in block (0,0,0) 
=========  Address 0x000370e8 is out of bounds 

Answer 1

問題的根源是，你是not allowed to take the address of a device variable in ordinary host code：

vector_projection<double><<<1,10>>>(DEV_X, 10); 
            ^

儘管這看起來正確編譯，傳遞的實際地址是垃圾。

若要在主機代碼的設備變量的地址，我們可以使用cudaGetSymbolAddress

這裏是一個工作的例子，編譯和運行正常對我來說：

$ cat t577.cu 
#include <stdio.h> 

double X[10] = {1,-2,3,-4,5,-6,7,-8,9,-10}; 
double Y[10] = {0}; 
__device__ double DEV_X[10]; 

template<typename T> __global__ void vector_projection(T *dx, int n) { 
    int tid; 
    tid = threadIdx.x + blockIdx.x * blockDim.x; 
    if (tid < n) { 
     if (dx[tid] < 0) 
      dx[tid] = (T) 0; 
    } 
} 



int main(void) { 
    cudaMemcpyToSymbol(DEV_X, X,10*sizeof(double)); 
    double *my_dx; 
    cudaGetSymbolAddress((void **)&my_dx, DEV_X); 
    vector_projection<double><<<1,10>>>(my_dx, 10); 
    cudaMemcpyFromSymbol(Y, DEV_X, 10*sizeof(double)); 
    for (int i = 0; i < 10; i++) 
     printf("%d: %f\n", i, Y[i]); 
    return 0; 
} 
$ nvcc -arch=sm_35 -o t577 t577.cu 
$ cuda-memcheck ./t577 
========= CUDA-MEMCHECK 
0: 1.000000 
1: 0.000000 
2: 3.000000 
3: 0.000000 
4: 5.000000 
5: 0.000000 
6: 7.000000 
7: 0.000000 
8: 9.000000 
9: 0.000000 
========= ERROR SUMMARY: 0 errors 
$ 

這是不是唯一的方法解決這個問題。它是合法的取設備代碼的設備變量的地址，所以你可以用一條線像這樣修改你的內核：

T *dx = DEV_X; 

，並放棄該設備的變量傳遞爲內核參數。正如評論中所建議的那樣，您也可以修改代碼以使用Unified Memory。

關於錯誤檢查，如果您偏離proper cuda error checking並且對您的偏差沒有小心，結果可能會令人困惑。大多數cuda API調用除了由其自身行爲引發的錯誤之外，還可以返回由以前的一些CUDA異步活動（通常是內核調用）導致的錯誤。