博客来源：https://leimao.github.io/blog/CUDA-Driver-Runtime-Load-Run-Kernel/ ，来自Lei Mao，已获得作者转载授权。后续会转载一些Lei Mao的CUDA相关Blog，也是一个完整的专栏，Blog会从稍早一些的CUDA架构到当前最新的CUDA架构，也会包含实用工程技巧，底层指令分析，Cutlass分析等等多个课题，是一个时间线十分明确的专栏。
使用CUDA driver API在运行时加载CUDA kernel
简介
在某些情况下，我们希望动态打开一些库并运行库中的某些函数。如果某些库并不总是被使用，这可以为程序节省一些内存。要加载C和C++库和函数，我们可以使用dlopen和dlsym函数来打开共享库并获取用extern "C"指定的函数地址。
在CUDA中，有时CUDA kernel 没有被编译到库中，而是被编译成PTX、CUBIN或FATBIN文件。在这种情况下，我们可以使用CUDA driver API来加载PTX、CUBIN或FATBIN文件并运行CUDA kernel 。
构建PTX/CUBIN/FATBIN文件
CUDA kernel 没有头文件。CUDA kernel 必须用extern “C”指定。否则，CUDA driver API无法使用函数名找到CUDA kernel 。CUDA kernel 文件将被编译成PTX、CUBIN或FATBIN文件。
// Build commands
// ptx
// nvcc --ptx --gpu-architecture=compute_80 vector_add.cu -o vector_add.ptx
// cubin
// nvcc --cubin --gpu-architecture=compute_80 --gpu-code=sm_86 vector_add.cu -o
// vector_add.cubin
// fatbin
// nvcc --fatbin --gpu-architecture=compute_80 --gpu-code=sm_86 vector_add.cu -o
// vector_add.fatbin
#include <cuda_runtime.h>
// The extern "C" is necessary.
extern"C"__global__ void vector_add(int const* a, int const* b, int* c, unsigned int n)
{
unsignedintconst stride{blockDim.x * gridDim.x};
unsignedintconst start_idx{blockDim.x * blockIdx.x + threadIdx.x};
for (unsignedint i{start_idx}; i < n; i += stride)
{
c[i] = a[i] + b[i];
}
}
要将CUDA kernel 编译成PTX、CUBIN或FATBIN文件，请运行以下命令。请根据您使用的GPU架构和GPU SM调整--gpu-architecture和--gpu-code标志。
$ ls -l
-rw-rw-r-- 1 1000 1000 699 Feb 24 19:30 vector_add.cu
-rw-r--r-- 1 root root 2984 Feb 24 19:29 vector_add.cubin
-rw-r--r-- 1 root root 3072 Feb 24 19:29 vector_add.fatbin
-rw-r--r-- 1 root root 1279 Feb 24 19:29 vector_add.ptx
也可以使用CMake来构建PTX/CUBIN/FATBIN文件。在我们的情况下，我们可以使用以下CMakeLists.txt文件。
cmake_minimum_required(VERSION 3.28)
project(Build-CUBIN VERSION 0.0.1 LANGUAGES CXX CUDA)
set(CMAKE_CXX_STANDARD 14)
set(CMAKE_CXX_STANDARD_REQUIRED ON)
# Find CUDA Toolkit
find_package(CUDAToolkit REQUIRED)
# PTX
add_library(vector_add_ptx_sm86 OBJECT vector_add.cu)
set_target_properties(vector_add_ptx_sm86 PROPERTIES CUDA_PTX_COMPILATION ON)
set_target_properties(vector_add_ptx_sm86 PROPERTIES CUDA_ARCHITECTURES "86-virtual")
add_library(vector_add_ptx_sm80 OBJECT vector_add.cu)
set_target_properties(vector_add_ptx_sm80 PROPERTIES CUDA_PTX_COMPILATION ON)
set_target_properties(vector_add_ptx_sm80 PROPERTIES CUDA_ARCHITECTURES "80-virtual")
add_library(vector_add_ptx_sm70 OBJECT vector_add.cu)
set_target_properties(vector_add_ptx_sm70 PROPERTIES CUDA_PTX_COMPILATION ON)
set_target_properties(vector_add_ptx_sm70 PROPERTIES CUDA_ARCHITECTURES "70-virtual")
# CUBIN
add_library(vector_add_cubin_sm86 OBJECT vector_add.cu)
set_target_properties(vector_add_cubin_sm86 PROPERTIES CUDA_CUBIN_COMPILATION ON)
set_target_properties(vector_add_cubin_sm86 PROPERTIES CUDA_ARCHITECTURES "86-real")
add_library(vector_add_cubin_sm80 OBJECT vector_add.cu)
set_target_properties(vector_add_cubin_sm80 PROPERTIES CUDA_CUBIN_COMPILATION ON)
set_target_properties(vector_add_cubin_sm80 PROPERTIES CUDA_ARCHITECTURES "80-real")
add_library(vector_add_cubin_sm70 OBJECT vector_add.cu)
set_target_properties(vector_add_cubin_sm70 PROPERTIES CUDA_CUBIN_COMPILATION ON)
set_target_properties(vector_add_cubin_sm70 PROPERTIES CUDA_ARCHITECTURES "70-real")
# FATBIN
add_library(vector_add_fatbin OBJECT vector_add.cu)
set_target_properties(vector_add_fatbin PROPERTIES CUDA_FATBIN_COMPILATION ON)
set_target_properties(vector_add_fatbin PROPERTIES CUDA_ARCHITECTURES "80-virtual;80-real;80-real")
使用CUDA driver API加载和运行PTX/CUBIN/FATBIN文件
要在运行时从PTX、CUBIN或FATBIN文件中加载和运行CUDA kernel ，需要使用CUDA driver API。
// Build command
// g++ run_vector_add.cpp -I/usr/local/cuda/include -L/usr/local/cuda/lib64
// -lcuda -o run_vector_add
// 使用 CUDA 驱动 API 就足够了。
#include <cuda.h>
#include <fstream>
#include <iostream>
#include <string>
#include <vector>
#define CHECK_CUDA_DRIVER_ERROR(val) \
check_driver((val), #val, __FILE__, __LINE__)
void check_driver(CUresult err, char const* func, char const* file, int line)
{
if (err != CUDA_SUCCESS)
{
std::cerr << "CUDA Driver Error at: " << file << ":" << line
<< std::endl;
charconst* err_string{nullptr};
std::cerr << cuGetErrorString(err, &err_string) << std::endl;
std::cerr << err_string << " " << func << std::endl;
std::exit(EXIT_FAILURE);
}
}
int main(int argc, char** argv)
{
// 检查命令行参数的数量。
if (argc != 2)
{
std::cerr << "Usage: " << argv[0] << " <ptx_cubin_fatbin_file_path>"
<< std::endl;
return1;
}
// 从命令行读取 CUDA kernel 库文件路径。
std::stringconst ptx_cubin_fatbin_file_path{argv[1]};
// 使用 CUDA 设备 0。
CUdevice cuda_device{0};
CUfunction vector_add_kernel{};
CUmodule cuda_module{};
CUcontext cuda_context{};
CUstream cuda_stream{};
// 目前，此标志必须为 0。
unsignedintconst cuda_driver_init_flags{0};
// 初始化 CUDA 驱动 API。
CHECK_CUDA_DRIVER_ERROR(cuInit(cuda_driver_init_flags));
// 目前，此标志必须为 0。
unsignedintconst cuda_context_init_flags{0};
// 创建 CUDA 上下文。
CHECK_CUDA_DRIVER_ERROR(
cuCtxCreate(&cuda_context, cuda_context_init_flags, cuda_device));
// 创建 CUDA 流。
CHECK_CUDA_DRIVER_ERROR(
cuStreamCreate(&cuda_stream, CU_STREAM_NON_BLOCKING));
// 确保 cubin/ptx/fatbin 文件存在。
std::ifstream cuda_kernel_library_file{ptx_cubin_fatbin_file_path};
if (!cuda_kernel_library_file)
{
std::cerr << "Error: The cubin/ptx/fatbin file does not exist."
<< std::endl;
return1;
}
// 从 cubin/ptx/fatbin 文件创建 CUDA 模块。
CHECK_CUDA_DRIVER_ERROR(
cuModuleLoad(&cuda_module, ptx_cubin_fatbin_file_path.c_str()));
// 从 CUDA 模块获取 CUDA kernel 。
CHECK_CUDA_DRIVER_ERROR(
cuModuleGetFunction(&vector_add_kernel, cuda_module, "vector_add"));
// 向量中的元素数量。
unsignedint num_elements{8192};
std::vector<int> host_vector_a(num_elements, 1);
std::vector<int> host_vector_b(num_elements, 2);
std::vector<int> host_vector_c(num_elements, -1);
std::vector<int> host_vector_c_reference(num_elements, -2);
// 初始化输入向量。
for (size_t i{0}; i < num_elements; ++i)
{
host_vector_a.at(i) = i;
host_vector_b.at(i) = i;
}
// 计算参考结果。
for (size_t i{0}; i < num_elements; ++i)
{
host_vector_c_reference.at(i) =
host_vector_a.at(i) + host_vector_b.at(i);
}
// 为输入向量分配设备内存。
CUdeviceptr device_vector_a{};
CUdeviceptr device_vector_b{};
CUdeviceptr device_vector_c{};
CHECK_CUDA_DRIVER_ERROR(
cuMemAlloc(&device_vector_a, num_elements * sizeof(int)));
CHECK_CUDA_DRIVER_ERROR(
cuMemAlloc(&device_vector_b, num_elements * sizeof(int)));
CHECK_CUDA_DRIVER_ERROR(
cuMemAlloc(&device_vector_c, num_elements * sizeof(int)));
// 将输入向量从主机复制到设备。
CHECK_CUDA_DRIVER_ERROR(cuMemcpyHtoD(device_vector_a, host_vector_a.data(),
num_elements * sizeof(int)));
CHECK_CUDA_DRIVER_ERROR(cuMemcpyHtoD(device_vector_b, host_vector_b.data(),
num_elements * sizeof(int)));
// 设置 kernel 参数。
void* kernel_params[]{&device_vector_a, &device_vector_b, &device_vector_c,
&num_elements};
// 启动 CUDA kernel 。
unsignedintconst block_size_x{256};
unsignedintconst block_size_y{1};
unsignedintconst block_size_z{1};
unsignedintconst grid_size_x{(num_elements + block_size_x - 1) /
block_size_x};
unsignedintconst grid_size_y{1};
unsignedintconst grid_size_z{1};
unsignedintconst shared_memory_size{0};
CHECK_CUDA_DRIVER_ERROR(cuLaunchKernel(
vector_add_kernel, grid_size_x, grid_size_y, grid_size_z, block_size_x,
block_size_y, block_size_z, shared_memory_size, cuda_stream,
kernel_params, nullptr));
CHECK_CUDA_DRIVER_ERROR(cuStreamSynchronize(cuda_stream));
// 将结果向量从设备复制到主机。
CHECK_CUDA_DRIVER_ERROR(cuMemcpyDtoH(host_vector_c.data(), device_vector_c,
num_elements * sizeof(int)));
// 验证结果。
for (size_t i{0}; i < num_elements; ++i)
{
if (host_vector_c.at(i) != host_vector_c_reference.at(i))
{
std::cerr << "Error: The result is incorrect." << std::endl;
return1;
}
}
// 释放设备内存。
CHECK_CUDA_DRIVER_ERROR(cuMemFree(device_vector_a));
CHECK_CUDA_DRIVER_ERROR(cuMemFree(device_vector_b));
CHECK_CUDA_DRIVER_ERROR(cuMemFree(device_vector_c));
// 销毁 CUDA 流。
CHECK_CUDA_DRIVER_ERROR(cuStreamDestroy(cuda_stream));
// 销毁 CUDA 模块。
CHECK_CUDA_DRIVER_ERROR(cuModuleUnload(cuda_module));
// 销毁 CUDA 上下文。
CHECK_CUDA_DRIVER_ERROR(cuCtxDestroy(cuda_context));
return0;
}
要构建程序，请运行以下命令。
$ g++ -o run_vector_add run_vector_add.cpp -I/usr/local/cuda/include -L/usr/local/cuda/lib64 -lcuda
要运行程序，请运行以下命令。应该不会遇到错误。
$ ./run_vector_add vector_add.ptx
$ ./run_vector_add vector_add.cubin
$ ./run_vector_add vector_add.fatbin
使用CUDA driver API加载和运行PTX/CUBIN/FATBIN字符串
通常也会看到PTX、CUBIN或FATBIN字符串嵌入在C++源代码中。例如，TensorRT多头注意力 kernel cubins(https://github.com/NVIDIA/TensorRT/blob/97ff24489d0ea979c418c7a0847dfc14c8483846/plugin/bertQKVToContextPlugin/fused_multihead_attention/src/fused_multihead_attention_fp16_128_64_kernel.sm80.cpp)在C++源代码中被表示为字符串。在这种情况下，PTX、CUBIN或FATBIN字符串使用一些脚本从PTX、CUBIN或FATBIN文件转换而来。
例如，我们可以使用以下Python脚本将我们刚才构建的PTX、CUBIN或FATBIN文件转换为PTX、CUBIN或FATBIN字符串。
# Open a specified file and convert every byte to a list of unsigned chars in hexadecimal format,
# then write the list to a new text file.
# Usage: python convert_file_to_chars.py --input_bin_file <input_file> --output_char_file <output_file>
import argparse
import os
def main():
parser = argparse.ArgumentParser(
description='Convert a binary file to a list of unsigned chars')
parser.add_argument('--input_bin_file',
type=str,
help='The input binary file')
parser.add_argument('--output_char_file',
type=str,
help='The output char file')
args = parser.parse_args()
ifnot os.path.isfile(args.input_bin_file):
print('The input binary file does not exist')
return
with open(args.input_bin_file, 'rb') as f:
content = f.read()
chars = list(content)
# Convert each char to a string in hexadecimal format "0x??"
chars = ['0x{:02x}'.format(c) for c in chars]
with open(args.output_char_file, 'w') as f:
f.write(', '.join(chars))
if __name__ == "__main__":
main()
$ python3 convert_file_to_chars.py --help
usage: convert_file_to_chars.py [-h] [--input_bin_file INPUT_BIN_FILE]
[--output_char_file OUTPUT_CHAR_FILE]
Convert a binary file to a list of unsigned chars
options:
-h, --help show this help message and exit
--input_bin_file INPUT_BIN_FILE
The input binary file
--output_char_file OUTPUT_CHAR_FILE
The output char file
PTX、CUBIN或FATBIN字符串应该然后被复制到C++源代码中。
// vector_add_ptx_cubin_fatbin_sm86.hpp
extern unsigned char vector_add_int_ptx[];
extern unsigned char const vector_add_int_cubin[];
extern unsigned char vector_add_int_fatbin[];
// vector_add_ptx_cubin_fatbin_sm86.cpp
#include "vector_add_ptx_cubin_fatbin_sm86.hpp"
unsignedchar vector_add_int_ptx[] = {
0x2f, 0x2f, 0x0a, 0x2f, 0x2f, 0x20, 0x47, 0x65, 0x6e, 0x65, 0x72, 0x61,
0x74, 0x65, 0x64, 0x20, 0x62, 0x79, 0x20, 0x4e, 0x56, 0x49, 0x44, 0x49,
0x41, 0x20, 0x4e, 0x56, 0x56, 0x4d, 0x20, 0x43, 0x6f, 0x6d, 0x70, 0x69,
0x6c, 0x65, 0x72, 0x0a, 0x2f, 0x2f, 0x0a, 0x2f, 0x2f, 0x20, 0x43, 0x6f,
0x6d, 0x70, 0x69, 0x6c, 0x65, 0x72, 0x20, 0x42, 0x75, 0x69, 0x6c, 0x64,
0x20, 0x49, 0x44, 0x3a, 0x20, 0x43, 0x4c, 0x2d, 0x33, 0x34, 0x30, 0x39,
0x37, 0x39, 0x36, 0x37, 0x0a, 0x2f, 0x2f, 0x20, 0x43, 0x75, 0x64, 0x61,
0x20, 0x63, 0x6f, 0x6d, 0x70, 0x69, 0x6c, 0x61, 0x74, 0x69, 0x6f, 0x6e,
0x20, 0x74, 0x6f, 0x6f, 0x6c, 0x73, 0x2c, 0x20, 0x72, 0x65, 0x6c, 0x65,
0x61, 0x73, 0x65, 0x20, 0x31, 0x32, 0x2e, 0x34, 0x2c, 0x20, 0x56, 0x31,
0x32, 0x2e, 0x34, 0x2e, 0x31, 0x33, 0x31, 0x0a, 0x2f, 0x2f, 0x20, 0x42,
0x61, 0x73, 0x65, 0x64, 0x20, 0x6f, 0x6e, 0x20, 0x4e, 0x56, 0x56, 0x4d,
0x20, 0x37, 0x2e, 0x30, 0x2e, 0x31, 0x0a, 0x2f, 0x2f, 0x0a, 0x0a, 0x2e,
0x76, 0x65, 0x72, 0x73, 0x69, 0x6f, 0x6e, 0x20, 0x38, 0x2e, 0x34, 0x0a,
0x2e, 0x74, 0x61, 0x72, 0x67, 0x65, 0x74, 0x20, 0x73, 0x6d, 0x5f, 0x38,
0x30, 0x0a, 0x2e, 0x61, 0x64, 0x64, 0x72, 0x65, 0x73, 0x73, 0x5f, 0x73,
0x69, 0x7a, 0x65, 0x20, 0x36, 0x34, 0x0a, 0x0a, 0x09, 0x2f, 0x2f, 0x20,
0x2e, 0x67, 0x6c, 0x6f, 0x62, 0x6c, 0x09, 0x76, 0x65, 0x63, 0x74, 0x6f,
0x72, 0x5f, 0x61, 0x64, 0x64, 0x0a, 0x0a, 0x2e, 0x76, 0x69, 0x73, 0x69,
0x62, 0x6c, 0x65, 0x20, 0x2e, 0x65, 0x6e, 0x74, 0x72, 0x79, 0x20, 0x76,
0x65, 0x63, 0x74, 0x6f, 0x72, 0x5f, 0x61, 0x64, 0x64, 0x28, 0x0a, 0x09,
0x2e, 0x70, 0x61, 0x72, 0x61, 0x6d, 0x20, 0x2e, 0x75, 0x36, 0x34, 0x20,
0x76, 0x65, 0x63, 0x74, 0x6f, 0x72, 0x5f, 0x61, 0x64, 0x64, 0x5f, 0x70,
0x61, 0x72, 0x61, 0x6d, 0x5f, 0x30, 0x2c, 0x0a, 0x09, 0x2e, 0x70, 0x61,
0x72, 0x61, 0x6d, 0x20, 0x2e, 0x75, 0x36, 0x34, 0x20, 0x76, 0x65, 0x63,
0x74, 0x6f, 0x72, 0x5f, 0x61, 0x64, 0x64, 0x5f, 0x70, 0x61, 0x72, 0x61,
0x6d, 0x5f, 0x31, 0x2c, 0x0a, 0x09, 0x2e, 0x70, 0x61, 0x72, 0x61, 0x6d,
0x20, 0x2e, 0x75, 0x36, 0x34, 0x20, 0x76, 0x65, 0x63, 0x74, 0x6f, 0x72,
0x5f, 0x61, 0x64, 0x64, 0x5f, 0x70, 0x61, 0x72, 0x61, 0x6d, 0x5f, 0x32,
0x2c, 0x0a, 0x09, 0x2e, 0x70, 0x61, 0x72, 0x61, 0x6d, 0x20, 0x2e, 0x75,
0x33, 0x32, 0x20, 0x76, 0x65, 0x63, 0x74, 0x6f, 0x72, 0x5f, 0x61, 0x64,
0x64, 0x5f, 0x70, 0x61, 0x72, 0x61, 0x6d, 0x5f, 0x33, 0x0a, 0x29, 0x0a,
0x7b, 0x0a, 0x09, 0x2e, 0x72, 0x65, 0x67, 0x20, 0x2e, 0x70, 0x72, 0x65,
0x64, 0x20, 0x09, 0x25, 0x70, 0x3c, 0x33, 0x3e, 0x3b, 0x0a, 0x09, 0x2e,
0x72, 0x65, 0x67, 0x20, 0x2e, 0x62, 0x33, 0x32, 0x20, 0x09, 0x25, 0x72,
0x3c, 0x31, 0x34, 0x3e, 0x3b, 0x0a, 0x09, 0x2e, 0x72, 0x65, 0x67, 0x20,
0x2e, 0x62, 0x36, 0x34, 0x20, 0x09, 0x25, 0x72, 0x64, 0x3c, 0x31, 0x31,
0x3e, 0x3b, 0x0a, 0x0a, 0x0a, 0x09, 0x6c, 0x64, 0x2e, 0x70, 0x61, 0x72,
0x61, 0x6d, 0x2e, 0x75, 0x36, 0x34, 0x20, 0x09, 0x25, 0x72, 0x64, 0x34,
0x2c, 0x20, 0x5b, 0x76, 0x65, 0x63, 0x74, 0x6f, 0x72, 0x5f, 0x61, 0x64,
0x64, 0x5f, 0x70, 0x61, 0x72, 0x61, 0x6d, 0x5f, 0x30, 0x5d, 0x3b, 0x0a,
0x09, 0x6c, 0x64, 0x2e, 0x70, 0x61, 0x72, 0x61, 0x6d, 0x2e, 0x75, 0x36,
0x34, 0x20, 0x09, 0x25, 0x72, 0x64, 0x35, 0x2c, 0x20, 0x5b, 0x76, 0x65,
0x63, 0x74, 0x6f, 0x72, 0x5f, 0x61, 0x64, 0x64, 0x5f, 0x70, 0x61, 0x72,
0x61, 0x6d, 0x5f, 0x31, 0x5d, 0x3b, 0x0a, 0x09, 0x6c, 0x64, 0x2e, 0x70,
0x61, 0x72, 0x61, 0x6d, 0x2e, 0x75, 0x36, 0x34, 0x20, 0x09, 0x25, 0x72,
0x64, 0x36, 0x2c, 0x20, 0x5b, 0x76, 0x65, 0x63, 0x74, 0x6f, 0x72, 0x5f,
0x61, 0x64, 0x64, 0x5f, 0x70, 0x61, 0x72, 0x61, 0x6d, 0x5f, 0x32, 0x5d,
0x3b, 0x0a, 0x09, 0x6c, 0x64, 0x2e, 0x70, 0x61, 0x72, 0x61, 0x6d, 0x2e,
0x75, 0x33, 0x32, 0x20, 0x09, 0x25, 0x72, 0x36, 0x2c, 0x20, 0x5b, 0x76,
0x65, 0x63, 0x74, 0x6f, 0x72, 0x5f, 0x61, 0x64, 0x64, 0x5f, 0x70, 0x61,
0x72, 0x61, 0x6d, 0x5f, 0x33, 0x5d, 0x3b, 0x0a, 0x09, 0x6d, 0x6f, 0x76,
0x2e, 0x75, 0x33, 0x32, 0x20, 0x09, 0x25, 0x72, 0x31, 0x2c, 0x20, 0x25,
0x6e, 0x74, 0x69, 0x64, 0x2e, 0x78, 0x3b, 0x0a, 0x09, 0x6d, 0x6f, 0x76,
0x2e, 0x75, 0x33, 0x32, 0x20, 0x09, 0x25, 0x72, 0x37, 0x2c, 0x20, 0x25,
0x63, 0x74, 0x61, 0x69, 0x64, 0x2e, 0x78, 0x3b, 0x0a, 0x09, 0x6d, 0x6f,
0x76, 0x2e, 0x75, 0x33, 0x32, 0x20, 0x09, 0x25, 0x72, 0x38, 0x2c, 0x20,
0x25, 0x74, 0x69, 0x64, 0x2e, 0x78, 0x3b, 0x0a, 0x09, 0x6d, 0x61, 0x64,
0x2e, 0x6c, 0x6f, 0x2e, 0x73, 0x33, 0x32, 0x20, 0x09, 0x25, 0x72, 0x31,
0x33, 0x2c, 0x20, 0x25, 0x72, 0x31, 0x2c, 0x20, 0x25, 0x72, 0x37, 0x2c,
0x20, 0x25, 0x72, 0x38, 0x3b, 0x0a, 0x09, 0x73, 0x65, 0x74, 0x70, 0x2e,
0x67, 0x65, 0x2e, 0x75, 0x33, 0x32, 0x20, 0x09, 0x25, 0x70, 0x31, 0x2c,
0x20, 0x25, 0x72, 0x31, 0x33, 0x2c, 0x20, 0x25, 0x72, 0x36, 0x3b, 0x0a,
0x09, 0x40, 0x25, 0x70, 0x31, 0x20, 0x62, 0x72, 0x61, 0x20, 0x09, 0x24,
0x4c, 0x5f, 0x5f, 0x42, 0x42, 0x30, 0x5f, 0x33, 0x3b, 0x0a, 0x0a, 0x09,
0x6d, 0x6f, 0x76, 0x2e, 0x75, 0x33, 0x32, 0x20, 0x09, 0x25, 0x72, 0x39,
0x2c, 0x20, 0x25, 0x6e, 0x63, 0x74, 0x61, 0x69, 0x64, 0x2e, 0x78, 0x3b,
0x0a, 0x09, 0x6d, 0x75, 0x6c, 0x2e, 0x6c, 0x6f, 0x2e, 0x73, 0x33, 0x32,
0x20, 0x09, 0x25, 0x72, 0x33, 0x2c, 0x20, 0x25, 0x72, 0x31, 0x2c, 0x20,
0x25, 0x72, 0x39, 0x3b, 0x0a, 0x09, 0x63, 0x76, 0x74, 0x61, 0x2e, 0x74,
0x6f, 0x2e, 0x67, 0x6c, 0x6f, 0x62, 0x61, 0x6c, 0x2e, 0x75, 0x36, 0x34,
0x20, 0x09, 0x25, 0x72, 0x64, 0x31, 0x2c, 0x20, 0x25, 0x72, 0x64, 0x34,
0x3b, 0x0a, 0x09, 0x63, 0x76, 0x74, 0x61, 0x2e, 0x74, 0x6f, 0x2e, 0x67,
0x6c, 0x6f, 0x62, 0x61, 0x6c, 0x2e, 0x75, 0x36, 0x34, 0x20, 0x09, 0x25,
0x72, 0x64, 0x32, 0x2c, 0x20, 0x25, 0x72, 0x64, 0x35, 0x3b, 0x0a, 0x09,
0x63, 0x76, 0x74, 0x61, 0x2e, 0x74, 0x6f, 0x2e, 0x67, 0x6c, 0x6f, 0x62,
0x61, 0x6c, 0x2e, 0x75, 0x36, 0x34, 0x20, 0x09, 0x25, 0x72, 0x64, 0x33,
0x2c, 0x20, 0x25, 0x72, 0x64, 0x36, 0x3b, 0x0a, 0x0a, 0x24, 0x4c, 0x5f,
0x5f, 0x42, 0x42, 0x30, 0x5f, 0x32, 0x3a, 0x0a, 0x09, 0x6d, 0x75, 0x6c,
0x2e, 0x77, 0x69, 0x64, 0x65, 0x2e, 0x75, 0x33, 0x32, 0x20, 0x09, 0x25,
0x72, 0x64, 0x37, 0x2c, 0x20, 0x25, 0x72, 0x31, 0x33, 0x2c, 0x20, 0x34,
0x3b, 0x0a, 0x09, 0x61, 0x64, 0x64, 0x2e, 0x73, 0x36, 0x34, 0x20, 0x09,
0x25, 0x72, 0x64, 0x38, 0x2c, 0x20, 0x25, 0x72, 0x64, 0x31, 0x2c, 0x20,
0x25, 0x72, 0x64, 0x37, 0x3b, 0x0a, 0x09, 0x61, 0x64, 0x64, 0x2e, 0x73,
0x36, 0x34, 0x20, 0x09, 0x25, 0x72, 0x64, 0x39, 0x2c, 0x20, 0x25, 0x72,
0x64, 0x32, 0x2c, 0x20, 0x25, 0x72, 0x64, 0x37, 0x3b, 0x0a, 0x09, 0x6c,
0x64, 0x2e, 0x67, 0x6c, 0x6f, 0x62, 0x61, 0x6c, 0x2e, 0x75, 0x33, 0x32,
0x20, 0x09, 0x25, 0x72, 0x31, 0x30, 0x2c, 0x20, 0x5b, 0x25, 0x72, 0x64,
0x39, 0x5d, 0x3b, 0x0a, 0x09, 0x6c, 0x64, 0x2e, 0x67, 0x6c, 0x6f, 0x62,
0x61, 0x6c, 0x2e, 0x75, 0x33, 0x32, 0x20, 0x09, 0x25, 0x72, 0x31, 0x31,
0x2c, 0x20, 0x5b, 0x25, 0x72, 0x64, 0x38, 0x5d, 0x3b, 0x0a, 0x09, 0x61,
0x64, 0x64, 0x2e, 0x73, 0x33, 0x32, 0x20, 0x09, 0x25, 0x72, 0x31, 0x32,
0x2c, 0x20, 0x25, 0x72, 0x31, 0x30, 0x2c, 0x20, 0x25, 0x72, 0x31, 0x31,
0x3b, 0x0a, 0x09, 0x61, 0x64, 0x64, 0x2e, 0x73, 0x36, 0x34, 0x20, 0x09,
0x25, 0x72, 0x64, 0x31, 0x30, 0x2c, 0x20, 0x25, 0x72, 0x64, 0x33, 0x2c,
0x20, 0x25, 0x72, 0x64, 0x37, 0x3b, 0x0a, 0x09, 0x73, 0x74, 0x2e, 0x67,
0x6c, 0x6f, 0x62, 0x61, 0x6c, 0x2e, 0x75, 0x33, 0x32, 0x20, 0x09, 0x5b,
0x25, 0x72, 0x64, 0x31, 0x30, 0x5d, 0x2c, 0x20, 0x25, 0x72, 0x31, 0x32,
0x3b, 0x0a, 0x09, 0x61, 0x64, 0x64, 0x2e, 0x73, 0x33, 0x32, 0x20, 0x09,
0x25, 0x72, 0x31, 0x33, 0x2c, 0x20, 0x25, 0x72, 0x31, 0x33, 0x2c, 0x20,
0x25, 0x72, 0x33, 0x3b, 0x0a, 0x09, 0x73, 0x65, 0x74, 0x70, 0x2e, 0x6c,
0x74, 0x2e, 0x75, 0x33, 0x32, 0x20, 0x09, 0x25, 0x70, 0x32, 0x2c, 0x20,
0x25, 0x72, 0x31, 0x33, 0x2c, 0x20, 0x25, 0x72, 0x36, 0x3b, 0x0a, 0x09,
0x40, 0x25, 0x70, 0x32, 0x20, 0x62, 0x72, 0x61, 0x20, 0x09, 0x24, 0x4c,
0x5f, 0x5f, 0x42, 0x42, 0x30, 0x5f, 0x32, 0x3b, 0x0a, 0x0a, 0x24, 0x4c,
0x5f, 0x5f, 0x42, 0x42, 0x30, 0x5f, 0x33, 0x3a, 0x0a, 0x09, 0x72, 0x65,
0x74, 0x3b, 0x0a, 0x0a, 0x7d, 0x0a, 0x0a};
unsignedcharconst vector_add_int_cubin[] = {
0x7f, 0x45, 0x4c, 0x46, 0x02, 0x01, 0x01, 0x33, 0x07, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x02, 0x00, 0xbe, 0x00, 0x7c, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x0b, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x08, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x56, 0x05, 0x50, 0x00, 0x40, 0x00, 0x38, 0x00, 0x03, 0x00, 0x40, 0x00,
0x0c, 0x00, 0x01, 0x00, 0x00, 0x2e, 0x73, 0x68, 0x73, 0x74, 0x72, 0x74,
0x61, 0x62, 0x00, 0x2e, 0x73, 0x74, 0x72, 0x74, 0x61, 0x62, 0x00, 0x2e,
0x73, 0x79, 0x6d, 0x74, 0x61, 0x62, 0x00, 0x2e, 0x73, 0x79, 0x6d, 0x74,
0x61, 0x62, 0x5f, 0x73, 0x68, 0x6e, 0x64, 0x78, 0x00, 0x2e, 0x6e, 0x76,
0x2e, 0x69, 0x6e, 0x66, 0x6f, 0x00, 0x2e, 0x74, 0x65, 0x78, 0x74, 0x2e,
0x76, 0x65, 0x63, 0x74, 0x6f, 0x72, 0x5f, 0x61, 0x64, 0x64, 0x00, 0x2e,
0x6e, 0x76, 0x2e, 0x69, 0x6e, 0x66, 0x6f, 0x2e, 0x76, 0x65, 0x63, 0x74,
0x6f, 0x72, 0x5f, 0x61, 0x64, 0x64, 0x00, 0x2e, 0x6e, 0x76, 0x2e, 0x73,
0x68, 0x61, 0x72, 0x65, 0x64, 0x2e, 0x76, 0x65, 0x63, 0x74, 0x6f, 0x72,
0x5f, 0x61, 0x64, 0x64, 0x00, 0x2e, 0x6e, 0x76, 0x2e, 0x63, 0x6f, 0x6e,
0x73, 0x74, 0x61, 0x6e, 0x74, 0x30, 0x2e, 0x76, 0x65, 0x63, 0x74, 0x6f,
0x72, 0x5f, 0x61, 0x64, 0x64, 0x00, 0x2e, 0x72, 0x65, 0x6c, 0x2e, 0x6e,
0x76, 0x2e, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x61, 0x6e, 0x74, 0x30, 0x2e,
0x76, 0x65, 0x63, 0x74, 0x6f, 0x72, 0x5f, 0x61, 0x64, 0x64, 0x00, 0x2e,
0x64, 0x65, 0x62, 0x75, 0x67, 0x5f, 0x66, 0x72, 0x61, 0x6d, 0x65, 0x00,
0x2e, 0x72, 0x65, 0x6c, 0x2e, 0x64, 0x65, 0x62, 0x75, 0x67, 0x5f, 0x66,
0x72, 0x61, 0x6d, 0x65, 0x00, 0x2e, 0x72, 0x65, 0x6c, 0x61, 0x2e, 0x64,
0x65, 0x62, 0x75, 0x67, 0x5f, 0x66, 0x72, 0x61, 0x6d, 0x65, 0x00, 0x2e,
0x6e, 0x76, 0x2e, 0x63, 0x61, 0x6c, 0x6c, 0x67, 0x72, 0x61, 0x70, 0x68,
0x00, 0x2e, 0x6e, 0x76, 0x2e, 0x70, 0x72, 0x6f, 0x74, 0x6f, 0x74, 0x79,
0x70, 0x65, 0x00, 0x2e, 0x6e, 0x76, 0x2e, 0x72, 0x65, 0x6c, 0x2e, 0x61,
0x63, 0x74, 0x69, 0x6f, 0x6e, 0x00, 0x00, 0x2e, 0x73, 0x68, 0x73, 0x74,
0x72, 0x74, 0x61, 0x62, 0x00, 0x2e, 0x73, 0x74, 0x72, 0x74, 0x61, 0x62,
0x00, 0x2e, 0x73, 0x79, 0x6d, 0x74, 0x61, 0x62, 0x00, 0x2e, 0x73, 0x79,
0x6d, 0x74, 0x61, 0x62, 0x5f, 0x73, 0x68, 0x6e, 0x64, 0x78, 0x00, 0x2e,
0x6e, 0x76, 0x2e, 0x69, 0x6e, 0x66, 0x6f, 0x00, 0x2e, 0x74, 0x65, 0x78,
0x74, 0x2e, 0x76, 0x65, 0x63, 0x74, 0x6f, 0x72, 0x5f, 0x61, 0x64, 0x64,
0x00, 0x2e, 0x6e, 0x76, 0x2e, 0x69, 0x6e, 0x66, 0x6f, 0x2e, 0x76, 0x65,
0x63, 0x74, 0x6f, 0x72, 0x5f, 0x61, 0x64, 0x64, 0x00, 0x2e, 0x6e, 0x76,
0x2e, 0x73, 0x68, 0x61, 0x72, 0x65, 0x64, 0x2e, 0x76, 0x65, 0x63, 0x74,
0x6f, 0x72, 0x5f, 0x61, 0x64, 0x64, 0x00, 0x2e, 0x72, 0x65, 0x6c, 0x2e,
0x6e, 0x76, 0x2e, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x61, 0x6e, 0x74, 0x30,
0x2e, 0x76, 0x65, 0x63, 0x74, 0x6f, 0x72, 0x5f, 0x61, 0x64, 0x64, 0x00,
0x2e, 0x6e, 0x76, 0x2e, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x61, 0x6e, 0x74,
0x30, 0x2e, 0x76, 0x65, 0x63, 0x74, 0x6f, 0x72, 0x5f, 0x61, 0x64, 0x64,
0x00, 0x2e, 0x64, 0x65, 0x62, 0x75, 0x67, 0x5f, 0x66, 0x72, 0x61, 0x6d,
0x65, 0x00, 0x2e, 0x72, 0x65, 0x6c, 0x2e, 0x64, 0x65, 0x62, 0x75, 0x67,
0x5f, 0x66, 0x72, 0x61, 0x6d, 0x65, 0x00, 0x2e, 0x72, 0x65, 0x6c, 0x61,
0x2e, 0x64, 0x65, 0x62, 0x75, 0x67, 0x5f, 0x66, 0x72, 0x61, 0x6d, 0x65,
0x00, 0x2e, 0x6e, 0x76, 0x2e, 0x63, 0x61, 0x6c, 0x6c, 0x67, 0x72, 0x61,
0x70, 0x68, 0x00, 0x2e, 0x6e, 0x76, 0x2e, 0x70, 0x72, 0x6f, 0x74, 0x6f,
0x74, 0x79, 0x70, 0x65, 0x00, 0x2e, 0x6e, 0x76, 0x2e, 0x72, 0x65, 0x6c,
0x2e, 0x61, 0x63, 0x74, 0x69, 0x6f, 0x6e, 0x00, 0x76, 0x65, 0x63, 0x74,
0x6f, 0x72, 0x5f, 0x61, 0x64, 0x64, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x32, 0x00, 0x00, 0x00,
0x03, 0x00, 0x0b, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x8a, 0x00, 0x00, 0x00,
0x03, 0x00, 0x0a, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0xa3, 0x00, 0x00, 0x00,
0x03, 0x00, 0x04, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0xd3, 0x00, 0x00, 0x00,
0x03, 0x00, 0x07, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0xef, 0x00, 0x00, 0x00,
0x03, 0x00, 0x08, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0xfe, 0x00, 0x00, 0x00,
0x12, 0x10, 0x0b, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x02, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0xff, 0xff, 0xff, 0xff,
0x24, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0xff, 0xff, 0xff, 0xff,
0xff, 0xff, 0xff, 0xff, 0x03, 0x00, 0x04, 0x7c, 0xff, 0xff, 0xff, 0xff,
0x0f, 0x0c, 0x81, 0x80, 0x80, 0x28, 0x00, 0x08, 0xff, 0x81, 0x80, 0x28,
0x08, 0x81, 0x80, 0x80, 0x28, 0x00, 0x00, 0x00, 0xff, 0xff, 0xff, 0xff,
0x34, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x02, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x04, 0x04, 0x00, 0x00,
0x00, 0x04, 0x14, 0x00, 0x00, 0x00, 0x0c, 0x81, 0x80, 0x80, 0x28, 0x00,
0x04, 0x34, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x04, 0x2f, 0x08, 0x00, 0x06, 0x00, 0x00, 0x00, 0x0e, 0x00, 0x00, 0x00,
0x04, 0x12, 0x08, 0x00, 0x06, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x04, 0x11, 0x08, 0x00, 0x06, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x04, 0x12, 0x08, 0x00, 0x06, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x04, 0x37, 0x04, 0x00, 0x7c, 0x00, 0x00, 0x00, 0x01, 0x35, 0x00, 0x00,
0x04, 0x0a, 0x08, 0x00, 0x02, 0x00, 0x00, 0x00, 0x60, 0x01, 0x1c, 0x00,
0x03, 0x19, 0x1c, 0x00, 0x04, 0x17, 0x0c, 0x00, 0x00, 0x00, 0x00, 0x00,
0x03, 0x00, 0x18, 0x00, 0x00, 0xf0, 0x11, 0x00, 0x04, 0x17, 0x0c, 0x00,
0x00, 0x00, 0x00, 0x00, 0x02, 0x00, 0x10, 0x00, 0x00, 0xf0, 0x21, 0x00,
0x04, 0x17, 0x0c, 0x00, 0x00, 0x00, 0x00, 0x00, 0x01, 0x00, 0x08, 0x00,
0x00, 0xf0, 0x21, 0x00, 0x04, 0x17, 0x0c, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0xf0, 0x21, 0x00, 0x03, 0x1b, 0xff, 0x00,
0x04, 0x1c, 0x08, 0x00, 0x50, 0x00, 0x00, 0x00, 0x30, 0x01, 0x00, 0x00,
0x04, 0x1e, 0x04, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0xff, 0xff, 0xff, 0xff, 0x00, 0x00, 0x00, 0x00, 0xfe, 0xff, 0xff, 0xff,
0x00, 0x00, 0x00, 0x00, 0xfd, 0xff, 0xff, 0xff, 0x00, 0x00, 0x00, 0x00,
0xfc, 0xff, 0xff, 0xff, 0x00, 0x00, 0x00, 0x00, 0x73, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x11, 0x25, 0x00, 0x05, 0x36,
0x44, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x02, 0x00, 0x00, 0x00,
0x06, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x02, 0x7a, 0x01, 0x00, 0x00, 0x0a, 0x00, 0x00, 0x00, 0x0f, 0x00, 0x00,
0x00, 0xe4, 0x0f, 0x00, 0x19, 0x79, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x25, 0x00, 0x00, 0x00, 0x28, 0x0e, 0x00, 0x19, 0x79, 0x03, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x21, 0x00, 0x00, 0x00, 0x24, 0x0e, 0x00,
0x24, 0x7a, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x03, 0x02, 0x8e, 0x07,
0x00, 0xca, 0x1f, 0x00, 0x0c, 0x7a, 0x00, 0x00, 0x00, 0x5e, 0x00, 0x00,
0x70, 0x60, 0xf0, 0x03, 0x00, 0xda, 0x0f, 0x00, 0x4d, 0x09, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x80, 0x03, 0x00, 0xea, 0x0f, 0x00,
0xb9, 0x7a, 0x04, 0x00, 0x00, 0x46, 0x00, 0x00, 0x00, 0x0a, 0x00, 0x00,
0x00, 0xe4, 0x0f, 0x00, 0x02, 0x78, 0x07, 0x00, 0x04, 0x00, 0x00, 0x00,
0x00, 0x0f, 0x00, 0x00, 0x00, 0xca, 0x1f, 0x00, 0x25, 0x76, 0x02, 0x00,
0x00, 0x58, 0x00, 0x00, 0x07, 0x00, 0x8e, 0x07, 0x00, 0xc8, 0x0f, 0x00,
0x25, 0x76, 0x04, 0x00, 0x00, 0x5a, 0x00, 0x00, 0x07, 0x00, 0x8e, 0x07,
0x00, 0xe4, 0x0f, 0x0c, 0x81, 0x79, 0x03, 0x02, 0x04, 0x00, 0x00, 0x00,
0x00, 0x19, 0x1e, 0x0c, 0x00, 0xa8, 0x0e, 0x00, 0x81, 0x79, 0x04, 0x04,
0x04, 0x00, 0x00, 0x00, 0x00, 0x19, 0x1e, 0x0c, 0x00, 0xa2, 0x0e, 0x00,
0x25, 0x76, 0x06, 0x00, 0x00, 0x5c, 0x00, 0x00, 0x07, 0x00, 0x8e, 0x07,
0x00, 0xe2, 0x0f, 0x00, 0x02, 0x7a, 0x0b, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x0f, 0x00, 0x00, 0x00, 0xca, 0x0f, 0x00, 0x24, 0x7a, 0x00, 0x0b,
0x00, 0x03, 0x00, 0x00, 0x00, 0x02, 0x8e, 0x07, 0x00, 0xca, 0x0f, 0x00,
0x0c, 0x7a, 0x00, 0x00, 0x00, 0x5e, 0x00, 0x00, 0x70, 0x60, 0xf0, 0x03,
0x00, 0xe4, 0x0f, 0x00, 0x10, 0x72, 0x09, 0x04, 0x03, 0x00, 0x00, 0x00,
0xff, 0xe0, 0xff, 0x07, 0x00, 0xca, 0x4f, 0x00, 0x86, 0x79, 0x00, 0x06,
0x09, 0x00, 0x00, 0x00, 0x04, 0x19, 0x10, 0x0c, 0x00, 0xec, 0x01, 0x00,
0x47, 0x89, 0x00, 0x00, 0x40, 0xff, 0xff, 0xff, 0xff, 0xff, 0x83, 0x03,
0x00, 0xea, 0x0f, 0x00, 0x4d, 0x79, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x80, 0x03, 0x00, 0xea, 0x0f, 0x00, 0x47, 0x79, 0x00, 0x00,
0xf0, 0xff, 0xff, 0xff, 0xff, 0xff, 0x83, 0x03, 0x00, 0xc0, 0x0f, 0x00,
0x18, 0x79, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0xc0, 0x0f, 0x00, 0x18, 0x79, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0xc0, 0x0f, 0x00, 0x18, 0x79, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0xc0, 0x0f, 0x00,
0x18, 0x79, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0xc0, 0x0f, 0x00, 0x18, 0x79, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0xc0, 0x0f, 0x00, 0x18, 0x79, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0xc0, 0x0f, 0x00,
0x18, 0x79, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0xc0, 0x0f, 0x00, 0x18, 0x79, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0xc0, 0x0f, 0x00, 0x18, 0x79, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0xc0, 0x0f, 0x00,
0x18, 0x79, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0xc0, 0x0f, 0x00, 0x18, 0x79, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0xc0, 0x0f, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x01, 0x00, 0x00, 0x00, 0x03, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x40, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0xfe, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x01, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x0b, 0x00, 0x00, 0x00, 0x03, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x3e, 0x01, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x09, 0x01, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x13, 0x00, 0x00, 0x00,
0x02, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x48, 0x02, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0xa8, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x02, 0x00, 0x00, 0x00, 0x06, 0x00, 0x00, 0x00, 0x08, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x18, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0xa3, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0xf0, 0x02, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x70, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x01, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x29, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x70,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x60, 0x03, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x30, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x03, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x04, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x43, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x70, 0x40, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x90, 0x03, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x74, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x03, 0x00, 0x00, 0x00, 0x0b, 0x00, 0x00, 0x00, 0x04, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0xd3, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x70, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x04, 0x04, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x20, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x03, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x04, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x08, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0xef, 0x00, 0x00, 0x00, 0x0b, 0x00, 0x00, 0x70,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x28, 0x04, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x10, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x08, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x08, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0xb0, 0x00, 0x00, 0x00,
0x09, 0x00, 0x00, 0x00, 0x40, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x38, 0x04, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x10, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x03, 0x00, 0x00, 0x00, 0x04, 0x00, 0x00, 0x00, 0x08, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x10, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x6d, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00, 0x42, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x48, 0x04, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x7c, 0x01, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x0b, 0x00, 0x00, 0x00,
0x04, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x32, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00,
0x06, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x06, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x02, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x03, 0x00, 0x00, 0x00,
0x06, 0x00, 0x00, 0x0e, 0x80, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x06, 0x00, 0x00, 0x00,
0x05, 0x00, 0x00, 0x00, 0x00, 0x0b, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0xa8, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0xa8, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x08, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00, 0x05, 0x00, 0x00, 0x00,
0x48, 0x04, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0xb8, 0x03, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0xb8, 0x03, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x08, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x01, 0x00, 0x00, 0x00, 0x05, 0x00, 0x00, 0x00, 0x00, 0x0b, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0xa8, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0xa8, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x08, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00};
unsignedchar vector_add_int_fatbin[] = {
0x50, 0xed, 0x55, 0xba, 0x01, 0x00, 0x10, 0x00, 0xf0, 0x0b, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x02, 0x00, 0x01, 0x01, 0x48, 0x00, 0x00, 0x00,
0xa8, 0x0b, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x07, 0x00, 0x01, 0x00, 0x56, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x11, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x7f, 0x45, 0x4c, 0x46, 0x02, 0x01, 0x01, 0x33,
0x07, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x02, 0x00, 0xbe, 0x00,
0x7c, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x0b, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x08, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x56, 0x05, 0x50, 0x00, 0x40, 0x00, 0x38, 0x00,
0x03, 0x00, 0x40, 0x00, 0x0c, 0x00, 0x01, 0x00, 0x00, 0x2e, 0x73, 0x68,
0x73, 0x74, 0x72, 0x74, 0x61, 0x62, 0x00, 0x2e, 0x73, 0x74, 0x72, 0x74,
0x61, 0x62, 0x00, 0x2e, 0x73, 0x79, 0x6d, 0x74, 0x61, 0x62, 0x00, 0x2e,
0x73, 0x79, 0x6d, 0x74, 0x61, 0x62, 0x5f, 0x73, 0x68, 0x6e, 0x64, 0x78,
0x00, 0x2e, 0x6e, 0x76, 0x2e, 0x69, 0x6e, 0x66, 0x6f, 0x00, 0x2e, 0x74,
0x65, 0x78, 0x74, 0x2e, 0x76, 0x65, 0x63, 0x74, 0x6f, 0x72, 0x5f, 0x61,
0x64, 0x64, 0x00, 0x2e, 0x6e, 0x76, 0x2e, 0x69, 0x6e, 0x66, 0x6f, 0x2e,
0x76, 0x65, 0x63, 0x74, 0x6f, 0x72, 0x5f, 0x61, 0x64, 0x64, 0x00, 0x2e,
0x6e, 0x76, 0x2e, 0x73, 0x68, 0x61, 0x72, 0x65, 0x64, 0x2e, 0x76, 0x65,
0x63, 0x74, 0x6f, 0x72, 0x5f, 0x61, 0x64, 0x64, 0x00, 0x2e, 0x6e, 0x76,
0x2e, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x61, 0x6e, 0x74, 0x30, 0x2e, 0x76,
0x65, 0x63, 0x74, 0x6f, 0x72, 0x5f, 0x61, 0x64, 0x64, 0x00, 0x2e, 0x72,
0x65, 0x6c, 0x2e, 0x6e, 0x76, 0x2e, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x61,
0x6e, 0x74, 0x30, 0x2e, 0x76, 0x65, 0x63, 0x74, 0x6f, 0x72, 0x5f, 0x61,
0x64, 0x64, 0x00, 0x2e, 0x64, 0x65, 0x62, 0x75, 0x67, 0x5f, 0x66, 0x72,
0x61, 0x6d, 0x65, 0x00, 0x2e, 0x72, 0x65, 0x6c, 0x2e, 0x64, 0x65, 0x62,
0x75, 0x67, 0x5f, 0x66, 0x72, 0x61, 0x6d, 0x65, 0x00, 0x2e, 0x72, 0x65,
0x6c, 0x61, 0x2e, 0x64, 0x65, 0x62, 0x75, 0x67, 0x5f, 0x66, 0x72, 0x61,
0x6d, 0x65, 0x00, 0x2e, 0x6e, 0x76, 0x2e, 0x63, 0x61, 0x6c, 0x6c, 0x67,
0x72, 0x61, 0x70, 0x68, 0x00, 0x2e, 0x6e, 0x76, 0x2e, 0x70, 0x72, 0x6f,
0x74, 0x6f, 0x74, 0x79, 0x70, 0x65, 0x00, 0x2e, 0x6e, 0x76, 0x2e, 0x72,
0x65, 0x6c, 0x2e, 0x61, 0x63, 0x74, 0x69, 0x6f, 0x6e, 0x00, 0x00, 0x2e,
0x73, 0x68, 0x73, 0x74, 0x72, 0x74, 0x61, 0x62, 0x00, 0x2e, 0x73, 0x74,
0x72, 0x74, 0x61, 0x62, 0x00, 0x2e, 0x73, 0x79, 0x6d, 0x74, 0x61, 0x62,
0x00, 0x2e, 0x73, 0x79, 0x6d, 0x74, 0x61, 0x62, 0x5f, 0x73, 0x68, 0x6e,
0x64, 0x78, 0x00, 0x2e, 0x6e, 0x76, 0x2e, 0x69, 0x6e, 0x66, 0x6f, 0x00,
0x2e, 0x74, 0x65, 0x78, 0x74, 0x2e, 0x76, 0x65, 0x63, 0x74, 0x6f, 0x72,
0x5f, 0x61, 0x64, 0x64, 0x00, 0x2e, 0x6e, 0x76, 0x2e, 0x69, 0x6e, 0x66,
0x6f, 0x2e, 0x76, 0x65, 0x63, 0x74, 0x6f, 0x72, 0x5f, 0x61, 0x64, 0x64,
0x00, 0x2e, 0x6e, 0x76, 0x2e, 0x73, 0x68, 0x61, 0x72, 0x65, 0x64, 0x2e,
0x76, 0x65, 0x63, 0x74, 0x6f, 0x72, 0x5f, 0x61, 0x64, 0x64, 0x00, 0x2e,
0x72, 0x65, 0x6c, 0x2e, 0x6e, 0x76, 0x2e, 0x63, 0x6f, 0x6e, 0x73, 0x74,
0x61, 0x6e, 0x74, 0x30, 0x2e, 0x76, 0x65, 0x63, 0x74, 0x6f, 0x72, 0x5f,
0x61, 0x64, 0x64, 0x00, 0x2e, 0x6e, 0x76, 0x2e, 0x63, 0x6f, 0x6e, 0x73,
0x74, 0x61, 0x6e, 0x74, 0x30, 0x2e, 0x76, 0x65, 0x63, 0x74, 0x6f, 0x72,
0x5f, 0x61, 0x64, 0x64, 0x00, 0x2e, 0x64, 0x65, 0x62, 0x75, 0x67, 0x5f,
0x66, 0x72, 0x61, 0x6d, 0x65, 0x00, 0x2e, 0x72, 0x65, 0x6c, 0x2e, 0x64,
0x65, 0x62, 0x75, 0x67, 0x5f, 0x66, 0x72, 0x61, 0x6d, 0x65, 0x00, 0x2e,
0x72, 0x65, 0x6c, 0x61, 0x2e, 0x64, 0x65, 0x62, 0x75, 0x67, 0x5f, 0x66,
0x72, 0x61, 0x6d, 0x65, 0x00, 0x2e, 0x6e, 0x76, 0x2e, 0x63, 0x61, 0x6c,
0x6c, 0x67, 0x72, 0x61, 0x70, 0x68, 0x00, 0x2e, 0x6e, 0x76, 0x2e, 0x70,
0x72, 0x6f, 0x74, 0x6f, 0x74, 0x79, 0x70, 0x65, 0x00, 0x2e, 0x6e, 0x76,
0x2e, 0x72, 0x65, 0x6c, 0x2e, 0x61, 0x63, 0x74, 0x69, 0x6f, 0x6e, 0x00,
0x76, 0x65, 0x63, 0x74, 0x6f, 0x72, 0x5f, 0x61, 0x64, 0x64, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x32, 0x00, 0x00, 0x00, 0x03, 0x00, 0x0b, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x8a, 0x00, 0x00, 0x00, 0x03, 0x00, 0x0a, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0xa3, 0x00, 0x00, 0x00, 0x03, 0x00, 0x04, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0xd3, 0x00, 0x00, 0x00, 0x03, 0x00, 0x07, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0xef, 0x00, 0x00, 0x00, 0x03, 0x00, 0x08, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0xfe, 0x00, 0x00, 0x00, 0x12, 0x10, 0x0b, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x02, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0xff, 0xff, 0xff, 0xff, 0x24, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0x03, 0x00, 0x04, 0x7c,
0xff, 0xff, 0xff, 0xff, 0x0f, 0x0c, 0x81, 0x80, 0x80, 0x28, 0x00, 0x08,
0xff, 0x81, 0x80, 0x28, 0x08, 0x81, 0x80, 0x80, 0x28, 0x00, 0x00, 0x00,
0xff, 0xff, 0xff, 0xff, 0x34, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x02, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x04, 0x04, 0x00, 0x00, 0x00, 0x04, 0x14, 0x00, 0x00, 0x00, 0x0c, 0x81,
0x80, 0x80, 0x28, 0x00, 0x04, 0x34, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x04, 0x2f, 0x08, 0x00, 0x06, 0x00, 0x00, 0x00,
0x0e, 0x00, 0x00, 0x00, 0x04, 0x12, 0x08, 0x00, 0x06, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x04, 0x11, 0x08, 0x00, 0x06, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x04, 0x12, 0x08, 0x00, 0x06, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x04, 0x37, 0x04, 0x00, 0x7c, 0x00, 0x00, 0x00,
0x01, 0x35, 0x00, 0x00, 0x04, 0x0a, 0x08, 0x00, 0x02, 0x00, 0x00, 0x00,
0x60, 0x01, 0x1c, 0x00, 0x03, 0x19, 0x1c, 0x00, 0x04, 0x17, 0x0c, 0x00,
0x00, 0x00, 0x00, 0x00, 0x03, 0x00, 0x18, 0x00, 0x00, 0xf0, 0x11, 0x00,
0x04, 0x17, 0x0c, 0x00, 0x00, 0x00, 0x00, 0x00, 0x02, 0x00, 0x10, 0x00,
0x00, 0xf0, 0x21, 0x00, 0x04, 0x17, 0x0c, 0x00, 0x00, 0x00, 0x00, 0x00,
0x01, 0x00, 0x08, 0x00, 0x00, 0xf0, 0x21, 0x00, 0x04, 0x17, 0x0c, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0xf0, 0x21, 0x00,
0x03, 0x1b, 0xff, 0x00, 0x04, 0x1c, 0x08, 0x00, 0x50, 0x00, 0x00, 0x00,
0x30, 0x01, 0x00, 0x00, 0x04, 0x1e, 0x04, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0xff, 0xff, 0xff, 0xff, 0x00, 0x00, 0x00, 0x00,
0xfe, 0xff, 0xff, 0xff, 0x00, 0x00, 0x00, 0x00, 0xfd, 0xff, 0xff, 0xff,
0x00, 0x00, 0x00, 0x00, 0xfc, 0xff, 0xff, 0xff, 0x00, 0x00, 0x00, 0x00,
0x73, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x11,
0x25, 0x00, 0x05, 0x36, 0x44, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x02, 0x00, 0x00, 0x00, 0x06, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x02, 0x7a, 0x01, 0x00, 0x00, 0x0a, 0x00, 0x00,
0x00, 0x0f, 0x00, 0x00, 0x00, 0xe4, 0x0f, 0x00, 0x19, 0x79, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x25, 0x00, 0x00, 0x00, 0x28, 0x0e, 0x00,
0x19, 0x79, 0x03, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x21, 0x00, 0x00,
0x00, 0x24, 0x0e, 0x00, 0x24, 0x7a, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x03, 0x02, 0x8e, 0x07, 0x00, 0xca, 0x1f, 0x00, 0x0c, 0x7a, 0x00, 0x00,
0x00, 0x5e, 0x00, 0x00, 0x70, 0x60, 0xf0, 0x03, 0x00, 0xda, 0x0f, 0x00,
0x4d, 0x09, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x80, 0x03,
0x00, 0xea, 0x0f, 0x00, 0xb9, 0x7a, 0x04, 0x00, 0x00, 0x46, 0x00, 0x00,
0x00, 0x0a, 0x00, 0x00, 0x00, 0xe4, 0x0f, 0x00, 0x02, 0x78, 0x07, 0x00,
0x04, 0x00, 0x00, 0x00, 0x00, 0x0f, 0x00, 0x00, 0x00, 0xca, 0x1f, 0x00,
0x25, 0x76, 0x02, 0x00, 0x00, 0x58, 0x00, 0x00, 0x07, 0x00, 0x8e, 0x07,
0x00, 0xc8, 0x0f, 0x00, 0x25, 0x76, 0x04, 0x00, 0x00, 0x5a, 0x00, 0x00,
0x07, 0x00, 0x8e, 0x07, 0x00, 0xe4, 0x0f, 0x0c, 0x81, 0x79, 0x03, 0x02,
0x04, 0x00, 0x00, 0x00, 0x00, 0x19, 0x1e, 0x0c, 0x00, 0xa8, 0x0e, 0x00,
0x81, 0x79, 0x04, 0x04, 0x04, 0x00, 0x00, 0x00, 0x00, 0x19, 0x1e, 0x0c,
0x00, 0xa2, 0x0e, 0x00, 0x25, 0x76, 0x06, 0x00, 0x00, 0x5c, 0x00, 0x00,
0x07, 0x00, 0x8e, 0x07, 0x00, 0xe2, 0x0f, 0x00, 0x02, 0x7a, 0x0b, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x0f, 0x00, 0x00, 0x00, 0xca, 0x0f, 0x00,
0x24, 0x7a, 0x00, 0x0b, 0x00, 0x03, 0x00, 0x00, 0x00, 0x02, 0x8e, 0x07,
0x00, 0xca, 0x0f, 0x00, 0x0c, 0x7a, 0x00, 0x00, 0x00, 0x5e, 0x00, 0x00,
0x70, 0x60, 0xf0, 0x03, 0x00, 0xe4, 0x0f, 0x00, 0x10, 0x72, 0x09, 0x04,
0x03, 0x00, 0x00, 0x00, 0xff, 0xe0, 0xff, 0x07, 0x00, 0xca, 0x4f, 0x00,
0x86, 0x79, 0x00, 0x06, 0x09, 0x00, 0x00, 0x00, 0x04, 0x19, 0x10, 0x0c,
0x00, 0xec, 0x01, 0x00, 0x47, 0x89, 0x00, 0x00, 0x40, 0xff, 0xff, 0xff,
0xff, 0xff, 0x83, 0x03, 0x00, 0xea, 0x0f, 0x00, 0x4d, 0x79, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x80, 0x03, 0x00, 0xea, 0x0f, 0x00,
0x47, 0x79, 0x00, 0x00, 0xf0, 0xff, 0xff, 0xff, 0xff, 0xff, 0x83, 0x03,
0x00, 0xc0, 0x0f, 0x00, 0x18, 0x79, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0xc0, 0x0f, 0x00, 0x18, 0x79, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0xc0, 0x0f, 0x00,
0x18, 0x79, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0xc0, 0x0f, 0x00, 0x18, 0x79, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0xc0, 0x0f, 0x00, 0x18, 0x79, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0xc0, 0x0f, 0x00,
0x18, 0x79, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0xc0, 0x0f, 0x00, 0x18, 0x79, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0xc0, 0x0f, 0x00, 0x18, 0x79, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0xc0, 0x0f, 0x00,
0x18, 0x79, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0xc0, 0x0f, 0x00, 0x18, 0x79, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0xc0, 0x0f, 0x00, 0x18, 0x79, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0xc0, 0x0f, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00, 0x03, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x40, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0xfe, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x0b, 0x00, 0x00, 0x00,
0x03, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x3e, 0x01, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x09, 0x01, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x13, 0x00, 0x00, 0x00, 0x02, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x48, 0x02, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0xa8, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x02, 0x00, 0x00, 0x00, 0x06, 0x00, 0x00, 0x00,
0x08, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x18, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0xa3, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0xf0, 0x02, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x70, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x29, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x70, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x60, 0x03, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x30, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x03, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x04, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x43, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x70, 0x40, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x90, 0x03, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x74, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x03, 0x00, 0x00, 0x00, 0x0b, 0x00, 0x00, 0x00,
0x04, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0xd3, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x70,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x04, 0x04, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x20, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x03, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x04, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x08, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0xef, 0x00, 0x00, 0x00,
0x0b, 0x00, 0x00, 0x70, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x28, 0x04, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x10, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x08, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x08, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0xb0, 0x00, 0x00, 0x00, 0x09, 0x00, 0x00, 0x00, 0x40, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x38, 0x04, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x10, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x03, 0x00, 0x00, 0x00, 0x04, 0x00, 0x00, 0x00,
0x08, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x10, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x6d, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00,
0x42, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x48, 0x04, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x7c, 0x01, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x0b, 0x00, 0x00, 0x00, 0x04, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x32, 0x00, 0x00, 0x00,
0x01, 0x00, 0x00, 0x00, 0x06, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x06, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x02, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x03, 0x00, 0x00, 0x00, 0x06, 0x00, 0x00, 0x0e, 0x80, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x06, 0x00, 0x00, 0x00, 0x05, 0x00, 0x00, 0x00, 0x00, 0x0b, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0xa8, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0xa8, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x08, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00,
0x05, 0x00, 0x00, 0x00, 0x48, 0x04, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0xb8, 0x03, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0xb8, 0x03, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x08, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00, 0x05, 0x00, 0x00, 0x00,
0x00, 0x0b, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0xa8, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0xa8, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x08, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00};
最后，这些PTX、CUBIN和FATBIN字节码可以使用cuModuleLoadData函数加载到CUDA driver API中，类似于我们之前使用的cuModuleLoad函数。
// Build command
// g++ run_vector_add_sm86.cpp vector_add_ptx_cubin_fatbin_sm86.cpp
// -I/usr/local/cuda/include -L/usr/local/cuda/lib64 -lcuda -o
// run_vector_add_sm86
// Using CUDA driver API is sufficient.
#include <cuda.h>
#include <fstream>
#include <iostream>
#include <string>
#include <vector>
#include "vector_add_ptx_cubin_fatbin_sm86.hpp"
#define CHECK_CUDA_DRIVER_ERROR(val) \
check_driver((val), #val, __FILE__, __LINE__)
void check_driver(CUresult err, char const* func, char const* file, int line)
{
if (err != CUDA_SUCCESS)
{
std::cerr << "CUDA Driver Error at: " << file << ":" << line
<< std::endl;
charconst* err_string{nullptr};
std::cerr << cuGetErrorString(err, &err_string) << std::endl;
std::cerr << err_string << " " << func << std::endl;
std::exit(EXIT_FAILURE);
}
}
int main(int argc, char** argv)
{
// 用户可以指定是否使用PTX、CUBIN或FATBIN字节码
if (argc != 2)
{
std::cerr << "Usage: " << argv[0] << " ptx|cubin|fatbin" << std::endl;
return1;
}
// 声明CUDA driver API所需的各种句柄和对象
CUdevice cuda_device{0}; // 使用CUDA设备0
CUfunction vector_add_kernel{}; // CUDA kernel 函数句柄
CUmodule cuda_module{}; // CUDA模块句柄
CUcontext cuda_context{}; // CUDA上下文句柄
CUstream cuda_stream{}; // CUDA流句柄
// 初始化CUDA driver API，目前此标志必须为0
unsignedintconst cuda_driver_init_flags{0};
CHECK_CUDA_DRIVER_ERROR(cuInit(cuda_driver_init_flags));
// 创建CUDA上下文，目前此标志必须为0
unsignedintconst cuda_context_init_flags{0};
CHECK_CUDA_DRIVER_ERROR(
cuCtxCreate(&cuda_context, cuda_context_init_flags, cuda_device));
// 创建非阻塞CUDA流
CHECK_CUDA_DRIVER_ERROR(
cuStreamCreate(&cuda_stream, CU_STREAM_NON_BLOCKING));
// 根据命令行参数从不同格式的字节码创建CUDA模块
if (std::string{argv[1]} == "ptx")
{
// 从PTX字节码加载模块
CHECK_CUDA_DRIVER_ERROR(cuModuleLoadData(
&cuda_module, static_cast<voidconst*>(vector_add_int_ptx)));
}
elseif (std::string{argv[1]} == "cubin")
{
// 从CUBIN字节码加载模块
CHECK_CUDA_DRIVER_ERROR(cuModuleLoadData(
&cuda_module, static_cast<voidconst*>(vector_add_int_cubin)));
}
elseif (std::string{argv[1]} == "fatbin")
{
// 从FATBIN字节码加载模块
CHECK_CUDA_DRIVER_ERROR(cuModuleLoadData(
&cuda_module, static_cast<voidconst*>(vector_add_int_fatbin)));
}
else
{
std::cerr << "Error: The byte type is not supported." << std::endl;
return1;
}
// 从CUDA模块中获取名为"vector_add"的 kernel 函数
CHECK_CUDA_DRIVER_ERROR(
cuModuleGetFunction(&vector_add_kernel, cuda_module, "vector_add"));
// 设置向量元素数量
unsignedint num_elements{8192};
// 在主机端创建输入和输出向量
std::vector<int> host_vector_a(num_elements, 1); // 输入向量A
std::vector<int> host_vector_b(num_elements, 2); // 输入向量B
std::vector<int> host_vector_c(num_elements, -1); // 输出向量C
std::vector<int> host_vector_c_reference(num_elements, -2); // 参考结果向量
// 初始化输入向量，设置为索引值
for (size_t i{0}; i < num_elements; ++i)
{
host_vector_a.at(i) = i;
host_vector_b.at(i) = i;
}
// 计算参考结果，用于验证 kernel 计算的正确性
for (size_t i{0}; i < num_elements; ++i)
{
host_vector_c_reference.at(i) =
host_vector_a.at(i) + host_vector_b.at(i);
}
// 在设备端分配内存
CUdeviceptr device_vector_a{};
CUdeviceptr device_vector_b{};
CUdeviceptr device_vector_c{};
CHECK_CUDA_DRIVER_ERROR(
cuMemAlloc(&device_vector_a, num_elements * sizeof(int)));
CHECK_CUDA_DRIVER_ERROR(
cuMemAlloc(&device_vector_b, num_elements * sizeof(int)));
CHECK_CUDA_DRIVER_ERROR(
cuMemAlloc(&device_vector_c, num_elements * sizeof(int)));
// 将输入向量从主机复制到设备
CHECK_CUDA_DRIVER_ERROR(cuMemcpyHtoD(device_vector_a, host_vector_a.data(),
num_elements * sizeof(int)));
CHECK_CUDA_DRIVER_ERROR(cuMemcpyHtoD(device_vector_b, host_vector_b.data(),
num_elements * sizeof(int)));
// 设置 kernel 参数数组，包含设备内存指针和元素数量
void* kernel_params[]{&device_vector_a, &device_vector_b, &device_vector_c,
&num_elements};
// 配置 kernel 启动参数
unsignedintconst block_size_x{256}; // 线程块大小：256个线程
unsignedintconst block_size_y{1};
unsignedintconst block_size_z{1};
// 计算网格大小，确保覆盖所有元素
unsignedintconst grid_size_x{(num_elements + block_size_x - 1) /
block_size_x};
unsignedintconst grid_size_y{1};
unsignedintconst grid_size_z{1};
unsignedintconst shared_memory_size{0}; // 不使用共享内存
// 启动CUDA kernel
CHECK_CUDA_DRIVER_ERROR(cuLaunchKernel(
vector_add_kernel, grid_size_x, grid_size_y, grid_size_z, block_size_x,
block_size_y, block_size_z, shared_memory_size, cuda_stream,
kernel_params, nullptr));
// 等待 kernel 执行完成
CHECK_CUDA_DRIVER_ERROR(cuStreamSynchronize(cuda_stream));
// 将结果向量从设备复制回主机
CHECK_CUDA_DRIVER_ERROR(cuMemcpyDtoH(host_vector_c.data(), device_vector_c,
num_elements * sizeof(int)));
// 验证计算结果的正确性
for (size_t i{0}; i < num_elements; ++i)
{
if (host_vector_c.at(i) != host_vector_c_reference.at(i))
{
std::cerr << "Error: The result is incorrect." << std::endl;
return1;
}
}
// 清理资源：释放设备内存
CHECK_CUDA_DRIVER_ERROR(cuMemFree(device_vector_a));
CHECK_CUDA_DRIVER_ERROR(cuMemFree(device_vector_b));
CHECK_CUDA_DRIVER_ERROR(cuMemFree(device_vector_c));
// 销毁CUDA流
CHECK_CUDA_DRIVER_ERROR(cuStreamDestroy(cuda_stream));
// 卸载CUDA模块
CHECK_CUDA_DRIVER_ERROR(cuModuleUnload(cuda_module));
// 销毁CUDA上下文
CHECK_CUDA_DRIVER_ERROR(cuCtxDestroy(cuda_context));
return0;
}
要构建程序，请运行以下命令。
$ g++ run_vector_add_sm86.cpp vector_add_ptx_cubin_fatbin_sm86.cpp -I/usr/local/cuda/include -L/usr/local/cuda/lib64 -lcuda -o run_vector_add_sm86
要运行程序，请运行以下命令。应该不会遇到错误。
$ ./run_vector_add_sm86 ptx
$ ./run_vector_add_sm86 cubin
$ ./run_vector_add_sm86 fatbin
参考
-
CUDA Compilation(https://leimao.github.io/blog/CUDA-Compilation/) -
Simple Driver Runtime – CUDA Examples(https://github.com/NVIDIA/cuda-samples/tree/9c688d7ff78455ed42e345124d1495aad6bf66de/Samples/0_Introduction/simpleDrvRuntime)
（文：GiantPandaCV）