Получите указатель функции на стороне устройства ядра cuda (без подписи), используя только имя символа ядра, поскольку с ⇐ C++

[Anonymous]

Примечание. Если сообщение кажется длинным, можно сразу перейти к разделу, начинающемуся с «Мне было интересно..» в конце, на случай, если вы захотите пропустить сборку/контекст.

Сборка/контекст:
Для кода, упомянутого в сообщении, это — это функция __nv_cudaEntityRegisterCallback, внедренная компилятором nvcc, которая регистрирует имя/символ ядра на стороне хоста ((void (*)(double*, double*, double*, int))vecAdd) с помощью символ

Код: Выделить всё

_Z6vecAddPdS_S_i

на устройстве для разрешения графическим процессором с помощью некоторой таблицы поиска во время вызова API cudaLaunchKernel (чтобы получить указатель устройства), как указано здесь.
Аналогично, для определения указателя функции __device__, чтобы получить указатель устройства ядра vecAdd as:

Код: Выделить всё

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 
#include 

// 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

main.cu

Код: Выделить всё

#include 
#include 
#include 
#include 
#include 

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

В приведенном выше коде я загружаю кубин кода ядра 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 устанавливают это в основной исполняемый файл.

Мне интересно

Мне интересно

Мне интересно

Strong>,

• Есть ли способ, с помощью которого, используя только строку _Z6vecAddPdS_S_i, я мог бы получить указатель на стороне устройства ядра vecAdd< /код>. Точно так же, как при использовании cuModuleGetFunction, мы получаем указатель хоста ядра vecAdd.
• Или, если vecAdd.cu был скомпилирован как vecAdd.o с помощью nvcc -c vecAdd.cu -o vecAdd.o -rdc=true и main скомпилировать как nvcc main.cu vecAdd.o, чтобы определить

  Код: Выделить всё

  __device__

  указатель функции требует, чтобы мы знали подпись vecAdd (для внешней связи) в main.cu:

< pre class="lang-cuda beautifulprint-override">

Код: Выделить всё

typedef void (*fp)(double *, double *, double *, int);
extern __global__ void vecAdd(double *a, double *b, double *c, int n);
__device__ fp kernelPtr = vecAdd;

Есть ли выход, при котором я могу получить указатель функции устройства vecAdd без знания его сигнатуры (по линии динамического параллелизма cuda)?

Подробнее здесь: https://stackoverflow.com/questions/790 ... e-using-ju