Можем ли мы получить имя функции ядра cuda в cudaLaunchKernel?

Question

Я пытался вставить некоторые коды в cudaLaunchKernel и мне нужно сохранить имя его функции, но я не могу найти прямой API, который может помочь мне получить имя функции ядра. Я рассматривал CUPTI, но он использует функцию обратного вызова для получения информации, поэтому я не могу изменить поведение при запуске ядра (или мне нужно интенсивное межпроцессное взаимодействие, что некрасиво ...)

Есть ли какие-либо Как я могу получить имя функции в cudaLaunchKernel (может быть, по указателю на функцию?)?

Пример выглядит следующим образом.

cudaKernelLaunch(...) {
    kernel_id = getKernelNameBySomeMethods(); // it's what I want..
    send_to_other_processes(kernel_name);
    return ::cudaKernelLaunch(...);
}

// for other process
receive_kernel_name_from_other_process;
store_it;

Изменить: идентификатор также подходит. Я могу отправить идентификатор в другой процесс для сохранения, поэтому мне нужно классифицировать разные ядра cuda. ​​

Answer 1

Для этого нет API, ни publi c, ни private AFAIK. Компилятор генерирует множество stati c шаблонов на стороне хоста для выполнения API среды выполнения magi c, которое мы считаем само собой разумеющимся, это не выполняется самой библиотекой времени выполнения.

Однако природа этого Шаблон означает, что вы можете довольно легко создать свою собственную таблицу поиска - некоторые взломы во время обеденного перерыва дали мне это частичное доказательство концепции, которая делает то, что, я думаю, вы хотите:

#include <cstdio>
#include <map>
#include <string>
#include <iostream>

__global__ void kernel_1(float *in, float *out, int N)
{
    int tidx = threadIdx.x + blockDim.x * blockIdx.x;
    if (tidx == 0) printf("Running kernel_1\n");
    if (tidx < N) out[tidx] = in[tidx];
}


__global__ void kernel_2(float *in, float *out, int N)
{
    int tidx = threadIdx.x + blockDim.x * blockIdx.x;
    if (tidx == 0) printf("Running kernel_2\n");
    if (tidx < N) out[tidx] = 2.f * in[tidx];
}

__global__ void kernel_3(float *in, float *out, int N)
{
    int tidx = threadIdx.x + blockDim.x * blockIdx.x;
    if (tidx == 0) printf("Running kernel_3\n");
    if (tidx < N) out[tidx] = 3.f * in[tidx];
}


void notakernel(float *in, float *out, int N)
{
   printf("Someone bad happened\n");
}

std::map <void*, std::string> ktable = {
    { (void*)kernel_1, "kernel_1" },
    { (void*)kernel_2, "kernel_2" },
    { (void*)kernel_3, "kernel_3" } };


cudaError_t MyLaunchKernel (void* func, dim3 gridDim, dim3 blockDim, void** args, size_t sharedMem, cudaStream_t stream)
{
    auto it = ktable.find(func);
    if (it != ktable.end()) {
        std::cout << "Received request to call " << it->second << std::endl;
    } else {
        std::cout << "Received request to call unknown function!" << std::endl;
    }

    return cudaLaunchKernel(func, gridDim, blockDim, args, sharedMem, stream);
}

int main()
{

    int N = 100;
    float* a; cudaMalloc<float>(&a, N * sizeof(float));  
    float* b; cudaMalloc<float>(&b, N * sizeof(float));  
    void* args[] = { (void*)&a, (void*)&b, (void*)&N };

    MyLaunchKernel((void*)kernel_1, dim3(1), dim3(1), args, 0, NULL);
    cudaDeviceSynchronize();

    MyLaunchKernel((void*)kernel_2, dim3(1), dim3(1), args, 0, NULL);
    cudaDeviceSynchronize();

    MyLaunchKernel((void*)kernel_3, dim3(1), dim3(1), args, 0, NULL);
    cudaDeviceSynchronize();

    MyLaunchKernel((void*)notakernel, dim3(1), dim3(1), args, 0, NULL);
    cudaDeviceSynchronize();

    return 0;
}

, который, похоже, работает:

$ nvcc -std=c++11 -arch=sm_52  -o lookup lookup.cu
$ cuda-memcheck ./lookup
========= CUDA-MEMCHECK
Received request to call kernel_1
Running kernel_1
Received request to call kernel_2
Running kernel_2
Received request to call kernel_3
Running kernel_3
Received request to call unknown function!
========= Program hit cudaErrorInvalidDeviceFunction (error 98) due to "invalid device function" on CUDA API call to cudaLaunchKernel. 
=========     Saved host backtrace up to driver entry point at error
=========     Host Frame:/usr/lib/x86_64-linux-gnu/libcuda.so.1 [0x3b9803]
=========     Host Frame:./lookup [0x4ca95]
=========     Host Frame:./lookup [0x746c]
=========     Host Frame:./lookup [0x769f]
=========     Host Frame:/lib/x86_64-linux-gnu/libc.so.6 (__libc_start_main + 0xe7) [0x21b97]
=========     Host Frame:./lookup [0x722a]
=========
========= ERROR SUMMARY: 1 error

Очевидно, что в полной реализации для вашего варианта использования все должно быть немного сложнее - вам потребуется реализация обратного поиска для другого, вызываемого go от имени / идентификатора до указателя, и если у вас есть несколько исходных файлов, скомпилированных отдельно, вам потребуется вызов конкатенации списка для построения рабочего списка во время выполнения. Но важно помнить, что указатели на функции, которые вы передаете, на самом деле являются указателями хоста, а не указателями устройств (благодаря API среды выполнения magi c), поэтому стоимость и сложность настройки среды выполнения тривиальны, когда вы можете использовать предварительно запеченные Контейнеры стандартной библиотеки C ++, алгоритмы и адаптеры функций, которые берут на себя большую часть тяжелой работы.