CUDA C 指针给出令人困惑的值

Question

好吧，如果这与之前看到的类似，我们很抱歉。我有以下代码:

//kern.cu
#include <stdio.h>
#include <dlfcn.h>
#include <unistd.h>
extern "C"{
#include "Kernalize.h"
#include <stdio.h>
}

extern "C" {
    __device__ void *dat;
    __global__ void memManageDevice(void *data){
        dat=data;
    }
    void memManageD(void *data){
        printf("A%d",data);
        void *d;
        printf("B%d",d);
        cudaMemcpy(d,&data,sizeof(data),cudaMemcpyHostToDevice);
        memManageDevice<<<1,1>>>(data);
    }
    __global__ void MemManageC(void *r){//don't call this unless in this file.
        r=dat;
    }
    void* memManageH(void *s){
        printf("C%d",s);
        void *dr;
        cudaMalloc((void **)&dr, sizeof(s));
        void *hr;
        int size=sizeof(s);
        MemManageC<<<1,1>>>(dr);
        cudaMemcpy(&hr, dr, size, cudaMemcpyDeviceToHost);
        printf("D%d",hr);
        return hr;
    }
    __global__ void kernalize(void (*ptr)(void *)) {
        (*ptr)(dat);
    }
    void Start(int d1, int d2, void (*ptr)(void *), void *data) {//TODO: make arrays as to start many  kernels
        int size=sizeof(data);
        // void *ddata;
        // bool ab=true;
        // bool *coolbeans=&ab;
        // memManageD<<<1,1>>>(data);
        kernalize<<<d1,d2,d2*size>>>(ptr);
        // data=sdata;
        // coolbeans=false;
        //kernalize(ptr,data);
    }
}

我将其编译成 .so:

nvcc --ptxas-options=-v --compiler-options '-fPIC' -o libpar.so --shared kern.cu

然后从普通的 C 中引用它:

typedef void (*gFunc) ();
typedef void (*sFunc) (int,int,gFunc*,void *data);

typedef void* (*hFunc) (void *);
typedef void (*dFunc) (void *);
void toBe(void *data){
        data=12;
        while(1){}//side-expirement, don't think it's the stem of the issue.
}
int main() {
    printf("start");
    sFunc fS;
    hFunc hS;
    dFunc dS;
    void* hLibrary = dlopen("./libpar.so", RTLD_NOW | RTLD_GLOBAL);
    if(hLibrary == NULL) {
        fprintf(stderr, "%s\n", dlerror());
        return 1;
    }
    int i=42;
    *(void**)(&dS)=dlsym(hLibrary,"memManageD");
    (void) dS(i);
    sleep(1);
    printf("checkpoint");
    *(void**)(&fS)=dlsym(hLibrary,"Start");
    (void) fS(2,2,toBe,&i);
    sleep(1);
    *(void**)(&hS)=dlsym(hLibrary,"memManageH");
    int x=(void*) hS(&i);
    printf("%d", x);

    return 0;
}

正如您可能能够通过我的极其丑陋的代码看出的那样，函数 toBe 被传递到 CUDA C 内核，在执行时预计会将非类型变量指针“data”更改为 12。“data”是对普通 c 中“i”的引用，以 42 开头。不幸的是，我的输出是 1，而不是 12:

 startA42B431891052checkpointC-288453328D1

这实际上只是“A42”和“D1”之间的垃圾内存列表。我对 CUDA C 和 C 都比较陌生。 (我大部分时间都花在高级编程语言上。)所以问题实际上是我在哪里犯了一个愚蠢的错误，无论是我对 CUDA 的理解、我的 C 语法，还是我对我的设想的整体看法。

Answer 1

这是我对为什么会看到输出的理解:

• A42 正在按预期工作，因为 i 等于 42

• B431891052 因为指针 void* d 未初始化，并且可能包含垃圾值，而打印语句 printf("B %d",d) 仅打印其低 32 位。

• C-288453328 因为语句 printf("C%d",s) 打印变量 i 的低 32 位>。在32位中，数字-288453328对应于十六进制数0xEECE8D30；我认为是堆栈地址0x7FFFEECE8D30的低32位。

• D1 因为 hr 将包含 dr 指向的前 8 个字节。它尚未初始化，在本例中它指向的值恰好是 1。

我认为要分配D12函数memManageC可能需要更改为以下内容:

__global__ void MemManageC(int *r) {
    *r = dat;
}

因此，r指向的内容将具有dat的值。如果没有取消引用，该函数只会修改传入的参数，这是一个无操作，并且可能会被编译器优化。

要拥有 D12 还需要 toBe 正确设置 dat 的值，这似乎需要更多代码更改。 toBe 是一个主机端函数，在 __global__ 或 __kernel__ 函数中主机端函数的用法可以在这个答案中找到:Passing Host Function as a function pointer in __global__ OR __device__ function in CUDA