Cublas не работает в ядре один раз компилируется cubin используя флаг -g с NVCC

Question

У меня есть ядро ​​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; 


} 

, когда это ядро работает inline kernel<<<1,2>>>(), построенный и связанный (в пределах Eclipse Nsight), ядро ​​работает полностью нормально и out возвращает 0.02 как и ожидалось.

Если я скомпилировать ядро ​​в .cubin с помощью -G (генерировать отладочные устройства символов), функция cublas никогда не работает, а out всегда 0.0

Я могу поставить точки останова в том, когда .cubin работает и Я вижу, что данные верны в функции cublas, но похоже, что функция cublas никогда не запускается вообще. Функция cublas также всегда возвращает 0 CUDA_SUCCESS. Важно это происходит только при выполнении этого из .cubin

Для компиляции в cubin я использую с -G:

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

, который не возвращает никаких ошибок.

Почему функции cublas в пределах .cubin перестают работать, если добавлен параметр -G?

CUDA 7.0 линукс 14,04 x64 980GTX

Answer 1

FWIW, ваш код не работает правильно для меня с или без переключателя -G. Вы можете запустить свой код с помощью cuda-memcheck, чтобы помочь определить ошибки. (Кажется, что вы не делаете proper CUDA error checking, либо в коде хоста, либо в коде вашего устройства. С динамическим параллелизмом вы можете использовать аналогичную методологию в коде устройства. И API CUBLAS вызывает коды ошибок возврата, которые вы не видите . проверять)

Это неправильно:

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 
$