仅使用内核的符号名称作为字符串即可获取cuda内核的设备端函数指针（没有签名）

问题描述 投票：0回答：1

对于post中提到的代码，它是nvcc编译器注入的

__nv_cudaEntityRegisterCallback
函数，它将主机端内核的名称/符号
((void (*)(double*, double*, double*, int))vecAdd)
注册到设备上的符号
_Z6vecAddPdS_S_i
，以便GPU解析： 
cudaLaunchKernel
API 调用期间的一些查找表（以获取设备指针），如此处所述。

类似地，对于

__device__
函数指针定义来获取内核
vecAdd
的设备指针：

typedef void (*fp)(double *, double *, double *, int);
__device__ fp kernelPtrvecAdd = vecAdd;

甚至由 

__nv_cudaEntityRegisterCallback
函数处理。

static void __nv_cudaEntityRegisterCallback(void **__T23) { 
    // Save the fat binary handle for managed runtime
    __nv_save_fatbinhandle_for_managed_rt(__T23);
    ...
    ...
    // Register the vecAdd function
    __cudaRegisterFunction(
        __T23, 
        (const char*)((void (*)(double*, double*, double*, int))vecAdd), 
        "_Z6vecAddPdS_S_i", 
        "_Z6vecAddPdS_S_i", 
        -1, (uint3*)0, (uint3*)0, (dim3*)0, (dim3*)0, (int*)0
    );
 
     // Register the kernelPtrvecAdd variable
    __cudaRegisterVar(
        __T23, 
        (char*)&::kernelPtrvecAdd, 
        "kernelPtrvecAdd", 
        "kernelPtrvecAdd", 
        0, 8UL, 0, 0
    );
    ...
    ...
}

上面的代码片段可以通过编译

post.cu
文件这里来获得：

$ nvcc -cuda post.cu -o post.cu.cpp.ii

但是请考虑下面给出的示例中的情况：

我有以下设置： vecAdd.cu

// vecAdd.cu
#include <cuda_runtime.h>
#include <stdio.h>

// CUDA kernel that adds two vectors, each thread handles one element of c
__global__ void vecAdd(double *a, double *b, double *c, int n) {
    int id = blockIdx.x * blockDim.x + threadIdx.x;
    if (id < n) {
        c[id] = a[id] + b[id];
    }
}

编译为

$ nvcc -cubin -arch=sm_75 vecAdd.cu -o vecAdd.cubin

主.cu

#include <math.h>
#include <stdio.h>
#include <stdlib.h>
#include <cuda_runtime.h>
#include <cuda.h>

#define CUDA_SAFECALL(call)                                                 \
    {                                                                       \
        call;                                                               \
        cudaError err = cudaGetLastError();                                 \
        if (cudaSuccess != err) {                                           \
            fprintf(                                                        \
                stderr,                                                     \
                "Cuda error in function '%s' file '%s' in line %i : %s.\n", \
                #call, __FILE__, __LINE__, cudaGetErrorString(err));        \
            fflush(stderr);                                                 \
            exit(EXIT_FAILURE);                                             \
        }                                                                   \
    }

#define SAFECALL_DRV(call)                                                  \
    {                                                                       \
        CUresult err = call;                                                \
        if (err != CUDA_SUCCESS) {                                          \
            const char *errStr;                                             \
            cuGetErrorString(err, &errStr);                                 \
            fprintf(                                                        \
                stderr,                                                     \
                "CUDA Driver API error in function '%s' file '%s' in line %i : %s.\n", \
                #call, __FILE__, __LINE__, errStr);                         \
            fflush(stderr);                                                 \
            exit(EXIT_FAILURE);                                             \
        }                                                                   \
    }
    
int main(int argc, char *argv[]) {
    int n = 100000000;  // Size of the vectors
    if (argc > 1) n = atoi(argv[1]);

    // Initialize CUDA Driver API
    cuInit(0);

    // Get a CUDA device and create a context
    CUdevice device;
    CUcontext context;
    cuDeviceGet(&device, 0);
    cuCtxCreate(&context, 0, device);
    cuDevicePrimaryCtxRetain(&context, device);
    // Load the module from vecAdd.o
    CUmodule module;
    SAFECALL_DRV(cuModuleLoad(&module, "vecAdd.cubin"));

    // Create a CUDA stream for asynchronous execution
    cudaStream_t stream;
    cudaStreamCreate(&stream);

    // Host and device vectors
    double *h_a, *h_b, *h_c;
    double *d_a, *d_b, *d_c;
    size_t bytes = n * sizeof(double);

    // Allocate host memory
    h_a = (double *)malloc(bytes);
    h_b = (double *)malloc(bytes);
    h_c = (double *)malloc(bytes);

    // Initialize host vectors
    for (int i = 0; i < n; i++) {
        h_a[i] = sin(i) * sin(i);
        h_b[i] = cos(i) * cos(i);
        h_c[i] = 0;
    }

    CUfunction vecAddFunc;
    SAFECALL_DRV(cuModuleGetFunction(&vecAddFunc, module, "_Z6vecAddPdS_S_i"));
    printf("vecAdd: %p\n", vecAddFunc);

    // Allocate device memory
    cudaMallocAsync(&d_a, bytes, stream);
    cudaMallocAsync(&d_b, bytes, stream);
    cudaMallocAsync(&d_c, bytes, stream);

    // Copy data from host to device
    cudaMemcpyAsync(d_a, h_a, bytes, cudaMemcpyHostToDevice, stream);
    cudaMemcpyAsync(d_b, h_b, bytes, cudaMemcpyHostToDevice, stream);

    // Time the kernel execution
    cudaEvent_t start, stop;
    cudaEventCreate(&start);
    cudaEventCreate(&stop);

    for (int i = 0; i < 10; i++) {
        cudaEventRecord(start, stream);
        int gridSize = (int)ceil((float)n / 1024);
        void *args[] = { &d_a, &d_b, &d_c, &n };

        SAFECALL_DRV(cuLaunchKernel(
                        vecAddFunc,      // Kernel function
                        gridSize, 1, 1,  // Grid dimensions
                        1024, 1, 1,      // Block dimensions
                        0,               // Shared memory
                        stream,          // Stream
                        args,            // Kernel arguments
                        NULL             // Extra (not used)
                    ));

        cudaStreamSynchronize(stream);
        cudaEventRecord(stop, stream);
        cudaEventSynchronize(stop);

        float time = 0;
        cudaEventElapsedTime(&time, start, stop);
        printf("Iteration %d: Time vecAdd: %f ms\n", i, time);
    }

    // Copy array back to host using async memory copy
    cudaMemcpyAsync(h_c, d_c, bytes, cudaMemcpyDeviceToHost, stream);
    
    // Release device memory using async memory deallocation
    cudaFreeAsync(d_a, stream);
    cudaFreeAsync(d_b, stream);
    cudaFreeAsync(d_c, stream);

    // Synchronize the stream to ensure everything is done
    cudaStreamSynchronize(stream);

    // Sum up vector c and print result divided by n, this should equal 1 within error
    double sum = 0;
    for (int i = 0; i < n; i++) sum += h_c[i];
    printf("Final sum = %f; sum/n = %f (should be ~1)\n", sum, sum / n);

    // Clean up resources
    cudaStreamDestroy(stream);
    cudaEventDestroy(start);
    cudaEventDestroy(stop);
    free(h_a);
    free(h_b);
    free(h_c);
    cudaFree(d_a);
    cudaFree(d_b);
    cudaFree(d_c);

    // Destroy the CUDA context
    cuCtxDestroy(context);

    return 0;
}

$ nvcc  main.cu -lcuda

./a.out vecAdd: 0x56400fc49640
Iteration 0: Time vecAdd: 6.092896 ms
...
Iteration 9: Time vecAdd: 6.029056 ms
Final sum = 100000000.000000; sum/n = 1.000000 (should be ~1)

在上面的代码中，我将 

cubin
内核代码的 
vecAdd
加载到 
main.cu
文件中，然后获取 
0x56400fc49640
的主机端存根（从地址 
vecAdd
可以看出）内核使用 
cuModuleGetFunction
仅将内核函数的符号名称作为字符串传递（如 
_Z6vecAddPdS_S_i
中所示）（尽管在这种特殊情况下，管理的 cpp 样式名称中包含函数签名信息，但它可能并不总是是这种情况）使用它我可以使用 
cuLaunchKernel
.

启动内核 

nvcc -cuda main.cu main.cu.cpp.ii
输出文件（main.cu.cpp.ii）不包含
__nv_cudaEntityRegisterCallback
中的任何寄存器函数行，但是 
nvcc -cuda vecAdd.cu vecAdd.cu.cpp.ii
包含 
_Z6vecAddPdS_S_i
的寄存器函数条目。所以，我猜 
cuModuleLoad
和 
cuModuleGetFunction
正在将其安装在主可执行文件中。

我想知道，

  1. 有什么方法可以仅使用字符串
    _Z6vecAddPdS_S_i
    我就可以获取内核的设备端指针
    vecAdd
    。就像使用 
    cuModuleGetFunction
    一样，我们获得了 
    vecAdd
    内核的主机指针。
  2. 或者，如果使用 
    vecAdd.cu
    将 
    vecAdd.o
    编译为 
    nvcc -c vecAdd.cu -o vecAdd.o -rdc=true
    并且 main 编译为 
    nvcc main.cu vecAdd.o
    ，则定义 
    __device__
    函数指针需要我们知道 
    vecAdd
    的签名（用于外部链接） 
    main.cu
    ：
typedef void (*fp)(double *, double *, double *, int);
extern __global__ void vecAdd(double *a, double *b, double *c, int n);
__device__ fp kernelPtr = vecAdd;

有没有办法，我可以在不知道其签名的情况下获取 

vecAdd
的设备函数指针？

c++ compilation cuda function-pointers nvidia
1个回答
0
投票

据我所知，你永远不会获得内核的设备端函数指针。

您可能会得到什么，例如使用 cuModuleGetFunction 是一个 

CUfunction
，其[定义][1]为：

typedef CUfunc_st * CUfunction;

因此，指针不是指向函数，也不是指向函数存根，而是指向 CUDA 驱动程序使用的不透明结构。

对于与可执行文件一起“静态”编译并与其一起加载的函数，确实存在您使用或调用的主机端存根。 AFAICT，没有直接的机制可以通过（损坏的）名称来获取 CUfunction 的名称，但说实话，我并没有试图努力得到它。

但是，您可以使用 CUDA 版本 11.0 或更高版本从存根转到 CU 函数：

extern __host__ cudaError_t cudaGetFuncBySymbol(cudaFunction_t* functionPtr, const void* symbolPtr);

并且 

cudaFunction_t
实际上只是像 
CUfunc_st *
那样的 
CUfunction
。

最新问题
© www.soinside.com 2019 - 2024. All rights reserved.