Включение стандартных заголовков C в код CUDA NVRTC

Я пишу ядро ​​CUDA, которое компилируется во время выполнения с использованием NVRTC (CUDA версия 9.2 с NVRTC версия 7.5), для которого требуется stdint.h заголовок, чтобы иметь int32_t и т. д.

Если я пишу исходный код ядра без включения, он работает правильно. Например ядро

extern "C" __global__ void f() { ... }

Компилируется в код PTX, где f определяется как .visible .entry f,

Но если исходный код ядра

#include <stdint.h>
extern "C" __global__ void f() { ... }

это сообщает A function without execution space annotations (__host__/__device__/__global__) is considered a host function, and host functions are not allowed in JIT mode. (также без extern "C").

Переходя -default-device делает код PTX .visible .func f, поэтому функция не может быть вызвана с хоста.

Есть ли способ включить заголовки в исходный код, и при этом __global__ функция входа? Или, альтернативно, способ узнать, какое соглашение о целочисленных размерах используется компилятором NVRTC, чтобы int32_t и т.д. типы могут быть определены вручную?

Изменить: Пример программы, которая показывает проблему:

#include <cstdlib>
#include <string>
#include <vector>
#include <memory>
#include <cassert>
#include <iostream>

#include <cuda.h>
#include <cuda_runtime.h>
#include <nvrtc.h>

[[noreturn]] void fail(const std::string& msg, int code) {
    std::cerr << "error: " << msg << " (" << code << ')' << std::endl;
    std::exit(EXIT_FAILURE);
}


std::unique_ptr<char[]> compile_to_ptx(const char* program_source) {
    nvrtcResult rv;

    // create nvrtc program
    nvrtcProgram prog;
    rv = nvrtcCreateProgram(
        &prog,
        program_source,
        "program.cu",
        0,
        nullptr,
        nullptr
    );
    if(rv != NVRTC_SUCCESS) fail("nvrtcCreateProgram", rv);

    // compile nvrtc program
    std::vector<const char*> options = {
        "--gpu-architecture=compute_30"
    };
    //options.push_back("-default-device");
    rv = nvrtcCompileProgram(prog, options.size(), options.data());
    if(rv != NVRTC_SUCCESS) {
        std::size_t log_size;
        rv = nvrtcGetProgramLogSize(prog, &log_size);
        if(rv != NVRTC_SUCCESS) fail("nvrtcGetProgramLogSize", rv);

        auto log = std::make_unique<char[]>(log_size);
        rv = nvrtcGetProgramLog(prog, log.get());
        if(rv != NVRTC_SUCCESS) fail("nvrtcGetProgramLog", rv);
        assert(log[log_size - 1] == '\0');

        std::cerr << "Compile error; log:\n" << log.get() << std::endl;

        fail("nvrtcCompileProgram", rv);
    }

    // get ptx code
    std::size_t ptx_size;
    rv = nvrtcGetPTXSize(prog, &ptx_size);
    if(rv != NVRTC_SUCCESS) fail("nvrtcGetPTXSize", rv);

    auto ptx = std::make_unique<char[]>(ptx_size);
    rv = nvrtcGetPTX(prog, ptx.get());
    if(rv != NVRTC_SUCCESS) fail("nvrtcGetPTX", rv);
    assert(ptx[ptx_size - 1] == '\0');

    nvrtcDestroyProgram(&prog);

    return ptx;
}

const char program_source[] = R"%%%(
//#include <stdint.h>
extern "C" __global__ void f(int* in, int* out) {
    out[threadIdx.x] = in[threadIdx.x];
}
)%%%";

int main() {
    CUresult rv;

    // initialize CUDA
    rv = cuInit(0);
    if(rv != CUDA_SUCCESS) fail("cuInit", rv);

    // compile program to ptx
    auto ptx = compile_to_ptx(program_source);
    std::cout << "PTX code:\n" << ptx.get() << std::endl;
}

когда //#include <stdint.h> в исходном коде ядра он не комментируется и больше не компилируется. когда //options.push_back("-default-device"); не комментируется, компилируется, но не помечает функцию f как .entry,

CMakeLists.txt для его компиляции (требуется API драйвера CUDA + NVRTC)

cmake_minimum_required(VERSION 3.4)
project(cudabug CXX)

find_package(CUDA REQUIRED)

set(CMAKE_CXX_STANDARD 14)
set(CMAKE_CXX_STANDARD_REQUIRED 14)

add_executable(cudabug cudabug.cc)
include_directories(SYSTEM ${CUDA_INCLUDE_DIRS})
link_directories(${CUDA_LIBRARY_DIRS})
target_link_libraries(cudabug PUBLIC ${CUDA_LIBRARIES} nvrtc cuda)
Источник

2 ответа

Решение

[Предисловие: это очень хакерский ответ, специфичный для цепочки инструментов GNU (хотя я подозреваю, что проблема в этом вопросе также специфична для цепочки инструментов GNU)].

Казалось бы, проблема здесь со стандартным заголовком GNU features.h который втягивается в stdint.hи который в конечном итоге определяет множество функций-заглушек, которые имеют значение по умолчанию __host__ пространство компиляции и вызвать взрыв nvrtc. Также кажется, что -default-device Опция приведет к разрешенному набору функций компилятора glibC, что приведет к сбою всего компилятора nvrtc.

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

const char program_source[] = R"%%%(
#define __ASSEMBLER__
#define __extension__
#include <stdint.h>
extern "C" __global__ void f(int32_t* in, int32_t* out) {
    out[threadIdx.x] = in[threadIdx.x];
}
)%%%";

дайте мне это:

$ nvcc -std=c++14 -ccbin=g++-7 jit_header.cu -o jitheader -lnvrtc -lcuda
$ ./jitheader 
PTX code:
//
// Generated by NVIDIA NVVM Compiler
//
// Compiler Build ID: CL-24330188
// Cuda compilation tools, release 9.2, V9.2.148
// Based on LLVM 3.4svn
//

.version 6.2
.target sm_30
.address_size 64

    // .globl   f

.visible .entry f(
    .param .u64 f_param_0,
    .param .u64 f_param_1
)
{
    .reg .b32   %r<3>;
    .reg .b64   %rd<8>;


    ld.param.u64    %rd1, [f_param_0];
    ld.param.u64    %rd2, [f_param_1];
    cvta.to.global.u64  %rd3, %rd2;
    cvta.to.global.u64  %rd4, %rd1;
    mov.u32     %r1, %tid.x;
    mul.wide.u32    %rd5, %r1, 4;
    add.s64     %rd6, %rd4, %rd5;
    ld.global.u32   %r2, [%rd6];
    add.s64     %rd7, %rd3, %rd5;
    st.global.u32   [%rd7], %r2;
    ret;
}

Большое предостережение: это сработало в системе glibC, на которой я ее пробовал. Вероятно, он не будет работать с другими наборами инструментов или реализациями libC (если, действительно, у них есть эта проблема).

Источник

Другой альтернативой является создание замен для некоторых заголовков стандартной библиотеки. API NVRTC поддерживает указание вами содержимого файла заголовка в виде строк, связанных с именами заголовков, прежде чем он будет просматривать файловую систему для вас. Этот подход принят в NVIDIA JITify, и я сам применил его, работая над чем-то еще, что может быть выпущено или не выпущено.

Простой способ сделать это. Вы можете просто взять заглушки заголовков JITify для `climits отсюда, которые я также прикрепляю, так как они вполне нужны. Кроме того, вы можете сгенерировать эту заглушку самостоятельно, чтобы убедиться, что вы не упускаете ничего важного из стандарта. Вот как это работает:

  1. Начните с вашего файла (или cstdintфайл в зависимости от обстоятельств);

  2. Для каждой директивы include в файле (и рекурсивно для каждого включения в include и т. д.):

    2.1 Выясните, можете ли вы вообще пропустить включение файла (возможно, сделав несколько определений, которые, как известно, хранятся на графическом процессоре).2.2 Если вы не уверены, что можете пропустить файл, включите его полностью и перейдите к (2.) или оставьте его как отдельный заголовок (и примените к нему весь процесс в (1.)).

  3. Теперь у вас есть заголовочный файл, который включает только заголовочные файлы, безопасные для устройства (или вообще не содержит).

  4. Частично предварительно обработайте файл, удалив все, что не будет использоваться на графическом процессоре. Удалите строки, которые могут вызывать проблемы на графическом процессоре (например, #pragmas) и добавить __device__ __host__или просто __host__ в соответствии с каждым объявлением функции.

Важное примечание . Для этого необходимо обратить внимание на лицензии и авторские права. Вы будете создавать «производную работу» из вкладов glibc и/или JITify и/или StackOverflow и т. д.

Теперь обещанное от NVIDIA JITify. Я адаптировал их, чтобы не было пространств имен:

stdint.h:

      #pragma once
#include <limits.h>
typedef signed char      int8_t;
typedef signed short     int16_t;
typedef signed int       int32_t;
typedef signed long long int64_t;
typedef signed char      int_fast8_t;
typedef signed short     int_fast16_t;
typedef signed int       int_fast32_t;
typedef signed long long int_fast64_t;
typedef signed char      int_least8_t;
typedef signed short     int_least16_t;
typedef signed int       int_least32_t;
typedef signed long long int_least64_t;
typedef signed long long intmax_t;
typedef signed long      intptr_t; //optional
typedef unsigned char      uint8_t;
typedef unsigned short     uint16_t;
typedef unsigned int       uint32_t;
typedef unsigned long long uint64_t;
typedef unsigned char      uint_fast8_t;
typedef unsigned short     uint_fast16_t;
typedef unsigned int       uint_fast32_t;
typedef unsigned long long uint_fast64_t;
typedef unsigned char      uint_least8_t;
typedef unsigned short     uint_least16_t;
typedef unsigned int       uint_least32_t;
typedef unsigned long long uint_least64_t;
typedef unsigned long long uintmax_t;
#define INT8_MIN    SCHAR_MIN
#define INT16_MIN   SHRT_MIN
#if defined _WIN32 || defined _WIN64
#define WCHAR_MIN   SHRT_MIN
#define WCHAR_MAX   SHRT_MAX
typedef unsigned long long uintptr_t; //optional
#else
#define WCHAR_MIN   INT_MIN
#define WCHAR_MAX   INT_MAX
typedef unsigned long      uintptr_t; //optional
#endif
#define INT32_MIN   INT_MIN
#define INT64_MIN   LLONG_MIN
#define INT8_MAX    SCHAR_MAX
#define INT16_MAX   SHRT_MAX
#define INT32_MAX   INT_MAX
#define INT64_MAX   LLONG_MAX
#define UINT8_MAX   UCHAR_MAX
#define UINT16_MAX  USHRT_MAX
#define UINT32_MAX  UINT_MAX
#define UINT64_MAX  ULLONG_MAX
#define INTPTR_MIN  LONG_MIN
#define INTMAX_MIN  LLONG_MIN
#define INTPTR_MAX  LONG_MAX
#define INTMAX_MAX  LLONG_MAX
#define UINTPTR_MAX ULONG_MAX
#define UINTMAX_MAX ULLONG_MAX
#define PTRDIFF_MIN INTPTR_MIN
#define PTRDIFF_MAX INTPTR_MAX
#define SIZE_MAX    UINT64_MAX

limits.h:

      #pragma once
#if defined _WIN32 || defined _WIN64
 #define __WORDSIZE 32
#else
 #if defined __x86_64__ && !defined __ILP32__
  #define __WORDSIZE 64
 #else
  #define __WORDSIZE 32
 #endif
#endif
#define MB_LEN_MAX  16
#define CHAR_BIT    8
#define SCHAR_MIN   (-128)
#define SCHAR_MAX   127
#define UCHAR_MAX   255
enum {
  _JITIFY_CHAR_IS_UNSIGNED = (char)-1 >= 0,
  CHAR_MIN = _JITIFY_CHAR_IS_UNSIGNED ? 0 : SCHAR_MIN,
  CHAR_MAX = _JITIFY_CHAR_IS_UNSIGNED ? UCHAR_MAX : SCHAR_MAX,
};
#define SHRT_MIN    (-32768)
#define SHRT_MAX    32767
#define USHRT_MAX   65535
#define INT_MIN     (-INT_MAX - 1)
#define INT_MAX     2147483647
#define UINT_MAX    4294967295U
#if __WORDSIZE == 64
 # define LONG_MAX  9223372036854775807L
#else
 # define LONG_MAX  2147483647L
#endif
#define LONG_MIN    (-LONG_MAX - 1L)
#if __WORDSIZE == 64
 #define ULONG_MAX  18446744073709551615UL
#else
 #define ULONG_MAX  4294967295UL
#endif
#define LLONG_MAX  9223372036854775807LL
#define LLONG_MIN  (-LLONG_MAX - 1LL)
#define ULLONG_MAX 18446744073709551615ULL
Источник
Другие вопросы по тегам