Блог

Ситуация: я пытаюсь использовать cuModuleLoad для загрузки текущего двоичного файла (ELF), встроенного в cubin (и PTX), но он продолжает выдавать ошибку с кодом ошибки 200. Мой вопрос в том, если cubin встроен в конечный двоичный файл, почему я не могу использовать cuModuleLoad для динамической самостоятельной загрузки? Это работает, когда я компилирую отдельный fatbinary, но не при загрузке отдельного модуля PTX и, конечно, когда я пытаюсь загрузить конечный двоичный файл (a.out ). У меня есть несколько причин, по которым я хочу загрузить текущий исполняемый файл, от которых я откажусь, чтобы не отклоняться от темы. Я также ищу обходной путь, который поддерживает один файл без использования утилит (или системных вызовов).

В Linux:

 #include "cuda.h"
#include <cstdio>
#include <iostream>

using clock_value_t = long long;

__device__ void test(  )
{
  printf("Testing... : n");
}

__device__ void sleep(clock_value_t sleep_cycles)
{
    clock_value_t start = clock64();
    clock_value_t cycles_elapsed;
    do { cycles_elapsed = clock64() - start; }
    while (cycles_elapsed < sleep_cycles);
}

extern "C" __global__ void hello_world(  )
{
  printf("Hello World from Devicen");
  sleep( 1e9 );
  test();
}


int main(int argc, char * argv[])
{

  std::cout << argv[0] << std::endl;

  // Initialize input vectors    ...
  //Initialize
  cuInit(0);
  // Get number of devices supporting CUDA
  int deviceCount = 0;
  cuDeviceGetCount(amp;deviceCount);

  if (deviceCount == 0)
  {
    printf("There is no device supporting CUDA.n");
    exit (0);
  }
  else std::cout << "Number of device is "<< deviceCount << std::endl;

  // Get handle for device 0
  CUdevice cuDevice;
  cuDeviceGet(amp;cuDevice, 0);
  // Create context
  CUcontext cuContext;
  int ret = cuCtxCreate(amp;cuContext, 0, cuDevice);

  if( ret != CUDA_SUCCESS )
          std::cout << "Could not create context on device 0" << std::endl;

  // Create module from binary file
  CUmodule cuModule;

  ret = cuModuleLoad(amp;cuModule, argv[0]); // <---errors HERE
  
  if( ret != CUDA_SUCCESS )
  {
    std::cout << "Failed to load self fatbin : " << argv[0] << " : " << ret<< std::endl;
    return -1;
  }
}
  

Я был бы разочарован, если бы мне пришлось использовать отдельный файл или утилиту для динамического извлечения cubins или PTXs. В любом случае — спасибо вам, ребята, за ваше понимание заранее.

Найдено решение. В двух словах :

1. fopen(argv[0])
2. mmap (файл)
3. Прочитайте заголовки ELF и найдите раздел «.nv_fatbin»
4. Проанализируйте «.nv_fatbin», выровняв по последовательности байтов «50 ed 55 ba 01 00 10 00»
5. Найдите cubin, связанный с глобальным методом, который вы хотите cumodulegetфункцию
6. Вызовите cuModuleLoadFatBinary с базовым адресом .nv_fatbin определенное смещение cubin.
7. Получить функцию с помощью cuModuleGetFunction
8. Наконец, вызовите cuLaunchKernel

Смотрите Неаккуратный код ниже для справки:

 int main(int argc, char * argv[])
{
  std::cout << "Hello World from Host" << std::endl;
  std::cout << argv[0] << std::endl;
  void * start_ptr =NULL;
  struct stat sb;
  size_t sz =0;

  //read_elf_header( argv[0] );
  // Either Elf64_Ehdr or Elf32_Ehdr depending on architecture.
  ElfW(Ehdr) elf_header;
  ElfW(Shdr) header;

  std::cout << "opening elf file" << std::endl;
  FILE* file = fopen(argv[0], "rb");

  int fd = fileno( file );

  if (fd < 0)
  {
    printf("Could not open file for memory mapping, fd = %in", errno);
    exit(1);
  }

  std::cout << "getting file size" << std::endl;
  if (fstat(fd, amp;sb) == -1)          // To obtain file size
    printf("Could not find fstat");
  sz = sb.st_size;

  std::cout << "Mapping file to memory : " << sz << std::endl;
  start_ptr = mmap(NULL, sz, PROT_READ | PROT_WRITE, MAP_PRIVATE, fd, 0);

  //check if valid elf
  bool b = elf_is_elf64( file );
  fseek( file, 0, SEEK_SET );
  std::cout << "is ELF file : " << b << std::endl;
  if( b)
  {
    std::cout << "Found valid ELF file" << std::endl;
    //get ELF_Header
    b = elf64_get_elf_header(file, amp;elf_header);
    fseek( file, 0, SEEK_SET );

    if( b )
    {
      std::cout << "-Found valid ELF Header" << std::endl;
      b = elf64_get_section_header_by_name(file, (const Elf64_Ehdr *) amp;elf_header, ".nv_fatbin", amp;header);
      fseek( file, 0, SEEK_SET );

      if( b )
      {
        std::cout << "Found fatbin section" << std::endl;
        cuInit(0);
        // Get number of devices supporting CUDA
        int deviceCount = 0;
        cuDeviceGetCount(amp;deviceCount);

        if (deviceCount == 0)
        {
          printf("There is no device supporting CUDA.n");
          exit (0);
        }
        else std::cout << "Number of device is "<< deviceCount << std::endl;

        // Get handle for device 0
        CUdevice cuDevice;
        cuDeviceGet(amp;cuDevice, 0);
        // Create context
        CUcontext cuContext;
        int ret = cuCtxCreate(amp;cuContext, 0, cuDevice);
        if( ret != CUDA_SUCCESS )
          std::cout << "Could not create context on device 0" << std::endl;
        // Create module from binary file
        CUmodule cuModule;
        std::cout << "sh_addr = " <<  header.sh_addr << std::endl;
        unsigned long long offset = header.sh_addr;
        
        unsigned long long cuOffset = _find_cubin_offset( header, start_ptr, offset, "_Z11hello_worldv");

        const void * fatbin = amp;((unsigned char *) start_ptr)[cuOffset];
        
        std:: cout << "fat bin = " << fatbin << std::endl;

        ret = cuModuleLoadFatBinary(amp;cuModule, fatbin );

        if( ret != CUDA_SUCCESS )
        {
          std::cout << "Failed to load self fatbin : " << argv[0] << " : " << ret<< std::endl;
        }

        CUfunction khw;
        //ret = cuModuleGetFunction(amp;khw, cuModule, "hello_world");
        ret = cuModuleGetFunction(amp;khw, cuModule, "_Z11hello_worldv");
        if( ret != CUDA_SUCCESS )
        {
          std::cout << "Failed to get hello_world from " << argv[0] << " : " << ret <<  std::endl;
        }
        else ret = cuLaunchKernel(khw, 1, 1, 1, 1, 1, 1, 0, 0, NULL, 0);

        if( ret != CUDA_SUCCESS )
        {
          std::cout << "Failed to launch : hello_world "  << std::endl;
        }

        ret = cuModuleUnload(cuModule);

        if( ret != CUDA_SUCCESS )
        {
          std::cout << "Failed to unload self fatbin : " << argv[0] << std::endl;
          return -1;
        }

        if (cudaDeviceSynchronize() != cudaSuccess)
        {
          printf ("Cuda call failedn");
        }

        //unmap sutff
        munmap(start_ptr, sz);
        return 0;
      }
    }

  }

  fclose(file);

  return 0;