CUDA 5.0: CUBIN y CUBLAS_dispositivo, capacidad de cálculo 3.5

nvcc

Question

Estoy intentando compilar un núcleo que usa paralelismo dinámico para ejecutar CUBLAS en un archivo cubin. Cuando intento compilar el código usando el comando

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

Obtengo ptxas fatal : Unresolved extern function ''cublasCreate_v2

Si agrego la opción de compilación -rdc=true compila bien, pero cuando intento cargar el módulo usando cuModuleLoad obtengo el error 500: CUDA_ERROR_NOT_FOUND. Desde 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,

El código del kernel:

#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 );
}
}

NOTA: El alcance de A es local, por lo que los datos en el puntero dados a cublasIdamax_v2 no están definidos, por lo que jp termina como un valor más o menos aleatorio en este código. La forma correcta de hacerlo sería tener A en la memoria global.

Código de host:

#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;
}

El código de host se compila utilizando nvcc -lcuda test.cpp . Si reemplazo el núcleo con un núcleo simple (abajo) y lo -rdc=true sin -rdc=true , funciona bien.

Núcleo de trabajo simple

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

Gracias por adelantado

• Soren

Answer 1

Solo te falta -dlink en tu primer acercamiento:

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

También puedes hacer eso en dos pasos:

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