CUDA 5.0: CUBIN и CUBLAS_device, возможность вычисления 3.5

Question

Я пытаюсь скомпилировать ядро, которое использует динамический параллелизм для запуска CUBLAS в кубиновый файл. Когда я пытаюсь скомпилировать код, используя команду

nvcc -cubin -m64 -lcudadevrt -lcublas_device -gencode arch=compute_35,code=sm_35 -o test.cubin -c test.cu 

Я получаю ptxas fatal : Unresolved extern function 'cublasCreate_v2

Если добавить опцию -rdc=true компиляции он компилирует нормально, но когда я пытаюсь загрузить модуль с помощью cuModuleLoad я получаю сообщение об ошибке 500 : CUDA_ERROR_NOT_FOUND. От cuda.h:

/** 
* This indicates that a named symbol was not found. Examples of symbols 
* are global/constant variable names, texture names, and surface names. 
*/ 
CUDA_ERROR_NOT_FOUND      = 500, 

Код ядра:

#include <stdio.h> 
#include <cublas_v2.h> 
extern "C" { 
__global__ void a() { 
    cublasHandle_t cb_handle = NULL; 
    cudaStream_t stream; 
    if(threadIdx.x == 0) { 
     cublasStatus_t status = cublasCreate_v2(&cb_handle); 
     cublasSetPointerMode_v2(cb_handle, CUBLAS_POINTER_MODE_HOST); 
     if (status != CUBLAS_STATUS_SUCCESS) { 
      return; 
     } 
     cudaStreamCreateWithFlags(&stream, cudaStreamNonBlocking); 
     cublasSetStream_v2(cb_handle, stream); 
    } 
    __syncthreads(); 
    int jp; 
    double A[3]; 
    A[0] = 4.0f; 
    A[1] = 5.0f; 
    A[2] = 6.0f; 
    cublasIdamax_v2(cb_handle, 3, A, 1, &jp); 
} 
} 

Примечание: Объем A является локальным, так что данные на указатель на данный cublasIdamax_v2 не определен, и поэтому jp концы как более или менее случайное значение в этом коде. Правильный способ сделать это - иметь A в глобальной памяти.

код Ведущий:

#include <stdio.h> 
#include <cuda.h> 
#include <cuda_runtime_api.h> 

int main() { 
    CUresult error; 
    CUdevice cuDevice; 
    CUcontext cuContext; 
    CUmodule cuModule; 
    CUfunction testkernel; 
    // Initialize 
    error = cuInit(0); 
    if (error != CUDA_SUCCESS) printf("ERROR: cuInit, %i\n", error); 
    error = cuDeviceGet(&cuDevice, 0); 
    if (error != CUDA_SUCCESS) printf("ERROR: cuInit, %i\n", error); 
    error = cuCtxCreate(&cuContext, 0, cuDevice); 
    if (error != CUDA_SUCCESS) printf("ERROR: cuCtxCreate, %i\n", error); 
    error = cuModuleLoad(&cuModule, "test.cubin"); 
    if (error != CUDA_SUCCESS) printf("ERROR: cuModuleLoad, %i\n", error); 
    error = cuModuleGetFunction(&testkernel, cuModule, "a"); 
    if (error != CUDA_SUCCESS) printf("ERROR: cuModuleGetFunction, %i\n", error); 
    return 0; 
} 

код хоста составлен с использованием nvcc -lcuda test.cpp. Если я заменил ядро ​​простым ядром (ниже) и скомпилировал его без -rdc=true, он отлично работает.

Простой рабочий ядро ​​

#include <stdio.h> 
extern "C" { 
__global__ void a() { 
    printf("hello\n"); 
} 
} 

Заранее спасибо

• Сорен

Answer 1

Вы просто не хватает -dlink в первом подходе:

nvcc -cubin -m64 -lcudadevrt -lcublas_device -gencode arch=compute_35,code=sm_35 -o test.cubin -c test.cu -dlink 

Вы также можете сделать это в два этапа:

nvcc -m64 test.cu -gencode arch=compute_35,code=sm_35 -o test.o -dc 
nvcc -dlink test.o -arch sm_35 -lcublas_device -lcudadevrt -cubin -o test.cubin