使用带有 nvcc 的 -G 标志编译为 cubin 后,Cublas 无法在内核中工作
Posted
技术标签:
【中文标题】使用带有 nvcc 的 -G 标志编译为 cubin 后,Cublas 无法在内核中工作【英文标题】:Cublas not working within kernel once compiled to cubin using -G flag with nvcc 【发布时间】:2015-08-19 09:50:15 【问题描述】:我有一个如下所示的 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
选项,为什么.cubin 中的cublas 函数会停止工作?
CUDA 7.0 linux 14.04 x64 980GTX
【问题讨论】:
提供一个 MCVE,包括你用来加载和调用内核的代码。 已在上面编辑以提供 MCVE 【参考方案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) 没有。
当我修复该错误时,您的代码对我“有效”。请注意,您仍然存在将 exact same 参数传递给两个线程中的每个线程的 cublasSgemv
调用的情况。因此,每个调用都写入相同的输出位置。由于在这种情况下未指定线程执行/行为的顺序,因此您可能会看到相当多变的行为:即使其他 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
$
【讨论】:
我稍微修改了我的代码,只在第 0 个线程中运行 cublas。运行 cuda-memcheck 我得到这个: cuda-memcheck ./example ========= CUDA-MEMCHECK out:0 ========= 错误摘要:0 错误所以这仍然不起作用我..有什么想法吗? 没有。你修改后的代码对我来说可以正常工作(它显示out:0.02
),有或没有cuda-memcheck
,有或没有-G
。您使用的是哪个 CUDA 版本?你是在windows还是linux上运行?什么显卡?您可能需要添加一些代码来检查内核中 cublas 调用的返回值。
我刚刚完成并在此处编辑了代码,就我所见,代码还可以...运行 ubuntu 14.04、cuda toolkit 7.0 和 980gtx。 nvidia-346 驱动程序。我没有将它用作我的显示设备(如果这有什么不同的话)
一点更新,我用一个非常基本的cublasSdot
替换了cublasSgemv
,它似乎根本没有问题。使用-G
运行良好而不是以上是关于使用带有 nvcc 的 -G 标志编译为 cubin 后,Cublas 无法在内核中工作的主要内容,如果未能解决你的问题,请参考以下文章
CUDA compiler driver nvcc 散点 part 2