GPU 프로그래밍용 고수준 언어 Triton의 내부를 살펴보고, Triton 프로그램이 어떻게 CUDA 커널(정확히는 CUBIN)로 컴파일되는지 단계를 따라 정리합니다. nvcc 기반 CUDA 컴파일 경로, Triton의 MLIR 백엔드 전환, IR→PTX→CUBIN→SASS 추출까지 개요와 예시를 통해 설명합니다.
최근 GPU 프로그래밍을 위한 고수준 언어인 Triton을 살펴보고 있습니다. Triton은 고수준 추상화를 제공하여 고성능 GPU 코드를 더 쉽게 작성할 수 있게 해 주며, 이 코드는 효율적인 CUDA 네이티브 코드로 컴파일됩니다. 이 글에서는 Triton의 내부를 들여다보고, Triton 프로그램이 내부적으로 어떻게 CUDA 커널(정확히는 CUBIN)로 컴파일되는지 살펴보겠습니다. 이에 대한 정보가 꽤 부족하기 때문에, Triton 전체 컴파일 과정의 개요를 잘 정리해 보려 합니다.
이 글은 심층 탐구의 1부입니다. 연재의 다른 글:
Triton으로 들어가기 전에, nvcc를 이용한 CUDA 컴파일 과정을 이해해 두면 좋습니다. 아래 NVIDIA 문서의 도표는 CUDA가 어떻게 실행 가능한 코드로 컴파일되는지 전체 과정을 보여줍니다.
간단한 cuda -> ptx 예제를 보겠습니다. Compiler Explorer를 사용하면 손쉽게 확인할 수 있습니다. 시演용으로 단순한 예제입니다.
1
2
3
4
5
6
7
8
9
10
11
#include <stdio.h>
// Size of array
#define N 1024
// Kernel
__global__ void add_vectors(double *a, double *b, double *c)
{
int id = blockDim.x * blockIdx.x + threadIdx.x;
if(id < N) c[id] = a[id] + b[id];
}
원 논문에서 시작해 보겠습니다. 이 논문은 Triton JIT 과정의 높은 수준 개요를 보여줍니다.
높은 수준의 개념은 여전히 유효하지만, 백엔드 인프라는 2022년 말에 MLIR로 완전히 교체되었습니다.
업데이트된 Triton 컴파일 과정에 대한 문서는 거의 없다시피 합니다. 컴파일 세부를 다루는 가장 좋은 자료라고 생각한 것들은 아래에 있습니다:
코드와 단위 테스트를 읽고 발표들을 시청한 후, 전체 프로세스가 처음부터 끝까지 어떻게 동작하는지에 대한 저의 일반적인 요약은 다음과 같습니다.
이 과정을 거치면 CUDA 커널 코드가 들어 있는 .c와 .h 파일이 생성됩니다.
1
2
3
4
5
6
7
8
9
10
11
12
13
14
#ifndef TT_KERNEL_INCLUDES
#define TT_KERNEL_INCLUDES
#include <cuda.h>
#include <inttypes.h>
#include <stdint.h>
#include <stdio.h>
#endif
void unload_add_kernel_9969bdda_0123(void);
void load_add_kernel_9969bdda_0123(void);
// tt-linker: add_kernel_9969bdda_0123:CUdeviceptr x_ptr, CUdeviceptr y_ptr, CUdeviceptr output_ptr, int32_t n_elements:64_warps1xstages3
CUresult add_kernel_9969bdda_0123(CUstream stream, CUdeviceptr x_ptr, CUdeviceptr y_ptr, CUdeviceptr output_ptr, int32_t n_elements);
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
/* clang-format off */
#include <stdio.h>
#include <stdint.h>
#include <inttypes.h>
#include <string.h>
#include <cuda.h>
// helpers to check for cuda errors
#define CUDA_CHECK(ans) {\
gpuAssert((ans), __FILE__, __LINE__);\
}\
static inline void gpuAssert(CUresult code, const char *file, int line) {
if (code != CUDA_SUCCESS) {
const char *prefix = "Triton Error [CUDA]: ";
const char *str;
cuGetErrorString(code, &str);
char err[1024] = {0};
strcat(err, prefix);
strcat(err, str);
printf("%s\\n", err);
exit(code);
}
}
// globals
#define CUBIN_NAME add_kernel_9969bdda_0123_cubin
CUmodule add_kernel_9969bdda_0123_mod = NULL;
CUfunction add_kernel_9969bdda_0123_func = NULL;
unsigned char CUBIN_NAME[11472] = { 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, 0xc0, 0x15, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x80, 0x11, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x59, 0x05, 0x59, 0x00, 0x40, 0x00, 0x38, 0x00, 0x03, 0x00, 0x40, 0x00, 0x11, 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, 0x61, 0x64, 0x64, 0x5f, 0x6b, 0x65, 0x72, 0x6e, 0x65, 0x6c, 0x00, 0x2e, 0x6e, 0x76, 0x2e, 0x69, 0x6e, 0x66, 0x6f, 0x2e, 0x61, 0x64, 0x64, 0x5f, 0x6b, 0x65, 0x72, 0x6e, 0x65, 0x6c, 0x00, 0x2e, 0x6e, 0x76, 0x2e, 0x73, 0x68, 0x61, 0x72, 0x65, 0x64, 0x2e, 0x61, 0x64, 0x64, 0x5f, 0x6b, 0x65, 0x72, 0x6e, 0x65, 0x6c, 0x00, 0x2e, 0x6e, 0x76, 0x2e, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x61, 0x6e, 0x74, 0x30, 0x2e, 0x61, 0x64, 0x64, 0x5f, 0x6b, 0x65, 0x72, 0x6e, 0x65, 0x6c, 0x00, 0x2e, 0x72, 0x65, 0x6c, 0x2e, 0x6e, 0x76, 0x2e, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x61, 0x6e, 0x74, 0x30, 0x2e, 0x61, 0x64, 0x64, 0x5f, 0x6b, 0x65, 0x72, 0x6e, 0x65, 0x6c, 0x00, 0x2e, 0x64, 0x65, 0x62, 0x75, 0x67, 0x5f, 0x66, 0x72, 0x61, 0x6d, 0x65, 0x00, 0x2e, 0x64, 0x65, 0x62, 0x75, 0x67, 0x5f, 0x6c, 0x69, 0x6e, 0x65, 0x00, 0x2e, 0x72, 0x65, 0x6c, 0x2e, 0x64, 0x65, 0x62, 0x75, 0x67, 0x5f, 0x6c, 0x69, 0x6e, 0x65, 0x00, 0x2e, 0x6e, 0x76, 0x5f, 0x64, 0x65, 0x62, 0x75, 0x67, 0x5f, 0x6c, 0x69, 0x6e, 0x65, 0x5f, 0x73, 0x61, 0x73, 0x73, 0x00, 0x2e, 0x72, 0x65, 0x6c, 0x2e, 0x6e, 0x76, 0x5f, 0x64, 0x65, 0x62, 0x75, 0x67, 0x5f, 0x6c, 0x69, 0x6e, 0x65, 0x5f, 0x73, 0x61, 0x73, 0x73, 0x00, 0x2e, 0x6e, 0x76, 0x5f, 0x64, 0x65, 0x62, 0x75, 0x67, 0x5f, 0x70, 0x74, 0x78, 0x5f, 0x74, 0x78, 0x74, 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, 0x61, 0x64, 0x64, 0x5f, 0x6b, 0x65, 0x72, 0x6e, 0x65, 0x6c, 0x00, 0x2e, 0x6e, 0x76, 0x2e, 0x69, 0x6e, 0x66, 0x6f, 0x2e, 0x61, 0x64, 0x64, 0x5f, 0x6b, 0x65, 0x72, 0x6e, 0x65, 0x6c, 0x00, 0x2e, 0x6e, 0x76, 0x2e, 0x73, 0x68, 0x61, 0x72, 0x65, 0x64, 0x2e, 0x61, 0x64, 0x64, 0x5f, 0x6b, 0x65, 0x72, 0x6e, 0x65, 0x6c, 0x00, 0x2e, 0x72, 0x65, 0x6c, 0x2e, 0x6e, 0x76, 0x2e, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x61, 0x6e, 0x74, 0x30, 0x2e, 0x61, 0x64, 0x64, 0x5f, 0x6b, 0x65, 0x72, 0x6e, 0x65, 0x6c, 0x00, 0x2e, 0x6e, 0x76, 0x2e, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x61, 0x6e, 0x74, 0x30, 0x2e, 0x61, 0x64, 0x64, 0x5f, 0x6b, 0x65, 0x72, 0x6e, 0x65, 0x6c, 0x00, 0x2e, 0x64, 0x65, 0x62, 0x75, 0x67, 0x5f, 0x66, 0x72, 0x61, 0x6d, 0x65, 0x00, 0x2e, 0x64, 0x65, 0x62, 0x75, 0x67, 0x5f, 0x6c, 0x69, 0x6e, 0x65, 0x00, 0x2e, 0x72, 0x65, 0x6c, 0x2e, 0x64, 0x65, 0x62, 0x75, 0x67, 0x5f, 0x6c, 0x69, 0x6e, 0x65, 0x00, 0x2e, 0x6e, 0x76, 0x5f, 0x64, 0x65, 0x62, 0x75, 0x67, 0x5f, 0x6c, 0x69, 0x6e, 0x65, 0x5f, 0x73, 0x61, 0x73, 0x73, 0x00, 0x2e, 0x72, 0x65, 0x6c, 0x2e, 0x6e, 0x76, 0x5f, 0x64, 0x65, 0x62, 0x75, 0x67, 0x5f, 0x6c, 0x69, 0x6e, 0x65, 0x5f, 0x73, 0x61, 0x73, 0x73, 0x00, 0x2e, 0x6e, 0x76, 0x5f, 0x64, 0x65, 0x62, 0x75, 0x67, 0x5f, 0x70, 0x74, 0x78, 0x5f, 0x74, 0x78, 0x74, 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, 0x61, 0x64, 0x64, 0x5f, 0x6b, 0x65, 0x72, 0x6e, 0x65, 0x6c, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 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, 0x10, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x8a, 0x00, 0x00, 0x00, 0x03, 0x00, 0x0f, 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, 0xb0, 0x00, 0x00, 0x00, 0x03, 0x00, 0x05, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0xcc, 0x00, 0x00, 0x00, 0x03, 0x00, 0x06, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0xf8, 0x00, 0x00, 0x00, 0x03, 0x00, 0x07, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x2d, 0x01, 0x00, 0x00, 0x03, 0x00, 0x0a, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x49, 0x01, 0x00, 0x00, 0x03, 0x00, 0x0b, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x58, 0x01, 0x00, 0x00, 0x12, 0x10, 0x10, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x80, 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, 0x80, 0x02, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x04, 0x04, 0x00, 0x00, 0x00, 0x04, 0x58, 0x00, 0x00, 0x00, 0x0c, 0x81, 0x80, 0x80, 0x28, 0x00, 0x04, 0x04, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0xb5, 0x00, 0x00, 0x00, 0x02, 0x00, 0x6c, 0x00, 0x00, 0x00, 0x01, 0x01, 0xfb, 0x0e, 0x0a, 0x00, 0x01, 0x01, 0x01, 0x01, 0x00, 0x00, 0x00, 0x01, 0x2f, 0x68, 0x6f, 0x6d, 0x65, 0x2f, 0x6b, 0x73, 0x68, 0x61, 0x72, 0x6d, 0x61, 0x2f, 0x64, 0x65, 0x76, 0x2f, 0x67, 0x69, 0x74, 0x2f, 0x6c, 0x65, 0x61, 0x72, 0x6e, 0x69, 0x6e, 0x67, 0x2d, 0x74, 0x72, 0x69, 0x74, 0x6f, 0x6e, 0x2f, 0x2e, 0x2e, 0x2f, 0x6d, 0x6c, 0x2d, 0x70, 0x72, 0x6f, 0x6a, 0x65, 0x63, 0x74, 0x73, 0x2f, 0x74, 0x72, 0x69, 0x74, 0x6f, 0x6e, 0x5f, 0x69, 0x6e, 0x74, 0x65, 0x72, 0x6e, 0x61, 0x6c, 0x73, 0x00, 0x00, 0x76, 0x65, 0x63, 0x74, 0x6f, 0x72, 0x5f, 0x61, 0x64, 0x64, 0x2e, 0x70, 0x79, 0x00, 0x01, 0x95, 0x92, 0xb1, 0xb5, 0x06, 0xde, 0x0e, 0x00, 0x00, 0x09, 0x02, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x04, 0x01, 0x03, 0x08, 0x01, 0x03, 0x0f, 0x02, 0x10, 0x01, 0xf4, 0x03, 0x75, 0x02, 0x30, 0x01, 0x03, 0x0b, 0x02, 0x10, 0x01, 0xea, 0x03, 0x02, 0x02, 0x20, 0x01, 0xf3, 0x03, 0x7a, 0x02, 0x10, 0x01, 0xf1, 0xf2, 0x03, 0x01, 0x02, 0x20, 0x01, 0xee, 0xf0, 0xf0, 0xf1, 0x03, 0x7e, 0x02, 0x20, 0x01, 0xf1, 0x02, 0xa0, 0x02, 0x00, 0x01, 0x01, 0x7b, 0x00, 0x00, 0x00, 0x02, 0x00, 0x10, 0x00, 0x00, 0x00, 0x01, 0x01, 0xfb, 0x0e, 0x0a, 0x00, 0x01, 0x01, 0x01, 0x01, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00, 0x09, 0x02, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x04, 0x00, 0x03, 0x11, 0x01, 0x03, 0x15, 0x02, 0x10, 0x01, 0x03, 0x0a, 0x02, 0x10, 0x01, 0xf4, 0x03, 0x5c, 0x02, 0x10, 0x01, 0x03, 0x0e, 0x02, 0x10, 0x01, 0x03, 0x1b, 0x02, 0x10, 0x01, 0x03, 0x6d, 0x02, 0x10, 0x01, 0xf1, 0xf2, 0x03, 0x12, 0x02, 0x10, 0x01, 0x03, 0x6c, 0x02, 0x10, 0x01, 0xf2, 0xf2, 0xf4, 0x03, 0x0e, 0x02, 0x10, 0x01, 0x03, 0x77, 0x02, 0x10, 0x01, 0x03, 0x0e, 0x02, 0x10, 0x01, 0xf3, 0xf2, 0xf4, 0x03, 0x79, 0x02, 0x10, 0x01, 0x03, 0x0b, 0x02, 0x10, 0x01, 0x03, 0x03, 0x02, 0x20, 0x01, 0x02, 0x80, 0x02, 0x00, 0x01, 0x01, 0x00, 0x00, 0x00, 0x00, 0x2e, 0x76, 0x65, 0x72, 0x73, 0x69, 0x6f, 0x6e, 0x20, 0x38, 0x2e, 0x34, 0x00, 0x2e, 0x74, 0x61, 0x72, 0x67, 0x65, 0x74, 0x20, 0x73, 0x6d, 0x5f, 0x38, 0x39, 0x00, 0x2e, 0x61, 0x64, 0x64, 0x72, 0x65, 0x73, 0x73, 0x5f, 0x73, 0x69, 0x7a, 0x65, 0x20, 0x36, 0x34, 0x00, 0x00, 0x00, 0x00, 0x2e, 0x76, 0x69, 0x73, 0x69, 0x62, 0x6c, 0x65, 0x20, 0x2e, 0x65, 0x6e, 0x74, 0x72, 0x79, 0x20, 0x61, 0x64, 0x64, 0x5f, 0x6b, 0x65, 0x72, 0x6e, 0x65, 0x6c, 0x28, 0x00, 0x2e, 0x70, 0x61, 0x72, 0x61, 0x6d, 0x20, 0x2e, 0x75, 0x36, 0x34, 0x20, 0x61, 0x64, 0x64, 0x5f, 0x6b, 0x65, 0x72, 0x6e, 0x65, 0x6c, 0x5f, 0x70, 0x61, 0x72, 0x61, 0x6d, 0x5f, 0x30, 0x2c, 0x00, 0x2e, 0x70, 0x61, 0x72, 0x61, 0x6d, 0x20, 0x2e, 0x75, 0x36, 0x34, 0x20, 0x61, 0x64, 0x64, 0x5f, 0x6b, 0x65, 0x72, 0x6e, 0x65, 0x6c, 0x5f, 0x70, 0x61, 0x72, 0x61, 0x6d, 0x5f, 0x31, 0x2c, 0x00, 0x2e, 0x70, 0x61, 0x72, 0x61, 0x6d, 0x20, 0x2e, 0x75, 0x36, 0x34, 0x20, 0x61, 0x64, 0x64, 0x5f, 0x6b, 0x65, 0x72, 0x6e, 0x65, 0x6c, 0x5f, 0x70, 0x61, 0x72, 0x61, 0x6d, 0x5f, 0x32, 0x2c, 0x00, 0x2e, 0x70, 0x61, 0x72, 0x61, 0x6d, 0x20, 0x2e, 0x75, 0x33, 0x32, 0x20, 0x61, 0x64, 0x64, 0x5f, 0x6b, 0x65, 0x72, 0x6e, 0x65, 0x6c, 0x5f, 0x70, 0x61, 0x72, 0x61, 0x6d, 0x5f, 0x33, 0x00, 0x29, 0x00, 0x2e, 0x72, 0x65, 0x71, 0x6e, 0x74, 0x69, 0x64, 0x20, 0x33, 0x32, 0x2c, 0x20, 0x31, 0x2c, 0x20, 0x31, 0x00, 0x7b, 0x00, 0x2e, 0x72, 0x65, 0x67, 0x20, 0x2e, 0x70, 0x72, 0x65, 0x64, 0x20, 0x09, 0x25, 0x70, 0x3c, 0x37, 0x3e, 0x3b, 0x00, 0x2e, 0x72, 0x65, 0x67, 0x20, 0x2e, 0x62, 0x33, 0x32, 0x20, 0x09, 0x25, 0x72, 0x3c, 0x31, 0x34, 0x3e, 0x3b, 0x00, 0x2e, 0x72, 0x65, 0x67, 0x20, 0x2e, 0x66, 0x33, 0x32, 0x20, 0x09, 0x25, 0x66, 0x3c, 0x37, 0x3e, 0x3b, 0x00, 0x2e, 0x72, 0x65, 0x67, 0x20, 0x2e, 0x62, 0x36, 0x34, 0x20, 0x09, 0x25, 0x72, 0x64, 0x3c, 0x31, 0x31, 0x3e, 0x3b, 0x00, 0x00, 0x24, 0x4c, 0x5f, 0x5f, 0x66, 0x75, 0x6e, 0x63, 0x5f, 0x62, 0x65, 0x67, 0x69, 0x6e, 0x30, 0x3a, 0x00, 0x00, 0x00, 0x6c, 0x64, 0x2e, 0x70, 0x61, 0x72, 0x61, 0x6d, 0x2e, 0x75, 0x36, 0x34, 0x20, 0x09, 0x25, 0x72, 0x64, 0x37, 0x2c, 0x20, 0x5b, 0x61, 0x64, 0x64, 0x5f, 0x6b, 0x65, 0x72, 0x6e, 0x65, 0x6c, 0x5f, 0x70, 0x61, 0x72, 0x61, 0x6d, 0x5f, 0x30, 0x5d, 0x3b, 0x00, 0x6c, 0x64, 0x2e, 0x70, 0x61, 0x72, 0x61, 0x6d, 0x2e, 0x75, 0x36, 0x34, 0x20, 0x09, 0x25, 0x72, 0x64, 0x38, 0x2c, 0x20, 0x5b, 0x61, 0x64, 0x64, 0x5f, 0x6b, 0x65, 0x72, 0x6e, 0x65, 0x6c, 0x5f, 0x70, 0x61, 0x72, 0x61, 0x6d, 0x5f, 0x31, 0x5d, 0x3b, 0x00, 0x24, 0x4c, 0x5f, 0x5f, 0x74, 0x6d, 0x70, 0x30, 0x3a, 0x00, 0x00, 0x00, 0x6d, 0x6f, 0x76, 0x2e, 0x75, 0x33, 0x32, 0x20, 0x25, 0x72, 0x31, 0x2c, 0x20, 0x25, 0x63, 0x74, 0x61, 0x69, 0x64, 0x2e, 0x78, 0x3b, 0x00, 0x00, 0x00, 0x73, 0x68, 0x6c, 0x2e, 0x62, 0x33, 0x32, 0x20, 0x09, 0x25, 0x72, 0x38, 0x2c, 0x20, 0x25, 0x72, 0x31, 0x2c, 0x20, 0x36, 0x3b, 0x00, 0x6c, 0x64, 0x2e, 0x70, 0x61, 0x72, 0x61, 0x6d, 0x2e, 0x75, 0x36, 0x34, 0x20, 0x09, 0x25, 0x72, 0x64, 0x39, 0x2c, 0x20, 0x5b, 0x61, 0x64, 0x64, 0x5f, 0x6b, 0x65, 0x72, 0x6e, 0x65, 0x6c, 0x5f, 0x70, 0x61, 0x72, 0x61, 0x6d, 0x5f, 0x32, 0x5d, 0x3b, 0x00, 0x6c, 0x64, 0x2e, 0x70, 0x61, 0x72, 0x61, 0x6d, 0x2e, 0x75, 0x33, 0x32, 0x20, 0x09, 0x25, 0x72, 0x39, 0x2c, 0x20, 0x5b, 0x61, 0x64, 0x64, 0x5f, 0x6b, 0x65, 0x72, 0x6e, 0x65, 0x6c, 0x5f, 0x70, 0x61, 0x72, 0x61, 0x6d, 0x5f, 0x33, 0x5d, 0x3b, 0x00, 0x00, 0x6d, 0x6f, 0x76, 0x2e, 0x75, 0x33, 0x32, 0x20, 0x09, 0x25, 0x72, 0x31, 0x30, 0x2c, 0x20, 0x25, 0x74, 0x69, 0x64, 0x2e, 0x78, 0x3b, 0x00, 0x61, 0x6e, 0x64, 0x2e, 0x62, 0x33, 0x32, 0x20, 0x20, 0x09, 0x25, 0x72, 0x31, 0x31, 0x2c, 0x20, 0x25, 0x72, 0x31, 0x30, 0x2c, 0x20, 0x33, 0x31, 0x3b, 0x00, 0x00, 0x6f, 0x72, 0x2e, 0x62, 0x33, 0x32, 0x20, 0x20, 0x09, 0x25, 0x72, 0x31, 0x32, 0x2c, 0x20, 0x25, 0x72, 0x38, 0x2c, 0x20, 0x25, 0x72, 0x31, 0x31, 0x3b, 0x00, 0x6f, 0x72, 0x2e, 0x62, 0x33, 0x32, 0x20, 0x20, 0x09, 0x25, 0x72, 0x31, 0x33, 0x2c, 0x20, 0x25, 0x72, 0x31, 0x32, 0x2c, 0x20, 0x33, 0x32, 0x3b, 0x00, 0x00, 0x73, 0x65, 0x74, 0x70, 0x2e, 0x6c, 0x74, 0x2e, 0x73, 0x33, 0x32, 0x20, 0x09, 0x25, 0x70, 0x31, 0x2c, 0x20, 0x25, 0x72, 0x31, 0x32, 0x2c, 0x20, 0x25, 0x72, 0x39, 0x3b, 0x00, 0x73, 0x65, 0x74, 0x70, 0x2e, 0x6c, 0x74, 0x2e, 0x73, 0x33, 0x32, 0x20, 0x09, 0x25, 0x70, 0x32, 0x2c, 0x20, 0x25, 0x72, 0x31, 0x33, 0x2c, 0x20, 0x25, 0x72, 0x39, 0x3b, 0x00, 0x00, 0x6d, 0x75, 0x6c, 0x2e, 0x77, 0x69, 0x64, 0x65, 0x2e, 0x73, 0x33, 0x32, 0x20, 0x09, 0x25, 0x72, 0x64, 0x31, 0x30, 0x2c, 0x20, 0x25, 0x72, 0x31, 0x32, 0x2c, 0x20, 0x34, 0x3b, 0x00, 0x61, 0x64, 0x64, 0x2e, 0x73, 0x36, 0x34, 0x20, 0x09, 0x25, 0x72, 0x64, 0x31, 0x2c, 0x20, 0x25, 0x72, 0x64, 0x37, 0x2c, 0x20, 0x25, 0x72, 0x64, 0x31, 0x30, 0x3b, 0x00, 0x61, 0x64, 0x64, 0x2e, 0x73, 0x36, 0x34, 0x20, 0x09, 0x25, 0x72, 0x64, 0x32, 0x2c, 0x20, 0x25, 0x72, 0x64, 0x31, 0x2c, 0x20, 0x31, 0x32, 0x38, 0x3b, 0x00, 0x00, 0x00, 0x6d, 0x6f, 0x76, 0x2e, 0x75, 0x33, 0x32, 0x20, 0x25, 0x72, 0x32, 0x2c, 0x20, 0x30, 0x78, 0x30, 0x3b, 0x00, 0x40, 0x25, 0x70, 0x31, 0x20, 0x6c, 0x64, 0x2e, 0x67, 0x6c, 0x6f, 0x62, 0x61, 0x6c, 0x2e, 0x62, 0x33, 0x32, 0x20, 0x7b, 0x20, 0x25, 0x72, 0x32, 0x20, 0x7d, 0x2c, 0x20, 0x5b, 0x20, 0x25, 0x72, 0x64, 0x31, 0x20, 0x2b, 0x20, 0x30, 0x20, 0x5d, 0x3b, 0x00, 0x00, 0x6d, 0x6f, 0x76, 0x2e, 0x62, 0x33, 0x32, 0x20, 0x09, 0x25, 0x66, 0x31, 0x2c, 0x20, 0x25, 0x72, 0x32, 0x3b, 0x00, 0x00, 0x6d, 0x6f, 0x76, 0x2e, 0x75, 0x33, 0x32, 0x20, 0x25, 0x72, 0x33, 0x2c, 0x20, 0x30, 0x78, 0x30, 0x3b, 0x00, 0x40, 0x25, 0x70, 0x32, 0x20, 0x6c, 0x64, 0x2e, 0x67, 0x6c, 0x6f, 0x62, 0x61, 0x6c, 0x2e, 0x62, 0x33, 0x32, 0x20, 0x7b, 0x20, 0x25, 0x72, 0x33, 0x20, 0x7d, 0x2c, 0x20, 0x5b, 0x20, 0x25, 0x72, 0x64, 0x32, 0x20, 0x2b, 0x20, 0x30, 0x20, 0x5d, 0x3b, 0x00, 0x00, 0x6d, 0x6f, 0x76, 0x2e, 0x62, 0x33, 0x32, 0x20, 0x09, 0x25, 0x66, 0x32, 0x2c, 0x20, 0x25, 0x72, 0x33, 0x3b, 0x00, 0x00, 0x61, 0x64, 0x64, 0x2e, 0x73, 0x36, 0x34, 0x20, 0x09, 0x25, 0x72, 0x64, 0x33, 0x2c, 0x20, 0x25, 0x72, 0x64, 0x38, 0x2c, 0x20, 0x25, 0x72, 0x64, 0x31, 0x30, 0x3b, 0x00, 0x61, 0x64, 0x64, 0x2e, 0x73, 0x36, 0x34, 0x20, 0x09, 0x25, 0x72, 0x64, 0x34, 0x2c, 0x20, 0x25, 0x72, 0x64, 0x33, 0x2c, 0x20, 0x31, 0x32, 0x38, 0x3b, 0x00, 0x00, 0x00, 0x6d, 0x6f, 0x76, 0x2e, 0x75, 0x33, 0x32, 0x20, 0x25, 0x72, 0x34, 0x2c, 0x20, 0x30, 0x78, 0x30, 0x3b, 0x00, 0x40, 0x25, 0x70, 0x31, 0x20, 0x6c, 0x64, 0x2e, 0x67, 0x6c, 0x6f, 0x62, 0x61, 0x6c, 0x2e, 0x62, 0x33, 0x32, 0x20, 0x7b, 0x20, 0x25, 0x72, 0x34, 0x20, 0x7d, 0x2c, 0x20, 0x5b, 0x20, 0x25, 0x72, 0x64, 0x33, 0x20, 0x2b, 0x20, 0x30, 0x20, 0x5d, 0x3b, 0x00, 0x00, 0x6d, 0x6f, 0x76, 0x2e, 0x62, 0x33, 0x32, 0x20, 0x09, 0x25, 0x66, 0x33, 0x2c, 0x20, 0x25, 0x72, 0x34, 0x3b, 0x00, 0x00, 0x6d, 0x6f, 0x76, 0x2e, 0x75, 0x33, 0x32, 0x20, 0x25, 0x72, 0x35, 0x2c, 0x20, 0x30, 0x78, 0x30, 0x3b, 0x00, 0x40, 0x25, 0x70, 0x32, 0x20, 0x6c, 0x64, 0x2e, 0x67, 0x6c, 0x6f, 0x62, 0x61, 0x6c, 0x2e, 0x62, 0x33, 0x32, 0x20, 0x7b, 0x20, 0x25, 0x72, 0x35, 0x20, 0x7d, 0x2c, 0x20, 0x5b, 0x20, 0x25, 0x72, 0x64, 0x34, 0x20, 0x2b, 0x20, 0x30, 0x20, 0x5d, 0x3b, 0x00, 0x00, 0x6d, 0x6f, 0x76, 0x2e, 0x62, 0x33, 0x32, 0x20, 0x09, 0x25, 0x66, 0x34, 0x2c, 0x20, 0x25, 0x72, 0x35, 0x3b, 0x00, 0x00, 0x61, 0x64, 0x64, 0x2e, 0x66, 0x33, 0x32, 0x20, 0x09, 0x25, 0x66, 0x35, 0x2c, 0x20, 0x25, 0x66, 0x31, 0x2c, 0x20, 0x25, 0x66, 0x33, 0x3b, 0x00, 0x61, 0x64, 0x64, 0x2e, 0x66, 0x33, 0x32, 0x20, 0x09, 0x25, 0x66, 0x36, 0x2c, 0x20, 0x25, 0x66, 0x32, 0x2c, 0x20, 0x25, 0x66, 0x34, 0x3b, 0x00, 0x00, 0x61, 0x64, 0x64, 0x2e, 0x73, 0x36, 0x34, 0x20, 0x09, 0x25, 0x72, 0x64, 0x35, 0x2c, 0x20, 0x25, 0x72, 0x64, 0x39, 0x2c, 0x20, 0x25, 0x72, 0x64, 0x31, 0x30, 0x3b, 0x00, 0x61, 0x64, 0x64, 0x2e, 0x73, 0x36, 0x34, 0x20, 0x09, 0x25, 0x72, 0x64, 0x36, 0x2c, 0x20, 0x25, 0x72, 0x64, 0x35, 0x2c, 0x20, 0x31, 0x32, 0x38, 0x3b, 0x00, 0x00, 0x6d, 0x6f, 0x76, 0x2e, 0x62, 0x33, 0x32, 0x20, 0x09, 0x25, 0x72, 0x36, 0x2c, 0x20, 0x25, 0x66, 0x35, 0x3b, 0x00, 0x00, 0x40, 0x25, 0x70, 0x31, 0x20, 0x73, 0x74, 0x2e, 0x67, 0x6c, 0x6f, 0x62, 0x61, 0x6c, 0x2e, 0x62, 0x33, 0x32, 0x20, 0x5b, 0x20, 0x25, 0x72, 0x64, 0x35, 0x20, 0x2b, 0x20, 0x30, 0x20, 0x5d, 0x2c, 0x20, 0x7b, 0x20, 0x25, 0x72, 0x36, 0x20, 0x7d, 0x3b, 0x00, 0x00, 0x6d, 0x6f, 0x76, 0x2e, 0x62, 0x33, 0x32, 0x20, 0x09, 0x25, 0x72, 0x37, 0x2c, 0x20, 0x25, 0x66, 0x36, 0x3b, 0x00, 0x00, 0x40, 0x25, 0x70, 0x32, 0x20, 0x73, 0x74, 0x2e, 0x67, 0x6c, 0x6f, 0x62, 0x61, 0x6c, 0x2e, 0x62, 0x33, 0x32, 0x20, 0x5b, 0x20, 0x25, 0x72, 0x64, 0x36, 0x20, 0x2b, 0x20, 0x30, 0x20, 0x5d, 0x2c, 0x20, 0x7b, 0x20, 0x25, 0x72, 0x37, 0x20, 0x7d, 0x3b, 0x00, 0x00, 0x00, 0x72, 0x65, 0x74, 0x3b, 0x00, 0x24, 0x4c, 0x5f, 0x5f, 0x74, 0x6d, 0x70, 0x31, 0x3a, 0x00, 0x24, 0x4c, 0x5f, 0x5f, 0x66, 0x75, 0x6e, 0x63, 0x5f, 0x65, 0x6e, 0x64, 0x30, 0x3a, 0x00, 0x00, 0x7d, 0x00, 0x00, 0x2e, 0x73, 0x65, 0x63, 0x74, 0x69, 0x6f, 0x6e, 0x09, 0x2e, 0x64, 0x65, 0x62, 0x75, 0x67, 0x5f, 0x61, 0x62, 0x62, 0x72, 0x65, 0x76, 0x00, 0x7b, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x7d, 0x00, 0x2e, 0x73, 0x65, 0x63, 0x74, 0x69, 0x6f, 0x6e, 0x09, 0x2e, 0x64, 0x65, 0x62, 0x75, 0x67, 0x5f, 0x69, 0x6e, 0x66, 0x6f, 0x00, 0x7b, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x7d, 0x00, 0x2e, 0x73, 0x65, 0x63, 0x74, 0x69, 0x6f, 0x6e, 0x09, 0x2e, 0x64, 0x65, 0x62, 0x75, 0x67, 0x5f, 0x6c, 0x6f, 0x63, 0x09, 0x7b, 0x09, 0x7d, 0x00, 0x00, 0x00, 0x00, 0x04, 0x2f, 0x08, 0x00, 0x09, 0x00, 0x00, 0x00, 0x10, 0x00, 0x00, 0x00, 0x04, 0x12, 0x08, 0x00, 0x09, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x04, 0x11, 0x08, 0x00, 0x09, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x04, 0x12, 0x08, 0x00, 0x09, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x04, 0x37, 0x04, 0x00, 0x7c, 0x00, 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, 0x60, 0x01, 0x00, 0x00, 0x80, 0x01, 0x00, 0x00, 0x04, 0x10, 0x0c, 0x00, 0x20, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00, 0x01, 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, 0x79, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x02, 0x00, 0x00, 0x00, 0x09, 0x00, 0x00, 0x00, 0x1d, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x02, 0x00, 0x00, 0x00, 0x09, 0x00, 0x00, 0x00, 0x44, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x02, 0x00, 0x00, 0x00, 0x09, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 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, 0x21, 0x00, 0x00, 0x00, 0x22, 0x0e, 0x00, 0x02, 0x78, 0x0b, 0x00, 0x04, 0x00, 0x00, 0x00, 0x00, 0x0f, 0x00, 0x00, 0x00, 0xe2, 0x0f, 0x00, 0x05, 0x78, 0x06, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0xff, 0x01, 0x00, 0x00, 0xe2, 0x0f, 0x00, 0xb9, 0x7a, 0x04, 0x00, 0x00, 0x46, 0x00, 0x00, 0x00, 0x0a, 0x00, 0x00, 0x00, 0xe2, 0x0f, 0x00, 0x19, 0x79, 0x03, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x25, 0x00, 0x00, 0x00, 0x62, 0x0e, 0x00, 0x05, 0x78, 0x08, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0xff, 0x01, 0x00, 0x00, 0xe2, 0x0f, 0x00, 0x12, 0x78, 0x00, 0x00, 0x1f, 0x00, 0x00, 0x00, 0xff, 0xc0, 0x8e, 0x07, 0x00, 0xc8, 0x1f, 0x00, 0x11, 0x72, 0x00, 0x03, 0x00, 0x00, 0x00, 0x00, 0xff, 0x30, 0x8e, 0x07, 0x00, 0xc8, 0x2f, 0x00, 0x0c, 0x7a, 0x00, 0x00, 0x00, 0x5e, 0x00, 0x00, 0x70, 0x62, 0xf0, 0x03, 0x00, 0xe2, 0x0f, 0x04, 0x25, 0x76, 0x04, 0x00, 0x00, 0x5a, 0x00, 0x00, 0x0b, 0x02, 0x8e, 0x07, 0x00, 0xe2, 0x0f, 0x04, 0x10, 0x78, 0x02, 0x00, 0x20, 0x00, 0x00, 0x00, 0xff, 0xe0, 0xff, 0x07, 0x00, 0xc8, 0x0f, 0x00, 0x0c, 0x7a, 0x00, 0x02, 0x00, 0x5e, 0x00, 0x00, 0x70, 0x62, 0xf2, 0x03, 0x00, 0xe2, 0x0f, 0x00, 0x25, 0x76, 0x02, 0x00, 0x00, 0x58, 0x00, 0x00, 0x0b, 0x02, 0x8e, 0x07, 0x00, 0xcc, 0x0f, 0x00, 0x81, 0x89, 0x06, 0x02, 0x04, 0x00, 0x00, 0x00, 0x00, 0x19, 0x1e, 0x0c, 0x00, 0xa8, 0x0e, 0x00, 0x81, 0x89, 0x07, 0x04, 0x04, 0x00, 0x00, 0x00, 0x00, 0x19, 0x1e, 0x0c, 0x00, 0xa8, 0x0e, 0x00, 0x81, 0x99, 0x08, 0x02, 0x04, 0x80, 0x00, 0x00, 0x00, 0x19, 0x1e, 0x0c, 0x00, 0xe8, 0x0e, 0x00, 0x81, 0x99, 0x09, 0x04, 0x04, 0x80, 0x00, 0x00, 0x00, 0x19, 0x1e, 0x0c, 0x00, 0xe2, 0x0e, 0x00, 0x21, 0x72, 0x0d, 0x07, 0x06, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0xe2, 0x4f, 0x00, 0x25, 0x76, 0x06, 0x00, 0x00, 0x5c, 0x00, 0x00, 0x0b, 0x02, 0x8e, 0x07, 0x00, 0xca, 0x0f, 0x00, 0x86, 0x89, 0x00, 0x06, 0x0d, 0x00, 0x00, 0x00, 0x04, 0x19, 0x10, 0x0c, 0x00, 0xe2, 0x01, 0x00, 0x21, 0x72, 0x09, 0x09, 0x08, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0xe2, 0x8f, 0x00, 0x4d, 0x19, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x80, 0x03, 0x00, 0xec, 0x0f, 0x00, 0x86, 0x79, 0x00, 0x06, 0x09, 0x80, 0x00, 0x00, 0x04, 0x19, 0x10, 0x0c, 0x00, 0xe2, 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, 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, 0x58, 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, 0x0b, 0x00, 0x00, 0x00, 0x03, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x98, 0x01, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x63, 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, 0x00, 0x03, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0xf0, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x02, 0x00, 0x00, 0x00, 0x09, 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, 0x03, 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, 0xb0, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x60, 0x04, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0xb9, 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, 0xcc, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x19, 0x05, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x7f, 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, 0xf8, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x98, 0x05, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x69, 0x06, 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, 0x04, 0x0c, 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, 0x34, 0x0c, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x78, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x03, 0x00, 0x00, 0x00, 0x10, 0x00, 0x00, 0x00, 0x04, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x2d, 0x01, 0x00, 0x00, 0x01, 0x00, 0x00, 0x70, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0xac, 0x0c, 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, 0x49, 0x01, 0x00, 0x00, 0x0b, 0x00, 0x00, 0x70, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0xd0, 0x0c, 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, 0xbc, 0x00, 0x00, 0x00, 0x09, 0x00, 0x00, 0x00, 0x40, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0xe0, 0x0c, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x10, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x03, 0x00, 0x00, 0x00, 0x05, 0x00, 0x00, 0x00, 0x08, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x10, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0xe0, 0x00, 0x00, 0x00, 0x09, 0x00, 0x00, 0x00, 0x40, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0xf0, 0x0c, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x10, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x03, 0x00, 0x00, 0x00, 0x06, 0x00, 0x00, 0x00, 0x08, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x10, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x0a, 0x01, 0x00, 0x00, 0x09, 0x00, 0x00, 0x00, 0x40, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x0d, 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, 0x10, 0x0d, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x7c, 0x01, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x10, 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, 0x0f, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x80, 0x02, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x03, 0x00, 0x00, 0x00, 0x09, 0x00, 0x00, 0x10, 0x80, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x06, 0x00, 0x00, 0x00, 0x05, 0x00, 0x00, 0x00, 0xc0, 0x15, 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, 0x10, 0x0d, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x70, 0x04, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x70, 0x04, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x08, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00, 0x05, 0x00, 0x00, 0x00, 0xc0, 0x15, 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 };
void unload_add_kernel_9969bdda_0123(void) {
CUDA_CHECK(cuModuleUnload(add_kernel_9969bdda_0123_mod));
}
// TODO: some code duplication with `runtime/backend/cuda.c`
void load_add_kernel_9969bdda_0123() {
int dev = 0;
void *bin = (void *)&CUBIN_NAME;
int shared = 0;
CUDA_CHECK(cuModuleLoadData(&add_kernel_9969bdda_0123_mod, bin));
CUDA_CHECK(cuModuleGetFunction(&add_kernel_9969bdda_0123_func, add_kernel_9969bdda_0123_mod, "add_kernel"));
// set dynamic shared memory if necessary
int shared_optin;
CUDA_CHECK(cuDeviceGetAttribute(&shared_optin, CU_DEVICE_ATTRIBUTE_MAX_SHARED_MEMORY_PER_BLOCK_OPTIN, dev));
if (shared > 49152 && shared_optin > 49152) {
CUDA_CHECK(cuFuncSetCacheConfig(add_kernel_9969bdda_0123_func, CU_FUNC_CACHE_PREFER_SHARED));
CUDA_CHECK(cuFuncSetAttribute(add_kernel_9969bdda_0123_func, CU_FUNC_ATTRIBUTE_MAX_DYNAMIC_SHARED_SIZE_BYTES, shared_optin))
}
}
/*
['BLOCK_SIZE=64', 'num_warps=1', 'num_stages=3']
*/
CUresult add_kernel_9969bdda_0123(CUstream stream, CUdeviceptr x_ptr, CUdeviceptr y_ptr, CUdeviceptr output_ptr, int32_t n_elements) {
if (add_kernel_9969bdda_0123_func == NULL)
load_add_kernel_9969bdda_0123();
unsigned int gX = 1024;
unsigned int gY = 1024;
unsigned int gZ = 1024;
void *args[4] = { &x_ptr, &y_ptr, &output_ptr, &n_elements };
// TODO: shared memory
if(gX * gY * gZ > 0)
return cuLaunchKernel(add_kernel_9969bdda_0123_func, gX, gY, gZ, 1 * 32, 1, 1, 0, stream, args, NULL);
}
여기부터는 Triton 컴파일 과정의 일부로 IR이 점진적으로 생성되는 과정을 쉽게 확인할 수 있습니다:
IR이 더 범용적이며 하드웨어 특화된 세부 정보가 없다는 것을 볼 수 있습니다.
이제 IR이 더 구체적이며, compute capability, CTA 수, 워프 수, 워프당 스레드 수 등과 같은 세부 정보가 포함됩니다. 이 IR을 보면 대체로 NVIDIA 비종속적으로 보이지만, triton_gpu.compute-capability처럼 일부 NVIDIA 특화 정보도 있는 듯합니다.
LLVM IR에는 nvvm.annotations, llvm.nvvm.read.ptx.sreg.tid.x() 등 NVIDIA 특화 속성이 있습니다. 여기서 다음 단계는 PTX입니다.
컴파일러 엔지니어가 아닌 사람들이 PTX의 세부까지 파고들 필요가 있는지는 모르겠습니다. 다만 성능 튜닝이 필요하다면 Triton이 필요 시 PTX 코드를 생성할 수 있다는 점을 알아두면 좋습니다.
마지막으로, CUBIN이 필요하다면 compiled_kernel.asm["cubin"]을 확인하면 됩니다. 그러면 GPU에서 실행 가능한 바이너리 코드를 얻을 수 있습니다.
CUBIN을 좀 더 사람이 읽기 쉬운 형식으로 보고 싶다면, 리눅스 유틸리티인 readelf로 cubin 파일을 읽을 수 있습니다.
ELF 헤더에는 바이너리 종류, 머신, 엔트리 포인트 등의 정보가 담겨 있음을 볼 수 있습니다. 다만 CUBIN을 읽을 줄 아는 것이 얼마나 유용한지는 잘 모르겠습니다.
Cubin 자체만으로는 크게 유용하지 않을 수도 있습니다. Triton 도구의 get_sass 메서드를 이용하면 이를 SASS로 변환할 수 있습니다.
이상으로 Triton 내부에 대한 심층 탐구를 마칩니다. Triton으로 간단한 커널을 컴파일한 뒤, 그로부터 PTX, CUBIN, SASS를 추출하는 방법을 살펴보았습니다. 개인적으로는 Triton이 내부에서 어떻게 동작하는지 이해하고자 했고, 많은 것을 배웠습니다.