CUBLAS内核中不工作编译一次使用-G标志与NVCC

Question

的cubin我有个CUDA核心，看起来像下面这样：

#include <cublas_v2.h> 
#include <math_constants.h> 
#include <stdio.h> 


extern "C" { 

    __device__ float ONE = 1.0f; 
    __device__ float M_ONE = -1.0f; 
    __device__ float ZERO = 0.0f; 

    __global__ void kernel(float *W, float *input, int i, float *output, int o) { 
     int idx = blockIdx.x*blockDim.x+threadIdx.x; 
     cublasHandle_t cnpHandle; 

     if(idx == 0) { 
      cublasCreate(&cnpHandle); 
      cublasStatus_t s = cublasSgemv(cnpHandle, CUBLAS_OP_N, o, i, &ONE, W, o, input, 1, &ZERO, output, 1); 
      printf("status %d\n", s); 
      cudaError_t e = cudaDeviceSynchronize(); 
      printf("sync %d\n", e); 
     } 

    } 

} 

主机代码：

#include <iostream> 
#include <numeric> 
#include <stdlib.h> 
#include <cstring> 
#include <cuda_runtime.h> 
#include <cuda.h> 
#include <cublas_v2.h> 

extern "C" { 
    __global__ void kernel(float *W, float *input, int i, float *output, int o); 
} 

#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); } 
inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort=true) 
{ 
    if (code != cudaSuccess) 
    { 
     fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line); 
     if (abort) exit(code); 
    } 
} 

int main(int argc, char* argv[]) 
{ 

    cuInit(0); 
    CUcontext pctx; 
    CUdevice dev; 
    cuDeviceGet(&dev, 0); 
    cuCtxCreate(&pctx, 0, dev); 

    CUmodule module; 
    CUresult t = cuModuleLoad(&module, "pathto/src/minimalKernel.cubin"); 

    CUfunction function; 
    CUresult r = cuModuleGetFunction(&function, module, "kernel"); 

    float *W = new float[2]; 
    W[0] = 0.1f; 
    W[1] = 0.1f; 
    float *input = new float[2]; 
    input[0] = 0.1f; 
    input[1] = 0.1f; 
    float *out = new float[1]; 
    out[0] = 0.0f; 

    int i = 2; 
    int o = 1; 

    float *d_W; 
    float *d_input; 
    float *d_out; 
    cudaMalloc((void**)&d_W, 2*sizeof(float)); 
    cudaMalloc((void**)&d_input, 2*sizeof(float)); 
    cudaMalloc((void**)&d_out, sizeof(float)); 
    cudaMemcpy(d_W, W, 2*sizeof(float), cudaMemcpyHostToDevice); 
    cudaMemcpy(d_input, input, 2*sizeof(float), cudaMemcpyHostToDevice); 
    cudaMemcpy(d_out, out, sizeof(float), cudaMemcpyHostToDevice); 
    //kernel<<<1, 2>>>(d_W, d_input, i, d_out, o); 

    //cudaMemcpy(out, d_out, sizeof(float), cudaMemcpyDeviceToHost); 

    //std::cout<<"out:"<<out[0]<<std::endl; 

    void * kernelParams[] { &d_W, &d_input, &i, &d_out, &o }; 

    CUresult k = cuLaunchKernel(function, 1, 1, 1, 2, 1, 1, 0, 0, (void**)kernelParams, 0); 
    gpuErrchk(cudaPeekAtLastError()); 
    gpuErrchk(cudaDeviceSynchronize()); 
    cudaMemcpy(out, d_out, sizeof(float), cudaMemcpyDeviceToHost); 

    std::cout<<"out:"<<out[0]<<std::endl; 


} 

当这个内核运行内嵌kernel<<<1,2>>>()，构建并链接（在eclipse Nsight中），内核运行完全正常，并且out返回0.02 如预期的那样。

如果我在编译内核到使用-G（生成设备调试符号）一.cubin的CUBLAS函数从不运行，而out总是0.0

我可以把断点时.cubin正在运行和我可以看到进入cublas函数的数据是正确的，但它看起来像cublas函数永远不会运行。 cublas函数也总是返回0 CUDA_SUCCESS。更重要的这从.cubin

运行这个时候要编译到我与-G使用的cubin只发生：

nvcc -G -cubin -arch=sm_52 --device-c kernel.cu -o kernel.cubin -dlink -lcudadevrt -lcublas_device 

返回没有错误。

如果添加-G选项，为什么cubc在.cubin中的功能停止工作？

CUDA 7.0 的Linux 14.04 64 980GTX

Answer 1

FWIW，您的代码不正确，我带或不带-G开关运行。您可以使用cuda-memcheck运行您的代码以帮助识别错误。 （您的主机代码或设备代码中似乎没有执行proper CUDA error checking。通过动态并行机制，您可以在设备代码中使用类似的方法，并且CUBLAS API调用会返回错误代码，但您不会看到。被检查）

这是错误的：

if(idx == 0) { 
     cublasCreate(&cnpHandle); 
    } 

这是一个线程局部变量：

cublasHandle_t cnpHandle; 

既然你推出一个内核2个线程：

CUresult k = cuLaunchKernel(function, 1, 1, 1, 2, 1, 1, 0, 0, (void**)kernelParams, 0); 

你的一个线程（0）被传递有效句柄cublasSgemv调用，其他线程（1）不是。

当我修复该错误时，您的代码对我“起作用”。请注意，您仍然存在这样的情况，即您将完全相同的参数传递给您的两个线程中的每个线程的cublasSgemv调用。因此，每个调用都将写入相同的输出位置。由于在这种情况下线程执行/行为的顺序是未指定的，因此可能会看到相当可变的行为：似乎获得有效的输出（因为一个线程作为成功的cublas调用写入了正确的值），即使其他cublas通话失败。我猜想，可能的是，-G开关可能会影响此排序，或者以某种方式影响此行为。

$ cat t889_kern.cu 
#include <cublas_v2.h> 
#include <math_constants.h> 
#include <stdio.h> 


extern "C" { 

    __device__ float ONE = 1.0f; 
    __device__ float M_ONE = -1.0f; 
    __device__ float ZERO = 0.0f; 

    __global__ void kernel(float *W, float *input, int i, float *output, int o) { 
//  int idx = blockIdx.x*blockDim.x+threadIdx.x; 
     cublasHandle_t cnpHandle; 

     cublasCreate(&cnpHandle); 

     cublasSgemv(cnpHandle, CUBLAS_OP_N, o, i, &ONE, W, o, input, 1, &ZERO, output, 1); 
     cudaDeviceSynchronize(); 
    } 

} 
$ cat t889.cpp 
#include <iostream> 
#include <numeric> 
#include <stdlib.h> 
#include <cstring> 
#include <cuda_runtime.h> 
#include <cuda.h> 
#include <cublas_v2.h> 

extern "C" { 
    __global__ void kernel(float *W, float *input, int i, float *output, int o); 
} 

int main(int argc, char* argv[]) 
{ 

    cuInit(0); 
    CUcontext pctx; 
    CUdevice dev; 
    cuDeviceGet(&dev, 0); 
    cuCtxCreate(&pctx, 0, dev); 

    CUmodule module; 
    CUresult t = cuModuleLoad(&module, "kernel.cubin"); 

    CUfunction function; 
    CUresult r = cuModuleGetFunction(&function, module, "kernel"); 

    float *W = new float[2]; 
    W[0] = 0.1f; 
    W[1] = 0.1f; 
    float *input = new float[2]; 
    input[0] = 0.1f; 
    input[1] = 0.1f; 
    float *out = new float[1]; 
    out[0] = 0.0f; 

    int i = 2; 
    int o = 1; 

    float *d_W; 
    float *d_input; 
    float *d_out; 
    cudaMalloc((void**)&d_W, 2*sizeof(float)); 
    cudaMalloc((void**)&d_input, 2*sizeof(float)); 
    cudaMalloc((void**)&d_out, sizeof(float)); 
    cudaMemcpy(d_W, W, 2*sizeof(float), cudaMemcpyHostToDevice); 
    cudaMemcpy(d_input, input, 2*sizeof(float), cudaMemcpyHostToDevice); 
    cudaMemcpy(d_out, out, sizeof(float), cudaMemcpyHostToDevice); 
    //kernel<<<1, 2>>>(d_W, d_input, i, d_out, o); 

    //cudaMemcpy(out, d_out, sizeof(float), cudaMemcpyDeviceToHost); 

    //std::cout<<"out:"<<out[0]<<std::endl; 

    void * kernelParams[] { &d_W, &d_input, &i, &d_out, &o }; 

    CUresult k = cuLaunchKernel(function, 1, 1, 1, 2, 1, 1, 0, 0, (void**)kernelParams, 0); 

    cudaMemcpy(out, d_out, sizeof(float), cudaMemcpyDeviceToHost); 

    std::cout<<"out:"<<out[0]<<std::endl; 


} 
$ nvcc -cubin -arch=sm_35 --device-c t889_kern.cu -o kernel.cubin -dlink -lcudadevrt -lcublas_device 
ptxas info : 'device-function-maxrregcount' is a BETA feature 
$ g++ -std=c++11 -I/usr/local/cuda/include t889.cpp -o t889 -L/usr/local/cuda/lib64 -lcuda -lcudart 
$ CUDA_VISIBLE_DEVICES="1" cuda-memcheck ./t889 
========= CUDA-MEMCHECK 
out:0.02 
========= ERROR SUMMARY: 0 errors 
$