从设备的cublas矩阵反转

我想从设备运行矩阵反转。如果从主机调用此逻辑，则工作正常。从设备的cublas矩阵反转

编译行如下（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日编辑使用正确的指针输入。这仍然报告内核中的非法内存访问。

来源

2014-11-23 Bob

您发布的代码中的第63行是空格。代码中发生的错误究竟在哪里？ – talonmies 2014-11-23 21:42:55

设备同步期间的第64行。我必须发布和更老的输出。我怀疑在调用cublasSgetrfBatched期间。 – Bob 2014-11-23 21:47:10

'（float **）a_i'看起来很可疑。当然，你的意思是传递'a_i'的地址而不是它的值？ – talonmies 2014-11-23 22:03:21

有关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, 
$

来源

2014-11-25 17:14:43

谢谢。这对我有用。我原来分配了p和info变量，但没有意识到我还需要分配a，aconst和c变量。阅读本地内存参考部分后，这是有道理的。我会想象n被分配给全局内存堆，因为它是内核调用参数的一部分。句柄变量可能不适用。 – Bob 2014-11-25 23:45:32

其他参数像'n'，'batch'等等，都是按值传递的。通过值传递的参数没有引用回调用环境。这是C/C++的特征，不是独特的CUDA概念。事实上，即使指针也是“按价值”传递的。但是，当这些指针值在子内核中被取消引用时，会发生不好的事情。对于非指针参数，在子内核中没有这样的解引用，并且一切正常。事实上，这个按值传递实际上发生在cublas函数调用中（并且随后，在后面发生的子内核启动时）。 – 2014-11-25 23:53:23

难道你运行的一些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

更改这些并没有帮助解决问题，但这些都是一些你可以看看探索的指针。

来源

2014-11-24 10:53:41 XapaJIaMnu

对，我的坏，对不起 – XapaJIaMnu 2014-11-24 16:59:33

当我用SM_50编译时，我得到“ptxas info：'device-function-maxrregcount'是测试版功能”。另一个警告消失。 – Bob 2014-11-25 00:08:33

从设备的cublas矩阵反转

回答

相关问题