從設備的cublas矩陣反轉

Question

我想從設備運行矩陣反轉。如果從主機調用此邏輯，則工作正常。

編譯行如下（Linux的）：

nvcc -ccbin g++ -arch=sm_35 -rdc=true simple-inv.cu -o simple-inv -lcublas_device -lcudadevrt 

我碰到下面的警告，我似乎無法化解。 （我的GPU是開普勒我不知道爲什麼它正試圖鏈接到麥克斯韋程序我有Cuda的6.5-14。）：

handle 0 n = 3 
simple-inv.cu:63 Error [an illegal memory access was encountered] 

測試：

nvlink warning : SM Arch ('sm_35') not found in '/usr/local/cuda/bin/../targets/x86_64-linux/lib/libcublas_device.a:maxwell_sm50_sgemm.o' 

程序與運行程序如下：

#include <stdio.h> 
#include <stdlib.h> 
#include <math.h> 
#include <cuda_runtime.h> 
#include <cublas_v2.h> 

#define PERR(call) \ 
    if (call) {\ 
    fprintf(stderr, "%s:%d Error [%s] on "#call"\n", __FILE__, __LINE__,\ 
     cudaGetErrorString(cudaGetLastError()));\ 
    exit(1);\ 
    } 
#define ERRCHECK \ 
    if (cudaPeekAtLastError()) { \ 
    fprintf(stderr, "%s:%d Error [%s]\n", __FILE__, __LINE__,\ 
     cudaGetErrorString(cudaGetLastError()));\ 
    exit(1);\ 
    } 

__global__ void 
inv_kernel(float *a_i, float *c_o, int n) 
{ 
    int p[3], info[1], batch; 
    cublasHandle_t hdl; 
    cublasStatus_t status = cublasCreate_v2(&hdl); 
    printf("handle %d n = %d\n", status, n); 

    info[0] = 0; 
    batch = 1; 
    float *a[] = {a_i}; 
    const float *aconst[] = {a_i}; 
    float *c[] = {c_o}; 
    // See 
    // http://docs.nvidia.com/cuda/pdf/CUDA_Dynamic_Parallelism_Programming_Guide.pdf 
    //http://stackoverflow.com/questions/27094612/cublas-matrix-inversion-from-device 

    status = cublasSgetrfBatched(hdl, n, a, n, p, info, batch); 
    __syncthreads(); 
    printf("rf %d info %d\n", status, info[0]); 
    status = cublasSgetriBatched(hdl, n, aconst, n, p, 
     c, n, info, batch); 
    __syncthreads(); 
    printf("ri %d info %d\n", status, info[0]); 

    cublasDestroy_v2(hdl); 
    printf("done\n"); 
} 
static void 
run_inv(float *in, float *out, int n) 
{ 
    float *a_d, *c_d; 

    PERR(cudaMalloc(&a_d, n*n*sizeof(float))); 
    PERR(cudaMalloc(&c_d, n*n*sizeof(float))); 
    PERR(cudaMemcpy(a_d, in, n*n*sizeof(float), cudaMemcpyHostToDevice)); 

    inv_kernel<<<1, 1>>>(a_d, c_d, n); 

    cudaDeviceSynchronize(); 
    ERRCHECK; 

    PERR(cudaMemcpy(out, c_d, n*n*sizeof(float), cudaMemcpyDeviceToHost)); 
    PERR(cudaFree(a_d)); 
    PERR(cudaFree(c_d)); 
} 

int 
main(int argc, char **argv) 
{ 
    float c[9]; 
    float a[] = { 
    1, 2, 3, 
    0, 4, 5, 
    1, 0, 6 }; 

    run_inv(a, c, 3); 
    return 0; 
} 

我也跟着指導在http://docs.nvidia.com/cuda/cublas/index.html#device-api第2.1.9節，但我懷疑我忽略了一些東西。

注意：11月24日編輯使用正確的指針輸入。這仍然報告內核中的非法內存訪問。

Answer 1

有關sm_50的警告是良性的。這是我說「在這種情況下可以安全地忽略它」的方式。

關於您當前發佈的代碼，問題與動態並行性文檔中介紹的有關使用線程本地內存here的內容有關。

簡而言之，父線程的本地內存在子內核啓動中「超出範圍」。雖然它不是很明顯，但來自設備代碼的cublas調用是（嘗試）啓動子內核。這意味着，這樣的聲明：

int p[3], info[1], 

將是有問題的，如果這些指針（例如p，info）被傳遞給子內核。指針本身的數值不會被破壞，但是它們不會指向子內核的內存空間中的任何「有意義」的東西。

有多種方法可以解決這個問題，但一種可能的解決方案是用「設備堆」中的分配替換此類型的任何堆棧/本地分配，這可以通過in-kernel malloc進行。

這是一個完全有效的代碼/示例，似乎對我來說正確工作。對於給定樣本矩陣的反演，輸出似乎是正確的：

$ cat t605.cu 
#include <stdio.h> 
#include <stdlib.h> 
#include <math.h> 
#include <cuda_runtime.h> 
#include <cublas_v2.h> 

#define PERR(call) \ 
    if (call) {\ 
    fprintf(stderr, "%s:%d Error [%s] on "#call"\n", __FILE__, __LINE__,\ 
     cudaGetErrorString(cudaGetLastError()));\ 
    exit(1);\ 
    } 
#define ERRCHECK \ 
    if (cudaPeekAtLastError()) { \ 
    fprintf(stderr, "%s:%d Error [%s]\n", __FILE__, __LINE__,\ 
     cudaGetErrorString(cudaGetLastError()));\ 
    exit(1);\ 
    } 

__global__ void 
inv_kernel(float *a_i, float *c_o, int n) 
{ 
    int *p = (int *)malloc(3*sizeof(int)); 
    int *info = (int *)malloc(sizeof(int)); 
    int batch; 
    cublasHandle_t hdl; 
    cublasStatus_t status = cublasCreate_v2(&hdl); 
    printf("handle %d n = %d\n", status, n); 

    info[0] = 0; 
    batch = 1; 
    float **a = (float **)malloc(sizeof(float *)); 
    *a = a_i; 
    const float **aconst = (const float **)a; 
    float **c = (float **)malloc(sizeof(float *)); 
    *c = c_o; 
    // See 
    // http://docs.nvidia.com/cuda/pdf/CUDA_Dynamic_Parallelism_Programming_Guide.pdf 
    //http://stackoverflow.com/questions/27094612/cublas-matrix-inversion-from-device 
    status = cublasSgetrfBatched(hdl, n, a, n, p, info, batch); 
    __syncthreads(); 
    printf("rf %d info %d\n", status, info[0]); 
    status = cublasSgetriBatched(hdl, n, aconst, n, p, 
     c, n, info, batch); 
    __syncthreads(); 
    printf("ri %d info %d\n", status, info[0]); 
    cublasDestroy_v2(hdl); 
    printf("done\n"); 
} 
static void 
run_inv(float *in, float *out, int n) 
{ 
    float *a_d, *c_d; 

    PERR(cudaMalloc(&a_d, n*n*sizeof(float))); 
    PERR(cudaMalloc(&c_d, n*n*sizeof(float))); 
    PERR(cudaMemcpy(a_d, in, n*n*sizeof(float), cudaMemcpyHostToDevice)); 

    inv_kernel<<<1, 1>>>(a_d, c_d, n); 

    cudaDeviceSynchronize(); 
    ERRCHECK; 

    PERR(cudaMemcpy(out, c_d, n*n*sizeof(float), cudaMemcpyDeviceToHost)); 
    PERR(cudaFree(a_d)); 
    PERR(cudaFree(c_d)); 
} 

int 
main(int argc, char **argv) 
{ 
    float c[9]; 
    float a[] = { 
    1, 2, 3, 
    0, 4, 5, 
    1, 0, 6 }; 

    run_inv(a, c, 3); 
    for (int i = 0; i < 3; i++){ 
    for (int j = 0; j < 3; j++) printf("%f, ",c[(3*i)+j]); 
    printf("\n");} 

    return 0; 
} 
$ nvcc -arch=sm_35 -rdc=true -o t605 t605.cu -lcublas_device -lcudadevrt 
nvlink warning : SM Arch ('sm_35') not found in '/shared/apps/cuda/CUDA-v6.5.14/bin/..//lib64/libcublas_device.a:maxwell_sgemm.asm.o' 
nvlink warning : SM Arch ('sm_35') not found in '/shared/apps/cuda/CUDA-v6.5.14/bin/..//lib64/libcublas_device.a:maxwell_sm50_sgemm.o' 
$ ./t605 
handle 0 n = 3 
rf 0 info 0 
ri 0 info 0 
done 
1.090909, -0.545455, -0.090909, 
0.227273, 0.136364, -0.227273, 
-0.181818, 0.090909, 0.181818, 
$ 

Answer 2

難道你運行的一些CUDA函數只支持不同的體系結構（即使文檔中提到的所有東西都是。）如果我使用-arch=sm_50進行編譯，我不會得到編譯器的警告。 sm_50設備能夠測試雖然...

此外，這些警告看起來像一些函數asm不適用於您的架構，因此它被鏈接到不同的架構asm，您的設備不支持，所以你會得到一些奇怪的錯誤。我認爲你應該接受這個nvidia開發人員誰更瞭解他們的編譯器正在做什麼。

我有權訪問Compute 3.5的設備，但不幸的是只有CUDA v 6.0和使用你的例子（略有修正，才能在第42行上編譯（const float *） - >（float *）），並且我沒有收到任何編譯警告（儘管結果相同）。

同樣如在評論中提到：

(float**)a_i 

不使A_I爲類型（浮動**）。你應該採取的地址： & a_i

更改這些並沒有幫助解決問題，但這些都是一些你可以看看探索的指針。