Skip to content

GitLab

Commit e4ae2422 authored by German Leon's avatar German Leon
Browse files

Initial commit

parents
Show whitespace changes
Inline Side-by-side
app_profiler.py 0 → 100755
View file @ e4ae2422
#!/usr/bin/env python3
import argparse
import os
import re
import time
import common_functions as cf
import common_parameters as cp
def generate_dict(sm_version, input_file_name):
with open(input_file_name, "r") as f:
# dictionary to store the number of allocated registers per static
kernel_reg = {}
kernel_name = "" # temporary variable to store the kernel_name
check_for_register_count = False
# process the input file created by capturing the stderr while compiling the
# application using -Xptxas -v options
for line in f: # for each line in the file
m = re.match(r".*Compiling entry function.*'(\S+)'.*for.*'{}'.*".format(sm_version), line)
if m:
kernel_name = m.group(1)
check_for_register_count = True
m = re.match(r".*Used[ ]+(\d+)[ ]+registers.*", line)
if check_for_register_count and m:
reg_num = m.group(1) # extract register number
if kernel_name not in kernel_reg:
# associate the extracted register number with the kernel name
kernel_reg[kernel_name] = int(reg_num.strip())
else:
print("Warning: {} exists in the kernel_reg dictionary. "
"Skipping this register count.".format(kernel_name))
check_for_register_count = False
return kernel_reg
"""
Function that calls the profiler based on the injection mode
"""
def profiler_caller(gdb_exec, kernel, benchmark_binary, benchmark_args,device,section,kernel_end):
acc_time = 0
acc_time_profiler=0
script = 'env CUDA_VISIBLE_DEVICES={} {} -ex \'py arg0 = {}\' -n -batch -x {}'
benchmark_args_striped = benchmark_args.replace('\\n', '').replace('\\', '')
print ("KERNEL"+kernel)
#init_string = '"file {}; set args {}"'.format(benchmark_binary, benchmark_args_striped)
init_string = '"{};{};file {}; set args {}; break {}"'.format(section,kernel_end,benchmark_binary, benchmark_args_striped,kernel)
profiler_cmd = script.format(device, gdb_exec, init_string, cp.PROFILER_SCRIPT)
print ("Profiler caller")
if cp.DEBUG:
print("PROFILER CMD: {}".format(profiler_cmd))
for i in range(0, cp.MAX_TIMES_TO_PROFILE):
start = time.time()
os.system(profiler_cmd)
end = time.time()
ret_profiler = cf.load_config_file("tmpxxx_return_profiler.conf")
acc_time_profiler+=float(ret_profiler.get('DEFAULT', 'Tiempo'))
acc_time += end - start
cf.kill_all("killall -9 {}; killall -9 {}".format(
os.path.basename(gdb_exec), os.path.basename(benchmark_binary)))
return acc_time_profiler / cp.MAX_TIMES_TO_PROFILE, acc_time / cp.MAX_TIMES_TO_PROFILE
"""
Function to generate the gold execution
"""
def generate_gold(gdb_exec, benchmark_binary, benchmark_args,device):
# Create tmp path and clean it if it exists
tmp_path = os.path.dirname(os.path.realpath(__file__)) + "/" + cp.LOGS_PATH + "/tmp"
os.system("mkdir -p " + tmp_path)
os.system("rm -rf " + tmp_path + "/*")
script = 'env CUDA_VISIBLE_DEVICES={} {} -ex \'py arg0 = {}\' -n -batch -x {} > {} 2> {}'
init_string = '"file {}; set args {}"'.format(benchmark_binary, benchmark_args)
profiler_cmd = script.format(device, gdb_exec, init_string, cp.PROFILER_SCRIPT, cp.GOLD_OUTPUT_PATH, cp.GOLD_ERR_PATH)
if cp.DEBUG:
print("PROFILER CMD: {}".format(profiler_cmd))
# Execute and save gold file
return os.system(profiler_cmd)
def main():
os.system("rm -f {}".format(cp.KERNEL_INFO_DIR))
parser = argparse.ArgumentParser()
parser.add_argument('-c', '--conf', dest="config_file", help='Configuration file', required=True)
parser.add_argument('-d', '--device', dest="device", help="The GPU to perform FI."
" Default is 0.", required=False, default=0, type=int)
args = parser.parse_args()
# Read the configuration file with data for all the apps that will be executed
conf = cf.load_config_file(args.config_file)
# First set env vars
cf.set_python_env()
########################################################################
# Profiler step
# Max time will be obtained by running
# it will also get app output for golden copy
# that is,
print("###################################################\n1 - Profiling application")
if 'benchmarkBinary_noverificar' in conf['DEFAULT']:
benchmark_binary = conf.get('DEFAULT', 'benchmarkBinary_noverificar')
else:
benchmark_binary = conf.get('DEFAULT', 'benchmarkBinary')
if 'benchmarkArgs_noverificar' in conf['DEFAULT']:
benchmark_args = conf.get('DEFAULT', 'benchmarkArgs_noverificar')
else:
benchmark_args = conf.get('DEFAULT', 'benchmarkArgs')
section= 'kernel_end' in conf['DEFAULT']
kernel_end=''
if (section):
kernel_end=conf.get('DEFAULT','kernel_end')
gdb_exec = conf.get("DEFAULT", "gdbExecName")
kernel=conf.get('DEFAULT', 'kernel')
[max_time_kernel,max_time_app] = profiler_caller(gdb_exec=gdb_exec,kernel=kernel, benchmark_binary=benchmark_binary, benchmark_args=benchmark_args,device=args.device,section=section,kernel_end=kernel_end)
print ("Time kernel= "+str(max_time_kernel)+ "Time app "+str(max_time_app))
# saving gold
print ("Saving gold");
generate_gold_result = generate_gold(gdb_exec=gdb_exec,
benchmark_binary=benchmark_binary, benchmark_args=benchmark_args,device=args.device)
if generate_gold_result != 0:
raise EnvironmentError("Gold generation did not finish well, the fault injection will not work")
# Remove trash GDB info from the std output and the err output
cf.remove_useless_information_from_output(cp.GOLD_OUTPUT_PATH)
cf.remove_useless_information_from_output(cp.GOLD_ERR_PATH)
# Save the kernel configuration txt file
cf.save_file(file_path=cp.KERNEL_INFO_DIR, data={'max_time': max_time_app,'max_time_kernel': max_time_kernel})
print("1 - Profile finished\n###################################################")
if __name__ == '__main__':
main()
app_profiler_old.py 0 → 100755
View file @ e4ae2422
#!/usr/bin/env python3
import argparse
import os
import re
import time
import common_functions as cf
import common_parameters as cp
def generate_dict(sm_version, input_file_name):
with open(input_file_name, "r") as f:
# dictionary to store the number of allocated registers per static
kernel_reg = {}
kernel_name = "" # temporary variable to store the kernel_name
check_for_register_count = False
# process the input file created by capturing the stderr while compiling the
# application using -Xptxas -v options
for line in f: # for each line in the file
m = re.match(r".*Compiling entry function.*'(\S+)'.*for.*'{}'.*".format(sm_version), line)
if m:
kernel_name = m.group(1)
check_for_register_count = True
m = re.match(r".*Used[ ]+(\d+)[ ]+registers.*", line)
if check_for_register_count and m:
reg_num = m.group(1) # extract register number
if kernel_name not in kernel_reg:
# associate the extracted register number with the kernel name
kernel_reg[kernel_name] = int(reg_num.strip())
else:
print("Warning: {} exists in the kernel_reg dictionary. "
"Skipping this register count.".format(kernel_name))
check_for_register_count = False
return kernel_reg
"""
Function that calls the profiler based on the injection mode
"""
def profiler_caller(gdb_exec, kernel, benchmark_binary, benchmark_args):
acc_time = 0
script = '{} -ex \'py arg0 = {}\' -n -batch -x {}'
benchmark_args_striped = benchmark_args.replace('\\n', '').replace('\\', '')
print ("KERNEL"+kernel)
init_string = '"file {}; set args {}"'.format(benchmark_binary, benchmark_args_striped)
#init_string = '"file {}; set args {}; break {}"'.format(benchmark_binary, benchmark_args_striped,kernel)
profiler_cmd = script.format(gdb_exec, init_string, cp.PROFILER_SCRIPT)
print ("Profiler caller")
if cp.DEBUG:
print("PROFILER CMD: {}".format(profiler_cmd))
for i in range(0, cp.MAX_TIMES_TO_PROFILE):
start = time.time()
os.system(profiler_cmd)
end = time.time()
acc_time += end - start
cf.kill_all("killall -9 {}; killall -9 {}".format(
os.path.basename(gdb_exec), os.path.basename(benchmark_binary)))
return acc_time / cp.MAX_TIMES_TO_PROFILE
"""
Function to generate the gold execution
"""
def generate_gold(gdb_exec, benchmark_binary, benchmark_args):
# Create tmp path and clean it if it exists
tmp_path = os.path.dirname(os.path.realpath(__file__)) + "/" + cp.LOGS_PATH + "/tmp"
if not os.path.exists(tmp_path):
os.mkdir(tmp_path)
os.system("rm -rf " + tmp_path + "/*")
script = '{} -ex \'py arg0 = {}\' -n -batch -x {} > {} 2> {}'
init_string = '"file {}; set args {}"'.format(benchmark_binary, benchmark_args)
profiler_cmd = script.format(gdb_exec, init_string, cp.PROFILER_SCRIPT, cp.GOLD_OUTPUT_PATH, cp.GOLD_ERR_PATH)
if cp.DEBUG:
print("PROFILER CMD: {}".format(profiler_cmd))
# Execute and save gold file
return os.system(profiler_cmd)
def main():
os.system("rm -f {}".format(cp.KERNEL_INFO_DIR))
parser = argparse.ArgumentParser()
parser.add_argument('-c', '--conf', dest="config_file", help='Configuration file', required=True)
args = parser.parse_args()
# Read the configuration file with data for all the apps that will be executed
conf = cf.load_config_file(args.config_file)
# First set env vars
cf.set_python_env()
########################################################################
# Profiler step
# Max time will be obtained by running
# it will also get app output for golden copy
# that is,
print("###################################################\n1 - Profiling application")
if 'benchmarkBinary_noverificar' in conf['DEFAULT']:
benchmark_binary = conf.get('DEFAULT', 'benchmarkBinary_noverificar')
else:
benchmark_binary = conf.get('DEFAULT', 'benchmarkBinary')
if 'benchmarkArgs_noverificar' in conf['DEFAULT']:
benchmark_args = conf.get('DEFAULT', 'benchmarkArgs_noverificar')
else:
benchmark_args = conf.get('DEFAULT', 'benchmarkArgs')
gdb_exec = conf.get("DEFAULT", "gdbExecName")
kernel=conf.get('DEFAULT', 'kernel')
max_time_app = profiler_caller(gdb_exec=gdb_exec,kernel=kernel, benchmark_binary=benchmark_binary, benchmark_args=benchmark_args)
# saving gold
print ("Saving gold");
generate_gold_result = generate_gold(gdb_exec=gdb_exec,
benchmark_binary=benchmark_binary, benchmark_args=benchmark_args)
if generate_gold_result != 0:
raise EnvironmentError("Gold generation did not finish well, the fault injection will not work")
# Remove trash GDB info from the std output and the err output
cf.remove_useless_information_from_output(cp.GOLD_OUTPUT_PATH)
cf.remove_useless_information_from_output(cp.GOLD_ERR_PATH)
# Save the kernel configuration txt file
cf.save_file(file_path=cp.KERNEL_INFO_DIR, data={'max_time': max_time_app})
print("1 - Profile finished\n###################################################")
if __name__ == '__main__':
main()
codes/mmElem/matrixMul.cu 0 → 100644
View file @ e4ae2422
/**
* Copyright 1993-2015 NVIDIA Corporation. All rights reserved.
*
* Please refer to the NVIDIA end user license agreement (EULA) associated
* with this source code for terms and conditions that govern your use of
* this software. Any use, reproduction, disclosure, or distribution of
* this software and related documentation outside the terms of the EULA
* is strictly prohibited.
*
*/
/**
* Matrix multiplication: C = A * B.
* Host code.
*
* This sample implements matrix multiplication as described in Chapter 3
* of the programming guide.
* It has been written for clarity of exposition to illustrate various CUDA
* programming principles, not with the goal of providing the most
* performant generic kernel for matrix multiplication.
*
* See also:
* V. Volkov and J. Demmel, "Benchmarking GPUs to tune dense linear algebra,"
* in Proc. 2008 ACM/IEEE Conf. on Supercomputing (SC '08),
* Piscataway, NJ: IEEE Press, 2008, pp. Art. 31:1-11.
*/
// System includes
#include <stdio.h>
#include <assert.h>
// CUDA runtime
#include <cuda_runtime.h>
// Helper functions and utilities to work with CUDA
#include <helper_functions.h>
#include <helper_cuda.h>
#include <omp.h>
#if BUILD_TIMER == 1
static double timer;
#endif
/**
* Matrix multiplication (CUDA Kernel) on the device: C = A * B
* wA is A's width and wB is B's width
* Every thread computes one element of C as a dot product
* C[i][j] = A[i][:] * B[:][j]
*/
__global__ void matrixMulCUDA(float *C, float *A, float *B, int ldA, int ldB, int ldC) {
// Thread global indexes
int i = blockIdx.y * blockDim.y + threadIdx.y;
int j = blockIdx.x * blockDim.x + threadIdx.x;
// printf("**C[%d][%d]\n", i, j);
float *ptrA = &A[i*ldA]; // Pointer to the first element of row i of A
float tmp = 0.0f;
for (int k = 0; k < ldA; k++) {
tmp += (*ptrA++) * B[k*ldB+j];
}
C[i*ldC+j] = tmp;
// printf("C[%d][%d] = %f\n", i, j, tmp);
}
void constantInit(float *data, int size, float val) {
for (int i = 0; i < size; ++i) {
data[i] = val;
}
}
double mysecond() {
struct timeval tp;
struct timezone tzp;
int i = gettimeofday(&tp, &tzp);
return ((double) tp.tv_sec + (double) tp.tv_usec * 1.e-6);
}
/**
* Run a simple test of matrix multiplication using CUDA
*/
int matrixMultiply(int argc, char **argv, int block_size, dim3 &dimsA,
dim3 &dimsB) {
// Allocate host memory for matrices A and B
unsigned int size_A = dimsA.x * dimsA.y;
unsigned int mem_size_A = sizeof(float) * size_A;
float *h_A = (float *) malloc(mem_size_A);
unsigned int size_B = dimsB.x * dimsB.y;
unsigned int mem_size_B = sizeof(float) * size_B;
float *h_B = (float *) malloc(mem_size_B);
// Initialize host memory
const float valB = 0.01f;
constantInit(h_A, size_A, 1.0f);
constantInit(h_B, size_B, valB);
// Allocate device memory
float *d_A, *d_B, *d_C;
// Allocate host matrix C
dim3 dimsC(dimsB.x, dimsA.y, 1);
unsigned int mem_size_C = dimsC.x * dimsC.y * sizeof(float);
float *h_C = (float *) malloc(mem_size_C);
if (h_C == NULL) {
fprintf(stderr, "Failed to allocate host matrix C!\n");
exit (EXIT_FAILURE);
}
cudaError_t error;
error = cudaMalloc((void **) &d_A, mem_size_A);
if (error != cudaSuccess) {
printf("cudaMalloc d_A returned error %s (code %d), line(%d)\n",
cudaGetErrorString(error), error, __LINE__);
exit (EXIT_FAILURE);
}
error = cudaMalloc((void **) &d_B, mem_size_B);
if (error != cudaSuccess) {
printf("cudaMalloc d_B returned error %s (code %d), line(%d)\n",
cudaGetErrorString(error), error, __LINE__);
exit (EXIT_FAILURE);
}
error = cudaMalloc((void **) &d_C, mem_size_C);
if (error != cudaSuccess) {
printf("cudaMalloc d_C returned error %s (code %d), line(%d)\n",
cudaGetErrorString(error), error, __LINE__);
exit (EXIT_FAILURE);
}
// copy host memory to device
error = cudaMemcpy(d_A, h_A, mem_size_A, cudaMemcpyHostToDevice);
if (error != cudaSuccess) {
printf("cudaMemcpy (d_A,h_A) returned error %s (code %d), line(%d)\n",
cudaGetErrorString(error), error, __LINE__);
exit (EXIT_FAILURE);
}
error = cudaMemcpy(d_B, h_B, mem_size_B, cudaMemcpyHostToDevice);
if (error != cudaSuccess) {
printf("cudaMemcpy (d_B,h_B) returned error %s (code %d), line(%d)\n",
cudaGetErrorString(error), error, __LINE__);
exit (EXIT_FAILURE);
}
// Setup execution parameters
dim3 threads(block_size, block_size);
dim3 grid(dimsB.x / threads.x, dimsA.y / threads.y);
// Create and start timer
printf("Computing result using CUDA Kernel...\n");
// Performs warmup operation using matrixMul CUDA kernel
// if (block_size == 16) {
// matrixMulCUDA<16> <<<grid, threads>>>(d_C, d_A, d_B, dimsA.x, dimsB.x);
// } else {
// matrixMulCUDA<32> <<<grid, threads>>>(d_C, d_A, d_B, dimsA.x, dimsB.x);
// }
// printf("done\n");
//
// cudaDeviceSynchronize();
// Allocate CUDA events that we'll use for timing
cudaEvent_t start;
error = cudaEventCreate(&start);
if (error != cudaSuccess) {
fprintf(stderr, "Failed to create start event (error code %s)!\n",
cudaGetErrorString(error));
exit (EXIT_FAILURE);
}
cudaEvent_t stop;
error = cudaEventCreate(&stop);
if (error != cudaSuccess) {
fprintf(stderr, "Failed to create stop event (error code %s)!\n",
cudaGetErrorString(error));
exit (EXIT_FAILURE);
}
// Record the start event
error = cudaEventRecord(start, NULL);
if (error != cudaSuccess) {
fprintf(stderr, "Failed to record start event (error code %s)!\n",
cudaGetErrorString(error));
exit (EXIT_FAILURE);
}
// Execute the kernel
int nIter = 1;
#if BUILD_TIMER == 1
printf("BEFORE START KERNEL %lf\n", mysecond() - timer);
double t1 = mysecond();
#endif
for (int j = 0; j < nIter; j++) {
//matrixMulCUDA<32> <<<grid, threads>>>(d_C, d_A, d_B, dimsA.x, dimsB.x);
matrixMulCUDA <<<grid, threads>>>(d_C, d_A, d_B, dimsA.x, dimsB.x, dimsC.x);
cudaDeviceSynchronize();
}
#if BUILD_TIMER == 1
double exec_time = mysecond() - t1;
printf("KERNEL EXECUTION TIME %lf\n", exec_time);
#endif
// Record the stop event
error = cudaEventRecord(stop, NULL);
if (error != cudaSuccess) {
fprintf(stderr, "Failed to record stop event (error code %s)!\n",
cudaGetErrorString(error));
exit (EXIT_FAILURE);
}
// Wait for the stop event to complete
error = cudaEventSynchronize(stop);
if (error != cudaSuccess) {
fprintf(stderr,
"Failed to synchronize on the stop event (error code %s)!\n",
cudaGetErrorString(error));
exit (EXIT_FAILURE);
}
float msecTotal = 0.0f;
error = cudaEventElapsedTime(&msecTotal, start, stop);
if (error != cudaSuccess) {
fprintf(stderr,
"Failed to get time elapsed between events (error code %s)!\n",
cudaGetErrorString(error));
exit (EXIT_FAILURE);
}
#if BUILD_TIMER == 1
// Compute and print the performance
float msecPerMatrixMul = msecTotal / nIter;
double flopsPerMatrixMul = 2.0 * (double) dimsA.x * (double) dimsA.y
* (double) dimsB.x;
double gigaFlops = (flopsPerMatrixMul * 1.0e-9f)
/ (msecPerMatrixMul / 1000.0f);
printf(
"Performance= %.2f GFlop/s, Time= %.3f msec, Size= %.0f Ops, WorkgroupSize= %u threads/block\n",
gigaFlops, msecPerMatrixMul, flopsPerMatrixMul,
threads.x * threads.y);
#endif
// Copy result from device to host
error = cudaMemcpy(h_C, d_C, mem_size_C, cudaMemcpyDeviceToHost);
if (error != cudaSuccess) {
printf("cudaMemcpy (h_C,d_C) returned error %s (code %d), line(%d)\n",
cudaGetErrorString(error), error, __LINE__);
exit (EXIT_FAILURE);
}
printf("Checking computed result for correctness: ");
bool correct = true;
// test relative error by the formula
// |<x, y>_cpu - <x,y>_gpu|/<|x|, |y|> < eps
double eps = 1.e-6; // machine zero
#if BUILD_TIMER == 1
t1 = mysecond();
#endif
#pragma omp parallel for shared(h_C, correct)
for (int i = 0; i < (int) (dimsC.x * dimsC.y); i++) {
float abs_err = fabs(h_C[i] - float(dimsA.x * valB));
float dot_length = dimsA.x;
float abs_val = fabs(h_C[i]);
float rel_err = abs_err / abs_val / dot_length;
if (rel_err > eps) {
printf("Error! Matrix[%05d]=%.8f, ref=%.8f error term is > %E\n", i,
h_C[i], dimsA.x * valB, eps);
#pragma omp critical
{
correct = false;
}
}
}
#if BUILD_TIMER == 1
exec_time = mysecond() - t1;
printf("CMP TIME %lf\n", exec_time);
#endif
printf("%s\n", correct ? "Result = PASS" : "Result = FAIL");
// Clean up memory
free(h_A);
free(h_B);
free(h_C);
cudaFree(d_A);
cudaFree(d_B);
cudaFree(d_C);
printf(
"\nNOTE: The CUDA Samples are not meant for performance measurements. "
"Results may vary when GPU Boost is enabled.\n");
if (correct) {
return EXIT_SUCCESS;
} else {
return EXIT_FAILURE;
}
}
/**
* Program main
*/
int main(int argc, char **argv) {
#if BUILD_TIMER == 1
timer = mysecond();
#endif
printf("[Matrix Multiply Using CUDA] - Starting...\n");
if (checkCmdLineFlag(argc, (const char **) argv, "help")
|| checkCmdLineFlag(argc, (const char **) argv, "?")) {
printf("Usage -device=n (n >= 0 for deviceID)\n");
printf(" -wA=WidthA -hA=HeightA (Width x Height of Matrix A)\n");
printf(" -wB=WidthB -hB=HeightB (Width x Height of Matrix B)\n");
printf(
" Note: Outer matrix dimensions of A & B matrices must be equal.\n");
exit (EXIT_SUCCESS);
}
// By default, we use device 0, otherwise we override the device ID based on what is provided at the command line
int devID = 0;
if (checkCmdLineFlag(argc, (const char **) argv, "device")) {
devID = getCmdLineArgumentInt(argc, (const char **) argv, "device");
cudaSetDevice(devID);
}
cudaError_t error;
cudaDeviceProp deviceProp;
error = cudaGetDevice(&devID);
if (error != cudaSuccess) {
printf("cudaGetDevice returned error %s (code %d), line(%d)\n",
cudaGetErrorString(error), error, __LINE__);
}
error = cudaGetDeviceProperties(&deviceProp, devID);
if (deviceProp.computeMode == cudaComputeModeProhibited) {
fprintf(stderr,
"Error: device is running in <Compute Mode Prohibited>, no threads can use ::cudaSetDevice().\n");
exit (EXIT_SUCCESS);
}
if (error != cudaSuccess) {
printf(
"cudaGetDeviceProperties returned error %s (code %d), line(%d)\n",
cudaGetErrorString(error), error, __LINE__);
} else {
printf("GPU Device %d: \"%s\" with compute capability %d.%d\n\n", devID,
deviceProp.name, deviceProp.major, deviceProp.minor);
}
// Use a larger block size for Fermi and above
int block_size = (deviceProp.major < 2) ? 16 : 32;
dim3 dimsA(5 * 2 * block_size, 5 * 2 * block_size, 1);
dim3 dimsB(5 * 4 * block_size, 5 * 2 * block_size, 1);
// width of Matrix A
if (checkCmdLineFlag(argc, (const char **) argv, "wA")) {
dimsA.x = getCmdLineArgumentInt(argc, (const char **) argv, "wA");
}
// height of Matrix A
if (checkCmdLineFlag(argc, (const char **) argv, "hA")) {
dimsA.y = getCmdLineArgumentInt(argc, (const char **) argv, "hA");
}
// width of Matrix B
if (checkCmdLineFlag(argc, (const char **) argv, "wB")) {
dimsB.x = getCmdLineArgumentInt(argc, (const char **) argv, "wB");
}
// height of Matrix B
if (checkCmdLineFlag(argc, (const char **) argv, "hB")) {
dimsB.y = getCmdLineArgumentInt(argc, (const char **) argv, "hB");
}
if (dimsA.x != dimsB.y) {
printf("Error: outer matrix dimensions must be equal. (%d != %d)\n",
dimsA.x, dimsB.y);
exit (EXIT_FAILURE);
}
printf("MatrixA(%d,%d), MatrixB(%d,%d)\n", dimsA.x, dimsA.y, dimsB.x,
dimsB.y);
int matrix_result = matrixMultiply(argc, argv, block_size, dimsA, dimsB);
exit(matrix_result);
}
codes/mmElem/matrixMul.cudafe1.c 0 → 100644
View file @ e4ae2422
# 1 "matrixMul.cu"
# 139 "/usr/include/stdio.h" 3
extern FILE *stderr;
# 74 "/usr/include/c++/9/iostream" 3
static struct _ZNSt8ios_base4InitE _ZN39_INTERNAL_17_matrixMul_cpp1_ii_9deaad98St8__ioinitE __attribute__((visibility("default"))) = {};
extern void *__dso_handle __attribute__((visibility("hidden")));
codes/mmElem/matrixMul.cudafe1.stub.c 0 → 100644
View file @ e4ae2422
#pragma GCC diagnostic push
#pragma GCC diagnostic ignored "-Wunused-function"
#pragma GCC diagnostic ignored "-Wcast-qual"
#define __NV_CUBIN_HANDLE_STORAGE__ static
#if !defined(__CUDA_INCLUDE_COMPILER_INTERNAL_HEADERS__)
#define __CUDA_INCLUDE_COMPILER_INTERNAL_HEADERS__
#endif
#include "crt/host_runtime.h"
#include "matrixMul.fatbin.c"
extern void __device_stub__Z13matrixMulCUDAPfS_S_iii(float *, float *, float *, int, int, int);
static void __nv_cudaEntityRegisterCallback(void **);
static void __sti____cudaRegisterAll(void) __attribute__((__constructor__));
void __device_stub__Z13matrixMulCUDAPfS_S_iii(float *__par0, float *__par1, float *__par2, int __par3, int __par4, int __par5){__cudaLaunchPrologue(6);__cudaSetupArgSimple(__par0, 0UL);__cudaSetupArgSimple(__par1, 8UL);__cudaSetupArgSimple(__par2, 16UL);__cudaSetupArgSimple(__par3, 24UL);__cudaSetupArgSimple(__par4, 28UL);__cudaSetupArgSimple(__par5, 32UL);__cudaLaunch(((char *)((void ( *)(float *, float *, float *, int, int, int))matrixMulCUDA)));}
# 50 "matrixMul.cu"
void matrixMulCUDA( float *__cuda_0,float *__cuda_1,float *__cuda_2,int __cuda_3,int __cuda_4,int __cuda_5)
# 50 "matrixMul.cu"
{__device_stub__Z13matrixMulCUDAPfS_S_iii( __cuda_0,__cuda_1,__cuda_2,__cuda_3,__cuda_4,__cuda_5);
# 66 "matrixMul.cu"
}
# 1 "matrixMul.cudafe1.stub.c"
static void __nv_cudaEntityRegisterCallback( void **__T3) { __nv_dummy_param_ref(__T3); __nv_save_fatbinhandle_for_managed_rt(__T3); __cudaRegisterEntry(__T3, ((void ( *)(float *, float *, float *, int, int, int))matrixMulCUDA), _Z13matrixMulCUDAPfS_S_iii, (-1)); }
static void __sti____cudaRegisterAll(void) { __cudaRegisterBinary(__nv_cudaEntityRegisterCallback); }
#pragma GCC diagnostic pop
codes/mmElem/matrixMul.fatbin.c 0 → 100644
View file @ e4ae2422
#ifndef __SKIP_INTERNAL_FATBINARY_HEADERS
#include "fatbinary_section.h"
#endif
#define __CUDAFATBINSECTION ".nvFatBinSegment"
#define __CUDAFATBINDATASECTION ".nv_fatbin"
asm(
".section .nv_fatbin, \"a\"\n"
".align 8\n"
"fatbinData:\n"
".quad 0x00100001ba55ed50,0x0000000000006da8,0x0000008801010001,0x0000000000002f10\n"
".quad 0x0000004000002f0f,0x0000004600070000,0x0000000c00000078,0x0000000000002013\n"
".quad 0x0000000000000000,0x00000000000136b7,0x0000002d00000048,0x6f642d2d20672d20\n"
".quad 0x656772656d2d746e,0x6c6263697361622d,0x722d2d20736b636f,0x74612d6e72757465\n"
".quad 0x00000020646e652d,0x754d78697274616d,0x0000000075632e6c,0x762e23f000010a13\n"
".quad 0x37206e6f69737265,0x677261742e0a302e,0x30375f6d73207465,0x0a6775626564202c\n"
".quad 0x737365726464612e,0x343620657a69735f,0x636e756603f00035,0x6d617261702e2820\n"
".quad 0x0012203233622e20,0x61767465725f03f9,0x736261662029306c,0x140000220a280a66\n"
".quad 0x305fef00115f1100,0x772e0a0a3b0a290a,0xa70e004e206b6165,0x6c6c614d61647563\n"
".quad 0x183436270052636f,0x00202c1f00570400,0x7b0a290a310df10b,0x206c61636f6c2e0a\n"
".quad 0x38206e67696c612e,0x155f5f2038622e20,0x6f7065645f01f200,0x2e0a3b5d385b3074\n"
".quad 0x53253b0048676572,0x8900104c15000f50,0x3e333c7225203233,0x3e353c6472810021\n"
".quad 0x322081006b0a0a3b,0x00d40a3120353720,0x3a306e6967656271,0x20302032096a0019\n"
".quad 0x766f6d0a63002530,0x00972c1b0065752e,0x2e617476630a3b71,0x8f2c130025040013\n"
".quad 0x180200fa646c2200,0x5b202c3164726e00,0x321f00295d1e0120,0x8a3b5d3136010029\n"
".quad 0x00062c3364724000,0x001374730a3b3151,0x5d302b5053255b71,0xf803002a33120016\n"
".quad 0x3b32220014341100,0x00f36578653d00ee,0x6d740a33203737a2,0x01410100623a3070\n"
".quad 0x4a393939202c3163,0x0049321000120100,0x625800af02007202,0x810302185b093233\n"
".quad 0x7465720a3b327000,0x00773a312300533b,0x0a7d0a3a30646e7f,0x6e754602fe1a0261\n"
".quad 0x7274744174654763,0x026c736574756269,0x2b0f02770e00230d,0x311f1b02820f0600\n"
".quad 0x02823038293d0282,0x30382f020282311f,0x210282311f0c0282,0x0e0f028d0f01360e\n"
".quad 0x751f2502980f013f,0x38220109060b0298,0x3f0298321f029832,0x0298311f02983316\n"
".quad 0x656369766544681e,0x00240e02990e010d,0x2c32332f00029a0f,0x1218002c311f0b00\n"
".quad 0x371d02810802c732,0x030a028133190281,0x0301030281321f05,0x00e60f0102490f05\n"
".quad 0x0f01b904024a090b,0x1f0034311e0b0034,0x027e321a0d003432,0x01021e06027e3214\n"
".quad 0x0279010012090291,0x381301000602780e,0x351d027834180510,0x351f004836100278\n"
".quad 0x1602783616100278,0x1e0278321f027835,0x6f0e010974654732,0x1f07ce07001b0502\n"
".quad 0x1a020e090d048f7b,0x02020e331f048f39,0x850d09020e30392f,0x01790b04019d0f00\n"
".quad 0x1f03f13913006b06,0x017937163f03f136,0x4f1afe1e0179331f,0x79636e6170756363\n"
".quad 0x766974634178614d,0x50736b636f6c4265,0x7069746c754d7265,0x726f737365636f72\n"
".quad 0xb90e16003b0f0199,0x04360e2500430f06,0x0086321f1e00430f,0x82351d0282331e2f\n"
".quad 0x04900a0282371902,0x04900f030282341f,0x3b5d2d22016e0f06,0x0f24004c321f004c\n"
".quad 0x1f2500970f0004bf,0x004c331f00009732,0x190007765d332f24,0x0e078a3514001431\n"
".quad 0x01ad020012010363,0x0603bf0b07c43614,0x38180538391301a8,0x830103bf331d03bf\n"
".quad 0x34161003bf331f00,0x341f03bf391603bf,0x687469578e4703bf,0x440f0a6967616c46\n"
".quad 0x004c0f03d10e2000,0x004c0c3003da0f27,0x98331f380098321f,0x091108c8341f3800\n"
".quad 0x1f06bb30312a0438,0x0f00260202043935,0x430f00b00c2f043a,0x044c0f00550c3204\n"
".quad 0x3204550f00540c31,0x0f00a9331e00550c,0x2204b2341f2e00fd,0xfc0f04b20b09e809\n"
".quad 0x3031230233060e09,0x3f09fe30312f0884,0x04c735130f0f3117,0x6c62697369761af6\n"
".quad 0x7972746e652e2065,0x74616d33315a5f20,0x55436c754d786972,0x5f535f5366504144\n"
".quad 0x0105000494696969,0x0f04780e0700280f,0x1c0030311f120030,0x300f012300044007\n"
".quad 0x1f1c0030331f0e00,0x043835171c003034,0x7025206465727077,0x663601b466100af3\n"
".quad 0x3d01c501044a393c,0x363126045c34323c,0x045c3520313a045d,0x045c311702293613\n"
".quad 0x344f045b0d002503,0x3f0f0e00b75b202c,0x1f110039351f0104,0x0039361f00003931\n"
".quad 0x0f02cc0404080a11,0x1f0003eb0f0f0038,0x0038341e11003837,0x5d352d110038381f\n"
".quad 0xaa3510017e0603a7,0xc7101d0903530103,0x6961746325202c39,0x317300173b792e64\n"
".quad 0x0017746e25202c30,0x732e6f6c2e6c7572,0x3900372c3123001a,0x0902000033303172\n"
".quad 0x2f64646134003203,0x32312200342c1500,0x33100094331603d3,0x8e33150043070094\n"
".quad 0x8e3414005b781900,0x2c3524008e781c00,0x0034343172390038,0x008f7818008f3613\n"
".quad 0x8f361200352c3226,0x20374b008f341600,0xc52c372400603631,0x732e401238361100\n"
".quad 0x72250942001c3436,0x730a3b3764001c64,0x202c3832030d6c68,0x350500773213001a\n"
".quad 0x0078381205850507,0x6303680201873512,0x0001306630202c34,0x3120394700913b15\n"
".quad 0x6622004600002634,0x365801da06003f34,0x384e00fa35312030,0x00770305e930202c\n"
".quad 0x1400980301fe3713,0x3818001c00062c35,0x19661300ca000093,0x322300603a392700\n"
".quad 0x30981269000a8f33,0x315f3642420a0a3a,0x4002002e3417008f,0x8500005a3b332706\n"
".quad 0x3212069f38662901,0x31a2005535120087,0x6c2e707465730a3a,0x4a2c317034016b74\n"
".quad 0x3a32327401b50200,0x3e0004b7746f6e0a,0x25400a3b3170d000,0x9220617262203270\n"
".quad 0x6172620a3b34a100,0x3222000f696e752e,0x3100a9321600a93b,0x3a3328004d362031\n"
".quad 0x0200912c332301ab,0x038d3a3432300169,0x1d0004680000ba02,0xa1321602875d1b00\n"
".quad 0x2c3225026f371900,0x31120240321d001e,0xcc0202413219001e,0x2602430b001c0102\n"
".quad 0x3b33313607a32c34,0x3123009636120096,0x3722001602009734,0x483666253200b12c\n"
".quad 0x822c3324002d0200,0x3a353700e0371201,0x00bf39322701bb0a,0x2f02060200d63513\n"
".quad 0x02080301023b3632,0x1e001903023b371d,0x392901e603023b38,0x331400650201a13a\n"
".quad 0x33364a0397060251,0x96393125016f3420,0x323500b03b382703,0x02016f0e001e2c30\n"
".quad 0x3201016e301a0017,0x27016d0b001b0000,0x08b53011094a2c31,0x312101680100dc00\n"
".quad 0x364503a70501445d,0xb2331708b2312036,0x00190008b2361c08,0x64344e5a5f2000fa\n"
".quad 0x6a6a453143336d69,0x110f001c0a0d396a,0x03002432332f000d,0x24321f100024311f\n"
".quad 0xf10e0010850f1000,0x20335907f0321607,0xb9371307f1313234,0x00260407f1331701\n"
".quad 0x0200880f020c4d0f,0x2c0f00890107e60a,0x0e000b53311f0600,0x331f0607cc0f002c\n"
".quad 0x060750331e05002c,0x31332100d10100dd,0x1402203a322001f3,0x090b595d13022075\n"
".quad 0x2b2100233819002a,0x2a0025321a002534,0x6a5d382300253534,0x6c32352600250805\n"
".quad 0x026c371f17950702,0x321f08026c321f0b,0x10026c321f10026c,0x6c321f10026c321f\n"
".quad 0x1f26026c381f3a02,0x026c321f18026c32,0x321f18026c321f18,0xf80400dd060b026c\n"
".quad 0x0a3a345801f30001,0x6d6574430136090a,0x3b67657278004570,0x000b0101e5090a0a\n"
".quad 0x093436420d753018,0xfe000d6f0300165b,0x3332332200330507,0x03001602162d0d00\n"
".quad 0x321f0100320f0033,0x1f00323213020032,0x0032331f00003232,0x4102f7302b333402\n"
".quad 0x0a1a06206c6c6163,0x0a280a202c5103a0,0x0009202c3032002d,0x0009321400093114\n"
".quad 0x0945023a290a3330,0x184d070330090a7d,0x0b13180f0330381e,0x20003474535a5f51\n"
".quad 0x00180402db0a282a,0x070d467b1a201807,0xc5343220344a1838,0xc5341701e8391302\n"
".quad 0x0502c50600260402,0x3c0e00715b1e0a02,0x3234325200540602,0x023b361f023b3120\n"
".quad 0x663a023b32332e18,0x0208661b023b3233,0x281401b907012c03,0x00b30a202c31013f\n"
".quad 0xdf0801a00001bb07,0x0c5b160898091100,0x0e00810a01bf0410,0x1710313266261031\n"
".quad 0x0f01e3391118b733,0xcf0f21e70e1321e8,0x960c00019b0f0e01,0x13001e7362613301\n"
".quad 0x011600cc0f012f32,0x09656c691cf500ba,0x656d6f682f222031,0x6d2f61696461622f\n"
".quad 0x662d6c6f72616379,0x2f7365646f632f69,0x1a2f6d656c456d6d,0x202c2275632ee00e\n"
".quad 0x3033393033303631,0x323835313153090d,0x752f22203281004e,0x762f101efe2f7273\n"
".quad 0x2f302e31312dd212,0x3c2f2e2e2f6e6962,0x3638782f7307f123,0x756e696c2d34365f\n"
".quad 0x64756c636e692f78,0x1ab0645f22002e65,0x656d69746e7572d0,0x0070682e6970615f\n"
".quad 0x3930333634393570,0x3037393431532373,0xce280070331f0070,0x745f726f74636576\n"
".quad 0x3343006573657079,0x6534130065343731,0x2f2b2b9200a90500,0x003674616d632f39\n"
".quad 0x39333030373037e4,0x3639303934202c35,0x645f28009b351f00,0xc006009b65766972\n"
".quad 0x2e0a0a3731303530,0x2e11249774636573,0x6f666e695f812485,0x39385002620a7b20\n"
".quad 0x06321121330a3833,0x0029030016301200,0x1976657262626161,0x070a312200250200\n"
".quad 0x333023000f301300,0x3123000831120008,0x1200183219001030,0x2200163312002f35\n"
".quad 0x2800150200073936,0x00513512001c3137,0x0a39250007363422,0x0054342038330096\n"
".quad 0x000f373922003f02,0x001f040008363123,0x0035321200083512,0x00ba371900273714\n"
".quad 0x001e393925006103,0x010634362b006905,0x14001e656e696c41,0x3129006b03004434\n"
".quad 0x270300ee04009231,0x302900a838392900,0x005337392a00a030,0x0098313228004303\n"
".quad 0x030071311c003c04,0x3519016e351b00c5,0x30190035391b0080,0x06002e31312900a6\n"
".quad 0x6b0401a5391800c5,0x0193301900360500,0x006b0801470a322c,0x00cd040183373928\n"
".quad 0x00a003000838302a,0x3f00520f00593919,0x01db353722012402,0x01f102002c373123\n"
".quad 0x4500113017070506,0x3842002630646e65,0x0a33230235353120,0x34070060321f0067\n"
".quad 0x00340300c2383838,0x3802240109005701,0x0701004535332300,0x0062351f00620b00\n"
".quad 0x2300333630353708,0x002d381302243431,0x2700bd0300083014,0x0357321900453137\n"
".quad 0x7b3539313f005102,0x0b029339362b1001,0x004c03001003029b,0x352f006b07039e00\n"
".quad 0x0600be093300620a,0x361f0174311300d0,0x001e333829110067,0x03250b0034393929\n"
".quad 0x310f01ad05018a0b,0x28054233372a3701,0x1902e30600453831,0x002636382f006435\n"
".quad 0x00b63118008a0502,0x1e371d4000b6311f,0x04d7040100270f04,0x2f381900ce35362b\n"
".quad 0x2b0205da39392f00,0x00d6301800553530,0x35302f4f018c321f,0x00ce37392b020091\n"
".quad 0x00c732322f002703,0xee331f1f02ef0f15,0x09004537362f3f00,0x382b00c702003606\n"
".quad 0x450a01c30b011c35,0x332c002f30312c00,0x293b00ce0f047b0a,0x0b001f05000f3038\n"
".quad 0x30311901b5060724,0x351800b838362b00,0x03003d38392f0084,0x00d60f04e630302e\n"
".quad 0x361f005638372b79,0x1f0401060f3802e8,0x29047e0c7e010636,0x009b391a00273431\n"
".quad 0x011533382904b404,0x36312a003d37392a,0x05371f01050a006c,0x08321a009406c301\n"
".quad 0x1f01010631302f00,0x4d37362b75072f38,0x331a04013c301f00,0x05005e34312900a3\n"
".quad 0x1f090699351f01f4,0xaf30382b7500f639,0x34302f0092361900,0x7600dd321f20084b\n"
".quad 0x0408820402a93319,0x08003e31312505aa,0x2b7500bf331f00bf,0x2a01a40a004d3237\n"
".quad 0x6535302b001f3038,0x082d0a0378361a00,0x362b7500e736312f,0x392900270a00a838\n"
".quad 0x1f2c00f60f001f39,0xa834382a7500f637,0x37312a00c0321b00,0xc8381f03640b0075\n"
".quad 0x36362a000800a800,0x0f0056303129004e,0x063639312f0c0a1c,0x312a002e3430297d\n"
".quad 0x0046313029000830,0x1b00c8351a02f30d,0x0f0154321b010534,0x011530322f0202dc\n"
".quad 0x39392a040dfa0f75,0x45361b05970300cf,0x3605003e34312900,0x59311f0027361b00\n"
".quad 0x3f010531322f0107,0x0400080500de361a,0x080205a0301c002f,0x600f084139372900\n"
".quad 0x2b00e03531290508,0x00df3619001f3739,0x2403c20f40075c0f,0x002f04000f303729\n"
".quad 0x3719031103009304,0x38302a01bd050092,0x33322f0a140a003d,0x19004e381b7700e6\n"
".quad 0x003e31312b04bf36,0xd0341f0200d0301f,0x30312f08dc0b4000,0x301f003e0e01017f\n"
".quad 0x049135322f571090,0x312a03027f301f76,0x1705015c0b021536,0x9d351f002f361b00\n"
".quad 0x7600f736322f0906,0x05900402004d381f,0x1433371f2c00e70f,0x31302903a5351a48\n"
".quad 0x373819073f0d0055,0x362f001f39392900,0x0d00b10c02004d39,0xcb0800fe311500f6\n"
".quad 0x17381b4109790f03,0x00c7381a008a0200,0xc503020b04001704,0x2f38302a074e0501\n"
".quad 0x473019003f301c00,0x2f04ab0a00180500,0x009a0b4801d63133,0x6d0f04cb0400270b\n"
".quad 0x31312905d0080f0e,0x5b04003d371a0008,0x1f010b4633302f00,0x0055351b510eaa33\n"
".quad 0x1902b40b0008321b,0x362a01b60c00e635,0x030235302b007d38,0x0589331f0205890f\n"
".quad 0x3129007b33382a76,0x010d37392a010e34,0x054004710f01ae0a,0x1804034d371f0150\n"
".quad 0x034337392f006c35,0x3c34312a02b10b02,0x3837290a161f0f00,0x2f0200fa301f003c\n"
".quad 0x0f34198500fa3334,0x2a00270405630500,0x1a00f50f00843731,0x02c40f8500f5341f\n"
".quad 0xf3351f1700f30f1d,0x09042a31302b4700,0x04af02001f0600eb,0x03e537392f19ca0b\n"
".quad 0x02ea0f013f351903,0x048c361f00e40807,0x3029002639392a4e,0xd706005c0301ee39\n"
".quad 0x362a0903a1351f01,0x014835302b007b38,0x1f060de60f008b0a,0x03011d0806032a36\n"
".quad 0x311b4909540f22d4,0x0f09e635302d0008,0x370502e6030900c9,0x76373129075a0400\n"
".quad 0x09560f0401650f00,0x3019410ed7351f04,0x31341f00820e000f,0xce0f00ef37190208\n"
".quad 0x09007337312b1b0e,0x0f4005c6351f02fa,0x2f002d37191b00be,0xba321b0304213032\n"
".quad 0x5b321b0091341901,0x5e301f0008311b00,0x000e33393133090a,0x01c1301f6c010a0f\n"
".quad 0x6a39392b05e20a03,0x07050f00dc321d00,0x00f3341f00f30904,0x1d34150073331bc0\n"
".quad 0x00fb351f00fb0904,0x02d30f0610fe0f45,0x381b7201e0361f60,0x05007a30312900a6\n"
".quad 0x650a0d05e70f01f8,0x34012b0f01330400,0x690c8e0f012b371e,0xe60b00270a02280d\n"
".quad 0x2b02012330372f00,0x00a2351b025d3939,0x0104361f15c1301a,0x00dd0f0017381946\n"
".quad 0x01a3381f009f0917,0x0479391f00c609a5,0x312a00be3631294d,0x047a373129020534\n"
".quad 0x00e405002e303729,0x00ce341f002e381b,0x257b019435352f02,0x05004d0800a73537\n"
".quad 0x3830290018040046,0x1f005d39302a0065,0x2890301f0200ed33,0x26cc0f17024b0f46\n"
".quad 0x1036190500100f02,0x2b89052325100f00,0x0291391a4629860f,0x08321b005b36362b\n"
".quad 0x0d23391a00450b00,0x1f00a80500263619,0x0501020f1102ed36,0x01ca0d450102351f\n"
".quad 0x038701120f02880a,0x3129471cd10f0054,0x0955373929000834,0x080249351f041d05\n"
".quad 0x1f0500c50f000805,0x2738302a4d00c537,0x005c391a00890300,0x032300c60f00d505\n"
".quad 0x5b0f4e07b50f0064,0x007b0c027f073e05,0xd139302f23075b0f,0x1b008235362a7b01\n"
".quad 0x02004d351f008a33,0x01e130312f01e10c,0x2202880f04130554,0xc2088600a731312f\n"
".quad 0x3419000f35312a01,0x0901c13739250008,0x059900dc321f00dc,0x30190268311b125e\n"
".quad 0x02199130312f0064,0x284500ed3331323f,0x0337362f0dd83936,0x02003b39392f0107\n"
".quad 0x002731302a000805,0xa738392a00883618,0xc834312f04200b03,0x0c005535312a5507\n"
".quad 0x01930b00de0d01c1,0x01c2050f66303028,0x23810b001039302a,0x2e0f730aa035312f\n"
".quad 0x53301a07d4052305,0x00ec09003d351501,0x7130154601cc361f,0x19001f0300080900\n"
".quad 0x002539392b00f336,0x094002180f00b60a,0x0b450105371f0105,0x00450301f104018d\n"
".quad 0x006404020030312a,0xaf321b00b730382a,0x5b037b00af381f00,0x29002e0400450438\n"
".quad 0x2900360412be3231,0x8a37362b01c43939,0x32160301b3301f00,0x53391f010c090018\n"
".quad 0x040300fd381f4e06,0x0102a337302f00a9,0x080502005539392f,0x361f002731302a00\n"
".quad 0x00f430322f180580,0xa4361f00e4341b46,0x1d01e239302b0201,0x1b11079d0f00b035\n"
".quad 0x0a0159371b001730,0x7c03a531322f02f6,0x9a34190400bf331f,0x3237313f01bc0a00\n"
".quad 0x05018d301f4638b6,0x012b31133701850f,0x343028471376321f,0x00c604001e04003d\n"
".quad 0x291a0f9c0f005c04,0x3d30312a07213338,0x290f013131312900,0x0ca00129341f3401\n"
".quad 0x01010335302f0309,0x0a0d0a026937392f,0xdc0345010a351f01,0x29001f0a04870600\n"
".quad 0x1900ce0600be3739,0x312b021e0c0f4330,0xd709023606003735,0x048f0f00d7361e00\n"
".quad 0x005d351b01050c6a,0x2f39392900653719,0x3129006432372900,0xcd0801e308005c30\n"
".quad 0x0104331f40740304,0xbc08008138302d47,0xcb38302900b00401,0x312f001e36312901\n"
".quad 0x00be351f0300be36,0x8c321f6004710f45,0x710b0300a00f4801,0x00fc000063391a02\n"
".quad 0x0f413df20f0c680e,0x00af35362a23022b,0x03470d006530302a,0x3deb381e359c351a\n"
".quad 0x2639372b601e120f,0x32302f0017361900,0x391f00fb0d30041b,0x050734382b7200fb\n"
".quad 0x08011b311d00b60c,0x7301c130392f00c6,0x31312f02001e331f,0x0401e7391a020b69\n"
".quad 0x2f0302fa391f00fc,0x009a0a060d423438,0x3900008c03014a0b,0x0132311e01320903\n"
".quad 0x0c35362a7f102a0f,0x391e002f34312901,0x00bf0900a90904a6,0x2f08e10c0401580f\n"
".quad 0x05c60e9101193239,0x00fa331e3000fa0f,0x00ae301a3b07420f,0x0401010a00ae371a\n"
".quad 0x45381902c9060025,0x2f02007330382f00,0x24040e0601383939,0x2011020f4100ec0f\n"
".quad 0x03145b351f043b09,0x2a007d0c0046351a,0x0f23c80e00303132,0x1b00a1351b4500ef\n"
".quad 0x0900b9341500b935,0x1c450091381f0091,0x02183c311f001738,0x000c039b0f04320a\n"
".quad 0xbe391f00be0900d6,0x1f001732372b4500,0x1f00a80601190235,0x0400d70402004639\n"
".quad 0x03001f38302a019d,0x0d0500aa391a003f,0x0609730308610b01,0x00280500940a0389\n"
".quad 0x301f1d930d005e0b,0x0f130b210f930135,0x4c0116311f360116,0x004c040016373929\n"
".quad 0x1903001f3931313f,0x002636312805b231,0x160089301f009804,0x4600eb321f00eb09\n"
".quad 0x01ac381f06e1331b,0x05001f0400550302,0x050301c2301f0573,0x361b00aa371a006d\n"
".quad 0x01d90b008c0401f9,0x7a371f4600ee331f,0x001f0500de0b0202,0x011530312b02a904\n"
".quad 0x0008301a0424301a,0x6e3116002f34312a,0x01d5341f00e70900,0x00cf0f0e11390f4d\n"
".quad 0x2f0b9b0902f0092d,0x24fa0401106b3939,0x2e06be037600f40f,0x700d140f248e3939\n"
".quad 0x381b00bb06006209,0x1f03016b341f003d,0x020231195609a837,0x37392f002f30312a\n"
".quad 0x0f002f38190403fe,0x0b0508ab0f1a011b,0x220c03c8341a0044,0x381b4601e7381f01\n"
".quad 0x2a02092c361f0017,0x2a35302a07573431,0x010a003036312a01,0x0f00b93036313807\n"
".quad 0x2b005c331a5100c0,0x0f00460a00083231,0x4500c0311f0e00c0,0x04b3060898333829\n"
".quad 0x1608b80f08d0361e,0xd90b009204008303,0x7400ce3236313f13,0x03700f004e38362a\n"
".quad 0x004e040c07140f04,0x038836312c00f604,0x0f00fd331e00fd09,0x740400970343046d\n"
".quad 0x0501ca0f009f0b02,0x0832312b008a331a,0x5a361b0084311b00,0x0c0f760c00ba0201\n"
".quad 0x2b05014a38190e6a,0x0228341f02280b01,0x3029002736312a4d,0x0159393028012b31\n"
".quad 0x003d32312b001604,0x4e22810f0045371e,0x011d371f5cfe3113,0x37392a2308840fa3\n"
".quad 0x03011b31302f005d,0x0f3719aa011b331f,0x33313e0c079b0f00,0x2aa703250f00ed34\n"
".quad 0x08d3361f00463930,0x1f01f80d003d0401,0x008b361aba010b35,0x0401050f0020301c\n"
".quad 0xfd351bb20418361f,0x1436312b04200a00,0x010c09011c311501,0x004503b1010c371f\n"
".quad 0x392b051d0b003705,0x1f0104361a010c39,0x2300d60f45010438,0x080b39362b006c04\n"
".quad 0x1f00c70c02df311a,0x8337382fa901cb39,0x382a05dc30190400,0x0a05008b0f04e834\n"
".quad 0x46011b30342f02e6,0x90393029025b341a,0x01063715002f0600,0x450099311f009909\n"
".quad 0x0639362a132a3f0f,0x0570391a01470419,0x01d205010089321f,0x3f2a10620f003503\n"
".quad 0x351f4601a3323431,0x3129001005022a5b,0x322f05220b158a31,0x015f361c4601a333\n"
".quad 0x378707240148351f,0x245c35312f003604,0x0f00ea0a12950409,0x001e30372a1100db\n"
".quad 0x3137290035303128,0x00e9050222040094,0x00e334190008361b,0x4e361b005c38392b\n"
".quad 0x00a80f00af351800,0x67230f11ae321695,0x2631156723311a09,0x0060041367230f00\n"
".quad 0x0034030086393825,0x0d6723311f67230e,0x37170a006139392f,0x67220f0252040032\n"
".quad 0x0f160f0d026f0f1f,0x00b7074a01d70f1c,0x050432169d00b00f,0x027f321a09027f0f\n"
".quad 0xf2381809027f321f,0x09007f0f04980d00,0x01023d383139384f,0xa30416023d38372f\n"
".quad 0xec36312f00f50b02,0x393425007e09021a,0xc33919007e0a215b,0x760e01a634312341\n"
".quad 0x1001e00f01ee0400,0x066938130701040f,0x0c00860f0341351e,0x05025b0f1803390f\n"
".quad 0x006a0f1b00c0381f,0xad30392f00e7025c,0x331f02ad331a0c02,0x00683030250802ad\n"
".quad 0x3602b50f1f00870f,0x002539181901b90f,0x002c323128002c0c,0x111f32312906d60b\n"
".quad 0x35362a09dc373929,0x1906010936190052,0xde36362b01290201,0xeb00011e7a311f03\n"
".quad 0x04003d06061d0506,0x2b311c11a80404aa,0x00373219007c0c04,0x06006d3119006d0d\n"
".quad 0x0f0103f4311f0008,0x6e8a32164dff015f,0x03a3341a0903a30f,0x0166030803a3341f\n"
".quad 0x1c0123391f00c706,0xac0f066f0e009e0e,0x1d42ce0a332a2508,0x0f00007e0f021b31\n"
".quad 0x1d007e311f0c092a,0x382f11010b38392f,0x0000a40f05453a33,0x1338372f0c068d0f\n"
".quad 0x1b011a31322a1e07,0x2b00b435180f1137,0x00103119001e3338,0x0571bb0f2200cb0f\n"
".quad 0x05a70f16016f321f,0x35302b11e9084fff,0x30290db934190066,0xef0a17db04011e38\n"
".quad 0x321792ff01a40f03,0x351a0906320f019b,0x055a0632351f0632,0x7f076306330f009f\n"
".quad 0x00a5068706340f00,0x2a250d480f06350e,0x01d13719011c3132,0x00ad39392a009704\n"
".quad 0x00cc39302f02cd0b,0x01f00c3606360f21,0xcb3838232103160f,0x01530f0277351e18\n"
".quad 0x0cbb351606b50a0c,0x313529002e203823,0x2b03d537392a0128,0x0472301f01383431\n"
".quad 0xe6050ece37362609,0x0400ae030511030d,0x35392f019103010a,0x000830312e00000e\n"
".quad 0xc80f00260300cf07,0xe6351300b202ae00,0x04fd393636344f01,0xfd361f04fd361a03\n"
".quad 0x07006037362f0704,0x0271263432393843,0x8e5a0f860a030032,0x00660816e130130c\n"
".quad 0x66311f3e0066351f,0x1e3e0066361f0000,0x6d0201f204006632,0x04db040c00dc0f03\n"
".quad 0x0076331f1b00760f,0x00760f0900ec0f10,0x2e0f100076341f23,0x35172300760f0902\n"
".quad 0x0000350a37220076,0x3328000b32158e46,0x1f0303120f74a131,0x04f8371f0c064b32\n"
".quad 0x0101590f005f050e,0x07a6032e0cdf331f,0x450f04660a095106,0xec0702bb37190202\n"
".quad 0x36190054636f2a80,0x004d0e00640004c9,0x004d381d14b53914,0x160180031eb32b19\n"
".quad 0x2e006e3019018036,0x21303629005e3730,0x04a63213005e0a01,0x2a1b1b040006301d\n"
".quad 0x0f01e40008563131,0x05ce0a312f011d78,0x3302020338372600,0x2805ed0900430a00\n"
".quad 0x2d1b590406023736,0x0f02040700083630,0x3120026c005c006f,0x1a0d05e60f89c736\n"
".quad 0x0005e6371f05e637,0xb0050b730c020b04,0x03026c0108470800,0xc00f0341361e01ee\n"
".quad 0x01d50575b8041408,0xab0e0900df30322f,0x006a0e1603b30f08,0x03be0f0d006a311f\n"
".quad 0x006a321f006a0e24,0xea301f25181a0f0c,0x5a02ea30352f3c02,0x1702ea0f1a006f0f\n"
".quad 0x02ea381f02ea381a,0x0f2f02c40a342f74,0x110044381f0e0044,0x0277321f3402780f\n"
".quad 0x3529015933382901,0x0e203739290ace32,0x005807001e353125,0x435034133e00510f\n"
".quad 0x00023a0f06d63119,0x023a391f023a391a,0xb6351e00f5331a00,0x01f10f0700700f04\n"
".quad 0x040c016934302f04,0x0f580d00f504028d,0x1900860d01233715,0x0c9039392b7a3834\n"
".quad 0x050c890300413018,0x0f0a0a04c0040edc,0x99040cb003001f04,0x1200860d08a50801\n"
".quad 0x025d35312303bd32,0x0800780f29008f0f,0x9a000044030d7307,0x0248081b25bb0f01\n"
".quad 0x4f3414008230382b,0xf4391f0201870f0b,0x910f6300f20f0308,0xa80401e3361a1312\n"
".quad 0x321f00c338382805,0x0f063e391aa200c3,0x0f0456391ba000c3,0x00253137296300c3\n"
".quad 0x170242301f00a506,0x023b099b00bc351f,0x02340a9a00bc361f,0x3338296400bc371f\n"
".quad 0x1900cb37392b056f,0x001737372b008131,0x2f003705032d3919,0x19381f3605963132\n"
".quad 0x1900c534382b4e01,0x0b0a383028010b36,0x08c905168c31312a,0x361f001f0500360b\n"
".quad 0x9c0300fc071d0139,0xb6371a4900fc0f0f,0x33382f10cf341900,0x00a730312f11172c\n"
".quad 0x301a0302bc371f4e,0x0b0081363129009f,0x1a6500a6311f3830,0x00b633302b004c32\n"
".quad 0x341900740602410a,0x312f31034c0f0020,0x9037362f4e010432,0x012a040114070d00\n"
".quad 0x331f02025836312f,0x01b33130295603fb,0x311a006c03002706,0x2b00e535362a0028\n"
".quad 0x03e6301a19c53530,0x00e808041531302d,0x03d831195008300f,0x312c0be403005c04\n"
".quad 0x0300e03419005531,0x0cec0e003e00007b,0x030054065003510f,0x1b301a0027040170\n"
".quad 0x1b03027434312f03,0x0400273519036931,0xb5371f05e10d0037,0x1b4f00fe361f0a01\n"
".quad 0x290091341a003635,0xc439362a01163130,0x02ea050106301902,0x00f60c04029b351f\n"
".quad 0x4b00ee0f87d8361e,0x053736312f033c0b,0x270e006c34312803,0xf138312f09560900\n"
".quad 0x0b301f0090055504,0x3c48fb32312f0106,0x2c5000e40f87ad0e,0x006c321a00993131\n"
".quad 0x37371c005e363129,0x322f0105e7301f00,0x443737284e00c030,0x4c34382a03220400\n"
".quad 0x3719003d30322b00,0x3934280028040045,0x45301b23040c008a,0xec31322f08850a00\n"
".quad 0x2800ec30352f9b00,0x32372aa900ec321f,0x129f0507e30c0045,0x9800f40f875a361e\n"
".quad 0xec341f2801e0311f,0x351f2401e00fa900,0x321e00910da900f4,0xc00f2587680f01e0\n"
".quad 0x311a005b361a7e03,0x00ce04003c0b0054,0x0122371f1c03020f,0x2a381f2403380fdf\n"
".quad 0x31291300360fe001,0x1f5011f30f012935,0x005c32180400fb33,0x35362a005c393929\n"
".quad 0x30312a01f70c0d82,0x0a001831302a0a1a,0x5e0ab330332f020f,0x312a02006b30312f\n"
".quad 0x0f0000b80f000834,0x332f01d50b200d8b,0x005b39184f00ef31,0x00a6040015393928\n"
".quad 0x2e0069301b1d820b,0x4b02890f89bc3030,0xd1353029005b3019,0xc435312a006a0412\n"
".quad 0x00a6331f00a60a43,0x001f0403004c0f65,0x0900bd311f002e05,0x5b34184f00bd341f\n"
".quad 0x2403004b39392900,0x6331312c00b40400,0x4b00ae0f8ae50e00,0x39302b003637372a\n"
".quad 0x362b054b341903c4,0x26146b311f037537,0xae0c56138336332f,0x38302f009f381900\n"
".quad 0x0f0334361e22010a,0x010437332f1e08ce,0x523718075e36194f,0x5a34302900700400\n"
".quad 0x60331f1318830f00,0x5c04002d34196e08,0x05008103002e0400,0x63090082301b0199\n"
".quad 0x008a381a00b70600,0x4a34312a003e301a,0x08351b0055391b02,0x123e0f0303d70f00\n"
".quad 0x0502b50300890b50,0x06ad0b0073020697,0x00180501f930312b,0x00b131312a03fd0a\n"
".quad 0x00e6070084303125,0x23034e00e630342f,0x2b00b80300a80630,0x0074311900103230\n"
".quad 0x000805000f353629,0x450b008506005603,0xf008012f35302d01,0x31196d0310311f00\n"
".quad 0x008203003d060017,0x190b3a0f0513540f,0x342f280dbc31302f,0x1400360fdf012232\n"
".quad 0x0dbb0f911d35312e,0x003c04005b3119a6,0x430f00440406650b,0x0fd7011a351f1c02\n"
".quad 0x132f36342f230ff7,0x00a10810020e0fa9,0x371f02010338302f,0x0f1801f70fa90103\n"
".quad 0xa9010b381f10010b,0x1001030f10153d0f,0x005b30195011df0f,0x092f04004c353029\n"
".quad 0x01d8040129393029,0x351f080eb030312f,0x2201aa341f500945,0x04361a00c730382b\n"
".quad 0x02cd0c03019b0f01,0x31302a0d6d30312a,0x63351f01cb0a0018,0x12060b0b31059c09\n"
".quad 0x1132312f0302110f,0xd0011332352f2007,0x1c01490f23097b0f,0xc40f07ff0149331f\n"
".quad 0x65015034352f220a,0x12176c0f006a331a,0x1f1c023a0f0beb0a,0xea30352f9900ea35\n"
".quad 0x0fa700ea361f2800,0x00f237352f230b4e,0x381f2801dc311f9a,0x1f2401dc0fa700ea\n"
".quad 0x1f008f0da700f239,0x8c30362f0901dc32,0xfc0f2306b30fa705,0xdd012031362f1b02\n"
".quad 0xff32362f2207d30f,0x1f500247361fa805,0x2305790fdd012033,0xde012f301f7a4204\n"
".quad 0x0fa0610e2103760f,0x0f220c9e0f90012e,0x180fa0860e1b036e,0xd90e5104bc0fc501\n"
".quad 0x03bb0ffb014e0fa0,0x90105b0fa1490e24,0x290248050208610f,0x0044391a002d3031\n"
".quad 0x0121371f1d02760f,0x01210f02060c0fa1,0x540fdd0121381f44,0x0fdd0129391f2508\n"
".quad 0x301f33f2090f1782,0x1f00f2371bb00120,0x03001f060905e932,0x20038230302f0074\n"
".quad 0x0ff001390f48b604,0x410f7ed60423039a,0x021b7137362f4f01,0x0a0037371a00d706\n"
".quad 0x54391b01050b0846,0x363129041ac90f08,0x36302a01c40b0811,0x0f7f0234312e0092\n"
".quad 0x1b0f04dc0ddb011a,0x0f54011b341f0501,0x321900f30c49093a,0x001f0601570b017e\n"
".quad 0x1f2404af0f007403,0x1733382f54013935,0x00de040046040200,0x00260506ad393029\n"
".quad 0x1a00180a01ff351d,0x351e009c05001036,0x570c0001440f0c2f,0x660a004736312a00\n"
".quad 0x49026e0f6d750b1f,0x4c313128032d311a,0x3729001e37392903,0x392a038d09025536\n"
".quad 0x011831302f057439,0x361a550118371f42,0x100f005339180110,0x196b0496381f7801\n"
".quad 0x00bd37392a00e533,0x3119047f0900b507,0x31322a00f40d038a,0x2a01008a311f0a4f\n"
".quad 0x2a03a10d018c3131,0x9831312a00563231,0x0f00080603c00a01,0x6b0169391f0405ff\n"
".quad 0x0c08e205004c321a,0x34190074361b00be,0x312f7a01530f0020,0x0f30312963015330\n"
".quad 0x3029005433302a00,0x0e1002ad0f011430,0x33195100d50fa63e,0x1201c337372f0165\n"
".quad 0x3c371a01ca31372b,0x00e404011b311932,0x00e40fa64630302e,0x3d094b00d5371f52\n"
".quad 0x7431312a00170300,0x362e016f32312800,0x1b52010a0fa66338,0x372a002f0a009e32\n"
".quad 0x312a01200a00e338,0xe30b006c0d007431,0xb20f00c835302802,0x54013334312f3c06\n"
".quad 0x302a030be033382f,0x0b61313029024d33,0x00de050074313129,0x301a003d0b0c640b\n"
".quad 0x0500d73130290104,0x31190010351b0083,0x730404059e0f009b,0x004d060536c50f16\n"
".quad 0x095101680fa7440e,0x38391f054f0c0044,0x0f000f3737290a01,0x3928008a07090565\n"
".quad 0x312f003d05000739,0x0f9af9040116f335,0xd60c0a00034f0109,0x0834312a006b0c01\n"
".quad 0x005604002f301a00,0x001e040092373728,0xc00401660b000f04,0x21371f2301210f00\n"
".quad 0x1a044a39302b6401,0x0317313029003737,0x050047301f01310c,0x04d00f050366361f\n"
".quad 0x44035b013e381f42,0x01dc351a00ce0500,0x530511063031302f,0x7b301a012b311b01\n"
".quad 0x37371b010c301a00,0x05320a9d34312f00,0xf20e0415301e0038,0x0c54017139312fe2\n"
".quad 0x0300bf35312f00d7,0x35240200c738312f,0xbe0fa83b0e14f837,0xbe0fa803311e9700\n"
".quad 0x321f00be321a9800,0x321f0008056404e9,0x4605006c030200c6,0x2a0200dd35302f00\n"
".quad 0x0704090a002e3637,0xf4331f00f4091070,0x381e00aa371bab00,0x1f040700b30f0037\n"
".quad 0x1f30014939392f00,0x6a37372f54014934,0x1f017536362b7511,0xe939372a01415c31\n"
".quad 0x09015700019b0b01,0x1a5b0291351f0148,0x720572900b01de39,0x0304800f04610500\n"
".quad 0x290020361901410d,0x0a00fc06007b3738,0x36322f097a0c0027,0x0f2707c10f730e3f\n"
".quad 0x011637322f410952,0x00f6060044381954,0x09d40f680cba311f,0x1309dc0f00100408\n"
".quad 0xdd34382f0400d60f,0x92381f02a80d1000,0x19027938362b5401,0x0036393929004634\n"
".quad 0x03e2040072373728,0xa00400f90b000f04,0x01a10f0401b90f00,0x311a009f30372a1a\n"
".quad 0x410e0801a10f004c,0x06520e5001560fc3,0x0300bd041705de0f,0x005b09007a060256\n"
".quad 0x2b0082381a00af06,0x006d341f040b3530,0x00273519012a0c03,0x34332f078c31312a\n"
".quad 0x00a032382f54013a,0x4e04002731302903,0x07b5301f00180500,0x014a3918014a0989\n"
".quad 0x353b6a720064ba07,0x0847800e00f73831,0x015c1b36312f6957,0x004e0d03a031312a\n"
".quad 0x1500d1311b03a904,0x060100d10f4c8530,0x30026343331a002f,0x0b021a0f5ab10862\n"
".quad 0x392f0200a730372f,0x0400c50411022139,0x2f002737312a0549,0x3635250105513130\n"
".quad 0x3034396967006a61,0x380f0122065aa938,0x0a019e33382f1202,0x009b0f0237363629\n"
".quad 0x29009e0578e6080a,0x2f00aa0869203531,0x3531290301943939,0x39034d00a20f006c\n"
".quad 0x080f0000a2381f00,0x0f001e3739290403,0x0f5e3036144e00a1,0x020bf1391f0100a2\n"
".quad 0x1a01fd0c02043419,0x150f00830002cc39,0x750b0800c0003304,0x57ee33332400c001\n"
".quad 0x0341311b0000c00f,0x2d0c0a0a5539302f,0x542f3814006d0e01,0x0183321a01006d0f\n"
".quad 0x29006d361a012e04,0x0301d90501173431,0x00850f00b40000f9,0x0085321fb29c040b\n"
".quad 0x3019005538392b00,0x31322f007c040193,0x4eaf32352443009b,0x008b391800009b0f\n"
".quad 0x31302a024439392a,0x30280244311a07b5,0x008835362f0d3931,0x00880e4b5738140d\n"
".quad 0x4938362b0402d50f,0x04003503013a0a03,0x33382900bf030113,0x00b60301670c00c6\n"
".quad 0x0d3c03cb0f00c604,0x0fd93a35362400f6,0x850a009c0c09029e,0x0d00c40f00200500\n"
".quad 0x001031190118341b,0x0504088d0f0ee20c,0xe30f07a7371e0259,0x0645ab3437240200\n"
".quad 0x35302907d40b00e3,0x004a0200540901c1,0xdc04078508084800,0x6239302f003b0500\n"
".quad 0x8b33120141030014,0x2b01603631343a00,0x3033372a00bb3032,0x003c0b004c381800\n"
".quad 0x3c341f12003c311f,0x1712003c321f0200,0x430905301f003c38,0xc23718090178351f\n"
".quad 0xa00400ad0a392400,0x7593060019321a09,0x353335384f00190b,0x1933313938030019\n"
".quad 0x2b05004b30322d00,0x0500760d02560400,0x39384c00440b73cf,0x2b36322800443533\n"
".quad 0x02fd4b7d0a303d00,0x0571bd7b2028fd22,0x52d404000f020273,0x6a0c06005f0a3827\n"
".quad 0x29003c0a001e331a,0x004b321a0c643631,0xf00a322900073019,0x2900d6030042044a\n"
".quad 0x2a00760a003b3436,0x00a4351b00943835,0x39d90400bc33372a,0x38190300b632312f\n"
".quad 0x065a06007b04000f,0x351e008a03000708,0x321a26009f0f0015,0x19002f31352a0070\n"
".quad 0x6306009d03000730,0x00910a0200530f03,0x0049351900073019,0x83311b010031311f\n"
".quad 0x0a302f0d00d20f00,0x0a0016341a000e62,0x6ab1045cde040062,0x1f005f371a002c0a\n"
".quad 0x2b00420a1101b731,0x5d01a80f18d30a38,0x533619470081391f,0x71391f0096301b00\n"
".quad 0x6704000730190800,0x82351f3903830f00,0x0e006b3231294f03,0xbb0f07013a0f00e9\n"
".quad 0x0f006133312e1301,0x1e3100800f160100,0x292d00530f008034,0x0200d40f59b50a35\n"
".quad 0x00080c004332362a,0x0405ec0400073019,0x420c0100520f01ac,0x00610b0c01360f00\n"
".quad 0x4000620f00f1371a,0x0c00620a0304850f,0x352b0c00620f0178,0x190010301a001f36\n"
".quad 0x01a039312a000730,0x0f01dd0f0074301e,0x450f59d206081402,0x1f0e0007301c0400\n"
".quad 0xf97b20636f6c5608,0x7a3813000b351682,0x177e270f6f830804,0x6703000b02004507\n"
".quad 0x3829010050371f0c,0x342900180519ef31,0x03004d341a25ab37,0x1f0800a8341f000b\n"
".quad 0x0b3432261400a838,0x391f080050371f00,0x0f8884371e150050,0x0a000730193200fc\n"
".quad 0x00ba0e0ca60784ab,0x09003d07077d9c0f,0x382f0200480f01aa,0x1f019a080c004834\n"
".quad 0x0048381f08009035,0x48381f000b35160d,0x1e0d0048391f0800,0x192200dc0f018a38\n"
".quad 0x3715013200000730,0x0f00a230322e000b,0x0b03003d0710807f,0x00480f027c311e00\n"
".quad 0x00027c0f02840404,0x0f00ee0e000b3116,0x0b03003d07118d0e,0x311f080048391f00\n"
".quad 0x0f0182391e0d0048,0x0e000730142a00e4,0x65676e61728f0ced,0x706100001b7b2073\n"
".quad 0x01001d6d616e6275,0x0055333932350d3c,0x694587b906008902,0x0a38120d746f666e\n"
".quad 0x1234333535334900,0x172d30302f0eb104,0x1f107a0600c90317,0x0005ef053c174436\n"
".quad 0x19f93336363f00b2,0x001f391900750527,0x4800ba044a00c20f,0x3928193930353334\n"
".quad 0x20044b5803344730,0x0500730f00c70311,0x010537372b135a04,0x0c35382812233818\n"
".quad 0x753218162c361b01,0x2f00000e35392f12,0x00d2050300083031,0x085686663530383f\n"
".quad 0x05343738334f0079,0x0c00253937281002,0xd4391f105f03002c,0x2915353231290202\n"
".quad 0x5235362a01773739,0x138b0b00e2361a00,0x1017cc0f0f17870f,0x0b007c0e0401ea0f\n"
".quad 0x03560a006d0c148c,0x0400373119000806,0x11016730392f0169,0x000838302a01240a\n"
".quad 0x2f1211301500820a,0x02810300023c3731,0x02ad32352900d704,0x1e3519037137392a\n"
".quad 0x3937334f005b0400,0x0113311b1100b737,0x00cf0520043a361f,0x4bff02923339303f\n"
".quad 0x172d3631293c7f0c,0x37392d0085303729,0x37362f01ae091a7e,0x04b939342f3504b9\n"
".quad 0x00000a7d0a305019,0x0000005001010002,0x0000000000003dc0,0x0000000000003dbf\n"
".quad 0x0000004600010007,0x0000000c00000040,0x0000000000002013,0x0000000000000000\n"
".quad 0x00000000000110e8,0x754d78697274616d,0x0000000075632e6c,0x010102464c457fa2\n"
".quad 0x0002660001000733,0x40310001006e00be,0x0804802200080110,0x40004605460df500\n"
".quad 0x2f00400003003800,0x7368732e00000100,0x082e006261747274,0x11f500086d792700\n"
".quad 0x2e0078646e68735f,0x006f666e692e766e,0x75632e747865742e,0x636f6c6c614d6164\n"
".quad 0x002505001409001a,0x6547636e754602fa,0x7562697274744174,0x001f0f0030736574\n"
".quad 0x76654468003b0500,0x003c0a001e656369,0x32003d050100200f,0x00340a0020746547\n"
".quad 0x1afa002b05001707,0x636e61707563634f,0x6974634178614d79,0x736b636f6c426576\n"
".quad 0x69746c754d726550,0x6f737365636f7270,0x0f1800370f004b72,0x687469578b1f006b\n"
".quad 0x400f012f67616c46,0x5f0bf6007d012100,0x697274616d33315a,0x414455436c754d78\n"
".quad 0x69695f535f536650,0x6f0c00240f006169,0x0026646572616873,0x6174736e6f639f0d\n"
".quad 0x723f0a002930746e,0x00c10317002d6c65,0x43336d6964344ec8,0x0a00b56a6a6a4531\n"
".quad 0x2d321f002d0a0018,0x65724f0018030400,0x00a100030032616c,0x880015030200190f\n"
".quad 0x6673626166347453,0x005808001406005c,0x110100150e001609,0xb4000f03004b0b00\n"
".quad 0x72665f6775626564,0x696c42000d656d61,0x360010080049656e,0x735f54000f5f766e\n"
".quad 0x00180e0162737361,0x5f78747077001405,0x007c000012747874,0x05002a6765725f4c\n"
".quad 0x008a707974350018,0x0076657262626186,0x110800f70300232e,0x07041101000c0200\n"
".quad 0x000b02000f0000b4,0x0e7365676e617264,0x6d616e6275708900,0x0a00150900497365\n"
".quad 0x0012020026070113,0x0204b30f1f04cd0f,0x1403037908001104,0x001c0f0d04b30f00\n"
".quad 0x0f001f0e00460800,0x0801001d0f0e04c8,0xe80f0000200f0053,0x0042080014070504\n"
".quad 0x0a4e04d60f001706,0xfb0f1700370f0079,0xab0e25003d0f2a04,0x04d20f1a00400f00\n"
".quad 0x0455060700210f10,0x1a056a0f2d05930f,0x61705f7f1500290f,0x030b057c006d6172\n"
".quad 0x00180900cb060015,0x0d00150a0405580f,0x05b80f001803003c,0x0600110805730d1f\n"
".quad 0x001a030014050067,0x6507003006006908,0x010f45ff05c90f00,0x2200000032840600\n"
".quad 0x0380210001002400,0x0000003d00670007,0x6211000104001803,0x11003000252c0030\n"
".quad 0x000100252c003078,0x010026250030b311,0xca00210006021100,0x20000100262c0030\n"
".quad 0x0100272400300107,0x1500300b67801200,0x000100272c003001,0x0100282400304011\n"
".quad 0x306e110060801300,0xd911000100282c00,0x0300010029250030,0x2c000b02102000c0\n"
".quad 0x0018a81100010029,0x0340200001002a2c,0x11000100232c0018,0x0001002b2c00187f\n"
".quad 0x01002c2c0018bb11,0x2d2c001804262000,0x012b044b20000100,0x0108030001002e25\n"
".quad 0x01002e2c00307c11,0x00042c0018971100,0x052c0018a4110001,0x2c0018c011000100\n"
".quad 0x0018ec1100010006,0x18fe11000100072c,0x1620000100082c00,0x000100092c001805\n"
".quad 0x01000a2c00182e11,0x000b2c00183c1100,0x0c2c001859110001,0x2c00188111000100\n"
".quad 0x00028d660001000d,0x6b80130198101200,0x0012000003707d00,0x1d0018ac1102102b\n"
".quad 0x0018041b2002b82c,0xffffff6602a02d1d,0x000100000c0540ff,0x80947c04000309f0\n"
".quad 0x00288080810c2880,0x808007288081ff08,0x8207103000050880,0x8410000583100005\n"
".quad 0x8610000585100005,0x0828808080500005,0x2880808208500033,0x2810002328100023\n"
".quad 0x2830002328100023,0x0005881000058707,0x00058a1000058910,0x00058c1000058b10\n"
".quad 0x00058e1000058d10,0x0005901000468f10,0x0005921000059110,0x2000970000059310\n"
".quad 0x0005961000059508,0x0005981000059710,0x00059a1000059910,0x00059c1000059b10\n"
".quad 0x00059e1000059d10,0x0005a01000559f10,0x0005a2100005a110,0x0005a4100019a310\n"
".quad 0x0005a6100005a510,0x0005a8100019a710,0x0005aa100005a910,0x0005ac100019ab10\n"
".quad 0x0005ae100005ad10,0x0005b0100019af10,0x0005b2100005b110,0x0005b4100019b310\n"
".quad 0x0005b6100005b510,0x0005b8100019b710,0x0005ba100005b910,0x0005bc100019bb10\n"
".quad 0x0005be100005bd10,0x0005c0100019bf10,0x0005c2100005c110,0x0005c4100019c310\n"
".quad 0x0005c6100005c510,0x0005c8100019c710,0x0005ca100005c910,0x0005cc100019cb10\n"
".quad 0x0005ce100005cd10,0x0005d0100019cf10,0x0005d2100005d110,0x0005d4100019d310\n"
".quad 0x0005d6100005d510,0x0005d8100019d710,0x0005da100005d910,0x0005dc100019db10\n"
".quad 0x0005de100005dd10,0x0005e0100019df10,0x0005e2100005e110,0x0005e4100019e310\n"
".quad 0x0005e6100005e510,0x0005e8100019e710,0x0005ea100005e910,0x0005ec100019eb10\n"
".quad 0x0005ee100005ed10,0x0005f0100019ef10,0x0005f2100005f110,0x0005f4100019f310\n"
".quad 0x0005f6100005f510,0x0005f8100019f710,0x0005fa100005f910,0x0005fc100019fb10\n"
".quad 0x0005fe100005fd10,0x02a880100019ff10,0x0005821000058110,0x0584082880818360\n"
".quad 0x0586100005851000,0x0588100019871000,0x058a100005891000,0x058c1000198b1000\n"
".quad 0x058e1000058d1000,0x05901000198f1000,0x0592100005911000,0x0594100019931000\n"
".quad 0x0596100005951000,0x0598100019971000,0x059a100005991000,0x059c1000199b1000\n"
".quad 0x059e1000059d1000,0x05a01000199f1000,0x05a2100005a11000,0x05a4100019a31000\n"
".quad 0x05a6100005a51000,0x05a8100019a71000,0x05aa100005a91000,0x05ac100019ab1000\n"
".quad 0x05ae100005ad1000,0x05b0100019af1000,0x05b2100005b11000,0x05b4100019b31000\n"
".quad 0x05b6100005b51000,0x05b8100019b71000,0x05ba100005b91000,0x05bc100019bb1000\n"
".quad 0x05be100005bd1000,0x05c0100019bf1000,0x05c2100005c11000,0x05c4100019c31000\n"
".quad 0x05c6100005c51000,0x05c8100019c71000,0x05ca100005c91000,0x05cc100019cb1000\n"
".quad 0x05ce100005cd1000,0x05d0100019cf1000,0x05d2100005d11000,0x05d4100019d31000\n"
".quad 0x05d6100005d51000,0x05d8100019d71000,0x05da100005d91000,0x05dc100019db1000\n"
".quad 0x05de100005dd1000,0x05e0100019df1000,0x05e2100005e11000,0x05e4100019e31000\n"
".quad 0x05e6100005e51000,0x05e8100019e71000,0x05ea100005e91000,0x05ec100019eb1000\n"
".quad 0x05ee100005ed1000,0x05f0100019ef1000,0x05f2100005f11000,0x05f4100019f31000\n"
".quad 0x05f6100005f51000,0x05f8100019f71000,0x05fa100005f91000,0x05fc100019fb1000\n"
".quad 0x81fe560005fd1000,0x00302f0540002880,0x0007037021030001,0x05590106bb040020\n"
".quad 0x880f000bb8040834,0x3f4dffffffffff05,0xffffff0588000588,0x101b000c046fffff\n"
".quad 0x042511e8f013110b,0x058060042f057d00,0x90394dffffffffff,0x0580701f00010010\n"
".quad 0xffffff0b00401f00,0x00010016294effff,0x00901f058002702e,0x294effffffffff0b\n"
".quad 0x0580f01f0001001b,0x0800080580a01f00,0x23002505810f1f00,0x8107102f05810f00\n"
".quad 0x8107100581071005,0x8107100581071005,0x8107100581071005,0x8107100581071005\n"
".quad 0x8107100581071005,0x8107100581071005,0x8107100581071005,0x0710050581071f05\n"
".quad 0x0710058107100581,0x10050581071f0581,0x1005810710058107,0x050581071f058107\n"
".quad 0x0581071005810710,0x0581071f05810710,0x8107100581071005,0x81071f0581071005\n"
".quad 0x0710058107100505,0x071f058107100581,0x1005810710050581,0x1f05810710058107\n"
".quad 0x0581071005058107,0x0581071005810710,0x810710050581071f,0x8107100581071005\n"
".quad 0x0710050581071f05,0x0710058107100581,0x10050581071f0581,0x1005810710058107\n"
".quad 0x050581071f058107,0x0581071005810710,0x0581071f05810710,0x8107100581071005\n"
".quad 0x81071f0581071005,0x0710058107100505,0x071f058107100581,0x1005810710050581\n"
".quad 0x1f05810710058107,0x0581071005058107,0x0581071005810710,0x810710050581071f\n"
".quad 0x8107100581071005,0x0710050581071f05,0x0710058107100581,0x10050581071f0581\n"
".quad 0x1005810710058107,0x050581071f058107,0x0581071005810710,0x0581071f05810710\n"
".quad 0x8107100581071005,0x81071f0581071005,0x0710058107100505,0x071f058107100581\n"
".quad 0x1005810710050581,0x1f05810710058107,0x0581071005058107,0x0581071005810710\n"
".quad 0x810710050581071f,0x8107100581071005,0x0710050581071f05,0x0710058107100581\n"
".quad 0x10050581071f0581,0x1005810710058107,0x050581071f058107,0x0581071005810710\n"
".quad 0x0581071f05810710,0x8107100581071005,0x1b88080581071605,0x7013000100211039\n"
".quad 0x12058502042b26c8,0xffff0b080f059986,0x2e8b981b4affffff,0x981f0583082c3804\n"
".quad 0x45ffffffffff0580,0x2c182f000100c822,0x00140420700d2690,0x16021104cd050000\n"
".quad 0x0b031204dd051000,0x000b041204ed0000,0x00000b051204fd00,0x1d00000b0612050d\n"
".quad 0xdd8211000b071205,0x82060000006c6126,0x0a9015000a000037,0x0a9215000a911500\n"
".quad 0x0a9415000a931500,0x1e08042500890000,0xffff0ba000002f06,0x0100702247ffffff\n"
".quad 0x0327280032383f00,0x090620101f062005,0x0f05f4381a05e902,0xffff11680f0905cc\n"
".quad 0x003800394affffff,0x60116800f02e0001,0x3d8b02530000001c,0x0efb01803ef9f510\n"
".quad 0xf005f2010101000a,0x2f656d6f682f0121,0x796d2f6169646162,0x69662d6c6f726163\n"
".quad 0x6d2f7365646f632f,0x752f006d656c456d,0x6c61636f6c2f7273,0x31312d1af2446c2f\n"
".quad 0x2e2f6e69622f302e,0x7465677261742f2e,0x34365f3638782f73,0x692f78756e696c2d\n"
".quad 0x00396564756c636e,0x2b2b632f85000d03,0x2ec043b20000392f,0xfcb5849b01007563\n"
".quad 0x645f2100575abe05,0x6e75725f0ef145e5,0x6970615f656d6974,0xf8b0c6ba0200682e\n"
".quad 0x6f7463657674fa05,0xfc00177315421972,0x6874616d6366f607,0xc605ecfbf8eb0300\n"
".quad 0x26657669726402ff,0x09000006b4b96300,0x0303fa180a001280,0x02020320020100ca\n"
".quad 0x01000290020104d0,0x0b001fcf1f001f01,0x0101f06d001fd413,0x1fd914001f01f002\n"
".quad 0xde1302001f001f00,0xb0020103b05e001f,0x031f003ee314001f,0x1e3103013200003e\n"
".quad 0x0201030104c0b100,0x4000060403010180,0x2010001202800202,0x7f03010390500005\n"
".quad 0x01c0020303500041,0x90020102b05c0006,0x0032a4030330004d,0x1b05001904e0023f\n"
".quad 0x00f3030420001906,0x0101a00201032090,0x5703d120011ac002,0x7200190257101d02\n"
".quad 0x1100382003003001,0x250006011300ac05,0x2400050600b200c0,0x6c00d000001b00f0\n"
".quad 0x00b9e002f0f0f001,0x3a004f00c503004f,0x00383012004fe813,0x0103480007f0f022\n"
".quad 0x003ac01e003ad002,0x006ef41401108810,0x01a7220c002ec01f,0x6d0200abe0180068\n"
".quad 0x4901902700100a00,0x49cd130049801e00,0x0604003900f02b00,0x9801902f000b0a00\n"
".quad 0x05030101f8520b00,0x15000600e02c002b,0x080047c014000530,0x14021f0005020126\n"
".quad 0x1f05019df0150200,0x00001ef101013801,0x001107001b060018,0x190050f111002201\n"
".quad 0x03f1f0f05c001ce0,0xf00100d04b003602,0x0c0200182014006d,0x0038090014031400\n"
".quad 0x0098031300082019,0x0e004a0210004a0e,0x8902eb0300420342,0x3c00010100c02701\n"
".quad 0x0241021200a30200,0x30022d0006f00224,0x2e0043038f220043,0x0037061000430280\n"
".quad 0x7f03f0f1f1012073,0x020110020b7d00a5,0x0047c014004702a0,0xf132003f0100b904\n"
".quad 0xfaf3f42f000c7e03,0x03e6032002550102,0x0300e42002240189,0x737265762ee20001\n"
".quad 0x00302e37206e6f69,0x5f6d73208105ea2e,0x02f0470e202c3037,0x7365726464612e00\n"
".quad 0x3620657a69735f73,0x636e756681003534,0x206049832e282020,0xa20013203233622e\n"
".quad 0x306c61767465725f,0x2200282948752029,0xef49b60200140000,0x00003b002900305f\n"
".quad 0x004f206b6165772e,0x270053074cfe060f,0x1f00580400183436,0x290031710b00202c\n"
".quad 0x02f106e92e007b00,0x206e67696c612e20,0x5f092038622e2038,0x7065645fb000165f\n"
".quad 0x153b5d385b30746f,0x25094c0049671248,0x00114c1500105053,0x3c7225092032339a\n"
".quad 0x3c64726000233e33,0x00ce0101663b3e35,0x3a306e6967656270,0x2e766f6d00640010\n"
".quad 0x007c202c2a004875,0xa561747663003b62,0x73202c2200260500,0x190300e1646c2200\n"
".quad 0x5b202c3164726e00,0x1f00002a5d1f0108,0x3b5d313701002a32,0x202c33647250008e\n"
".quad 0x147473003b420055,0x302b5053255b7100,0x04002c331200175d,0x32220015341100df\n"
".quad 0xe56578653700e03b,0x5a3a30706d745200,0x202c3163011f0200,0x0132020041393939\n"
".quad 0x02006b0200403210,0x5b093233625800aa,0x3b327000790301fb,0x2300553b74657200\n"
".quad 0x30646e7f006e3a31,0x0d1b0245007d003a,0x00230d02500e4efd,0x0f06002b0f025b0e\n"
".quad 0x450266311f1c0266,0x66311f0f0266311f,0x02710f011e0e2302,0x27027c0f01280e10\n"
".quad 0x7c31120c027c751f,0x331641027c321f02,0x0e1f027c311f027c,0x00240e027d0e5126\n"
".quad 0x2c32332f00027e0f,0x1218002c311f0b00,0x371e02620902ab32,0x32120262331c0262\n"
".quad 0xc90f0202280f0262,0x019f050229090b00,0x0035311f0b00350f,0x321b0d0035321f00\n"
".quad 0x0607025e3214025e,0x0100130a02720102,0x58321202580e0259,0x58351e0258341902\n"
".quad 0x58351f003f361002,0x3516025836161002,0x051f0258321f0258,0x001b05024f0e533c\n"
".quad 0x0f04507b1e077707,0x0c01ee331f0e01ee,0x0b05017b0f00680d,0x361f015533120155\n"
".quad 0x1f015537164103ad,0x16536d0f1f015533,0x0e16003b0f01750e,0xf20e2500430f0659\n"
".quad 0x86321f1e00430f03,0x1e00025e331f2f00,0x1f025e371c025e35,0x2201510f08025e34\n"
".quad 0x4d321f004d3b5d2e,0x990f01047b0f2400,0x1f010099321f2500,0x115d332f24004d33\n"
".quad 0x35140015311a0107,0x130100034d0f0726,0x0763361401950300,0x1903a1341203a10b\n"
".quad 0x0103a1331e03a138,0x161003a1331f007d,0x1f03a1391603a134,0x0f570e054803a134\n"
".quad 0xb30e0044052803aa,0x03bc0f27004c0f03,0x0098321f004c0c30,0x341f380098331f38\n"
".quad 0x351f041a0d130866,0x230f00910c35041a,0x042c0f00560c3304,0x3304350f00550c32\n"
".quad 0x0000ab331f00560c,0x0493341f2e01000f,0x0f04930c09870925,0x2f04a635120f099c\n"
".quad 0x72311741099d3031,0x7600ff04a835130e,0x2e20656c62697369,0x5aa6207972746e65\n"
".quad 0x0f00ff0004740607,0x300f04580e070028,0x071c0030311f1200,0x00300f0125000420\n"
".quad 0x341f1c0030331f0e,0x87041835171c0030,0x7025092064657270,0x663601b666110a91\n"
".quad 0x3e01c802042b393c,0x36312c043e34323c,0x4f04043f361f043f,0x0f0e009c5b202c34\n"
".quad 0x11003a351f020423,0x3a361f01003a311f,0x02b40503ec0a1100,0x0103cf0f0f00390f\n"
".quad 0x39341f110039371f,0x2d110039381f0000,0x0387361303875d35,0x25202c39c80f7c0a\n"
".quad 0x3b792e6469617463,0x25202c3031730018,0x2e6c75730018746e,0x3132001b732e6f6c\n"
".quad 0x3031723a0039202c,0x00340308bd000035,0x2c24003164646135,0x03b6323122003620\n"
".quad 0x33150039003a3339,0x34140052781a0087,0xab010087781d0087,0x3431723a003a0104\n"
".quad 0x7819008836130036,0x003702095b020088,0x0b09420000883612,0x00b3202c37330057\n"
".quad 0x36732e4011853611,0x1c02022700001d34,0x6c6873003b376500,0x001b202c383202db\n"
".quad 0x05070706006e3213,0x3512006f38120551,0x2c3463033a03017c,0x3b28000130663020\n"
".quad 0x6622003c00001b00,0x4f012c361b003534,0x030005a030202c38,0x850401df37130063\n"
".quad 0x001d0005e6351400,0x1300b90000803819,0x00643a3928001a66,0x11ac000a32333223\n"
".quad 0x5f3642425a143000,0x332805ee02008831,0x662a01620000523b,0x1200813212065238\n"
".quad 0x73003a31a2004d35,0x0154746c2e707465,0x016f341301267011,0x6e003a3265002200\n"
".quad 0xe0003f000471746f,0x32702540003b3170,0x008c092061726220,0x2e617262003b34a2\n"
".quad 0x3b32220010696e75,0x33323910080300a4,0x8a0000cf01018b3a,0x3a34323001530200\n"
".quad 0x04360000b5030357,0x1602625d1c001e00,0x010249371a009c32,0x25321d001f0200d4\n"
".quad 0x26321a001e311202,0x0c001d0102a90202,0x0752030062010228,0x3612009b3b333137\n"
".quad 0x03009c343123009b,0x00b7000287000017,0x2f03004b36662532,0x1201840200630000\n"
".quad 0xb90709a20100e837,0x01f50200d1351300,0xf70302022c36322f,0x001a03022c371e01\n"
".quad 0x01df0300022c381f,0x00690201983a392a,0x00036f0f02433314,0xab3b3828036f3915\n"
".quad 0x001f030481321000,0x301b00170201640e,0x001c000033010163,0xf10400600101620c\n"
".quad 0x00d8010862301108,0x01375d3121015d01,0x5533170855003b25,0x0019000855361c08\n"
".quad 0x0cbc0a62ab5f2029,0x2f000c940f001c0a,0x24311f0300243233,0x0f100024321f1000\n"
".quad 0x1c0007920f010fea,0x080bd0371f079132,0x0207850a0200690f,0x0b7f0f002d0e006a\n"
".quad 0x076b0f002d321e08,0x331e05002d331f07,0x33300408371206ed,0x01ec751501ec3a32\n"
".quad 0x2100160a0acb5d13,0x0018321d0018342b,0x021b0605085d3823,0x0b021b371f166607\n"
".quad 0x1b321f08021b321f,0x1f10021b321f1002,0x021b321f10021b32,0x321f11021b381f30\n"
".quad 0x19021b321f19021b,0x1b321f19021b321f,0x8c00021b38130b02,0x6d657443010b0804\n"
".quad 0x3b67657259003670,0x3018000b0101b900,0x165b093436420c9b,0x050764000c950300\n"
".quad 0x0d00333233220033,0x0033030016021516,0x0032321f0100320f,0x32321f0032321302\n"
".quad 0x34020032331f0000,0x61634102bd302b33,0x51678f0b05936c6c,0x32002d002800202c\n"
".quad 0x0931140009202c30,0x0033300009321400,0x2f2f7d02f6020c29,0x7165736c6c614320\n"
".quad 0x03123020646e4520,0x0312381e03123516,0x2a6720060c122d0f,0x0700180402be0028\n"
".quad 0x0a0d0c7a7b1b1e94,0x094e0602a7391b17,0x13021b0e00515b1e,0x16021b361f021b39\n"
".quad 0x33663a021b32332e,0x0301e8661b021b32,0x10281401990700fd,0x9b076858202c2201\n"
".quad 0x0200d10801800001,0x9f0f0f125b160967,0x450e008f311b0101,0x33170f453266260f\n"
".quad 0x440f01c33911178e,0x01af0f20430e1420,0x01980c01019d0f0f,0x3213001f73626134\n"
".quad 0x7d231600d00f0141,0x6365732e50000100,0x203f68a807213d74,0xffffffff0001007b\n"
".quad 0xffffffffffffffff,0x9fcfffffffffffff,0x313234203631622e,0x2d00180f13002700\n"
".quad 0x0f0100350f00010e,0x315f009201760001,0xffff000100393735,0xffffffffffffffff\n"
".quad 0x170209ffffffffff,0x001f003830344f0f,0x0a00200035313f09,0x36110c002031322f\n"
".quad 0x333f0a00240f000e,0x35381f0400190033,0x1f080099341e0700,0x08001d0032353f00\n"
".quad 0x2b030e000100382f,0x1200280035363f00,0x3b0800200034373f,0x0082070010363134\n"
".quad 0x2500010f0800120f,0x0f8dc10224fa7d1d,0x023b0e12ffff24fc,0x00ae7b20636f6c5f\n"
".quad 0x1b7b202f90650296,0x03060f905d040000,0x00324871007d26f9,0x004663000001000e\n"
".quad 0x60224e4405100008,0x03530015014e4b02,0x2a03001510050000,0x90b205108f400000\n"
".quad 0x8f3c00002901c025,0x01d0000001500581,0x0f002d1c00000100,0x1500005260052851\n"
".quad 0x5270050000055800,0x2a8fc33264723000,0x8fc03210002a8005,0x543311002a90052a\n"
".quad 0x543311002aa01a00,0x543411002ab01a00,0x2aa0000002305600,0x1540130054341100\n"
".quad 0x2906d001002a0200,0x23064c0100147005,0x0247310f00148005,0x46100e19ff013c0f\n"
".quad 0x501000eb05912305,0x601300eb011b000e,0x009600006d040015,0x006d00d026594b00\n"
".quad 0x0000008067009500,0x13005200011300e0,0x5200011305002990,0xa33311002aa01900\n"
".quad 0x1402938100b02201,0x14c0052907d70100,0x0015d01301650300,0x33002ae019016509\n"
".quad 0x00140601f5040035,0x0014201300143611,0x080a020f00459d09,0x0800802b6b520001\n"
".quad 0x00e0090015401301,0x00e000702b013200,0x2b013304002a601c,0x0015801300b800c0\n"
".quad 0x020014901901c908,0x870f0014a01301c9,0x2700d816002c1a3e,0x1300d80402f40080\n"
".quad 0x710000ae09001590,0xb01900ae01b02b00,0xc01703473111002a,0xa40a1302ca019559\n"
".quad 0x0b1302df01602701,0x0002df01702701a4,0x3433002af01a0092,0x33002a0703c40500\n"
".quad 0x2a801601cf060035,0x2a01d00700353300,0x00103010012c0190,0x02e7366420030905\n"
".quad 0x361000d00703ee01,0x92c01803ef019407,0x02002703ef0a1300,0x540703ef0b130092\n"
".quad 0x2a801a00fb351100,0x7e901c0443351100,0x002a0703da061300,0x002ac01900bc3611\n"
".quad 0x110014d019030602,0x0f0014e013001434,0x050018003b234083,0x0502008d00902902\n"
".quad 0x0205090015a01302,0x020501e02b01f000,0x130205045d32c017,0xf000020509022ed0\n"
".quad 0x01f0000205901c01,0xef091301b101a026,0x0110290219c01c01,0x01ef00021902003e\n"
".quad 0xdb301a021901a02b,0x0805de0402190201,0x02402b01ef00015d,0x0200e40160290219\n"
".quad 0x8df01b01ef000219,0x022d02003e801a04,0x1a022d0900159013,0xb01a022d02002aa0\n"
".quad 0x92c01c022d02002a,0x002a0701f0061300,0x03002af019022d02,0xe1030014021804e1\n"
".quad 0x073ea80f00140304,0x6022016b0059003b,0x0c1300c50506f600,0x1700c501e02704af\n"
".quad 0x1a0180011801fc05,0xaa08003634023bd0,0x15301300c5081b01,0x0180091300b10400\n"
".quad 0x0601563715001406,0xc019015538110014,0xd01a019464120014,0x2ae01a00cf020066\n"
".quad 0x2701400200353300,0x0002110000ba0410,0x15005a2e2011001a,0x1507470500393300\n"
".quad 0x060030314700ce90,0x0538030015030748,0x030015d000000246,0x02c00000025a0333\n"
".quad 0x12009102c028072d,0x000002f06507ed31,0x0001750200150310,0x7502001520160096\n"
".quad 0x15039a2303102201,0x0003306602b70200,0x0029000348035000,0x2937110068034029\n"
".quad 0x29d0000003605500,0x0011010015641200,0x6502f53711010e06,0x001503f0000003e0\n"
".quad 0x1500110000153811,0x003f386421012340,0x6231110015040029,0x57043027015e0001\n"
".quad 0x002a201c01620004,0x0000044800bc0e10,0x0300bd0f10002ae0,0x6625310001000015\n"
".quad 0x0000045055005334,0x601b006702001470,0x0010010090000014,0x2938312100660514\n"
".quad 0x04a0000004806600,0x14901b0014000402,0x2600100000e20000,0x6602300200a30520\n"
".quad 0x00a405f0000004b0,0x020016c013023102,0xf01800ce3811007d,0x00ce333272310014\n"
".quad 0x6934110015051028,0x019f101462030000,0x28501000ba326621,0xe310100577071700\n"
".quad 0x7706900000055700,0x004e0000e4111005,0x64082a7011002a04,0x0066e0000005c002\n"
".quad 0x7ae0022906fe7011,0x000657008f641300,0x00b900058b085000,0x1100900200152013\n"
".quad 0x0090069028007b35,0x0006a06501733212,0x213212001506b000,0x3707200000065601\n"
".quad 0x2500120000160201,0x0032313000690750,0x4000000730570137,0x001200058f020016\n"
".quad 0x0001a30200427017,0x0a0200166017003e,0x0042c01700120002,0x00ae077027071a02\n"
".quad 0x002600002a366621,0x001437110014d015,0x020014e015001000,0x3311003ce0190068\n"
".quad 0x08702600100000fe,0x0800270028000565,0x2801b635722101b6,0x1302330403690810\n"
".quad 0x301802330a001620,0x146018023303002c,0x0029801902330400,0x1b0010c010021f02\n"
".quad 0x520300150900290a,0x0a007c0a00140400,0x00160a00a804003f,0x002d00690a04a303\n"
".quad 0x084c058603001405,0x0011000571c00000,0x059b060028099021,0x2500158015001500\n"
".quad 0x0008d033003f3931,0x1900153032210015,0x062b3964210068d0,0x6909700000094065\n"
".quad 0x0950560655391100,0x3031210015600000,0x7fa0160012000016,0x7029004130312100\n"
".quad 0x130b5c3112005609,0x0000160115002c80,0x09b02800bd0000c1,0x42a013005800002c\n"
".quad 0x1306b000002c0600,0x73003bd10a0016b0,0x0cf9601a09660a00,0x020cf97019063302\n"
".quad 0x0250260b9400017b,0x140609a30313017b,0x14a0052915a80100,0x077c070031644300\n"
".quad 0x7908130066025027,0x2ba3670b00150009,0x0b52201a00a3000a,0x020963301900a302\n"
".quad 0x00a3040b3d0400a3,0x007933110b3c071d,0x0a0453010f3f601c,0x0f00030453030eeb\n"
".quad 0x1709221113006604,0x1707f61213010960,0x0007f51313010960,0x0134003c58060014\n"
".quad 0x26012c0413061105,0x0acd0213001400b0,0x7a04130360012026,0x1b0046020014000a\n"
".quad 0x000000206a004602,0x0df0401300320030,0x2f0001031711bf0b,0x28031d0710ab0202\n"
".quad 0x001f00040f970f00,0x41080ead0e000404,0x0038011b0e0f0f00,0x00480e0007030327\n"
".quad 0x18004a031a0c5b0f,0x810f005300000802,0x0038020027040c0a,0x3100010313000106\n"
".quad 0x020e81000e0e0e0e,0x000a101003030e02,0x33001b0e12001b05,0x00040000440e020e\n"
".quad 0x030003980f000102,0x1a0300030f0f0024,0x36024d0202023b00,0x0e1ff402140e0e0e\n"
".quad 0x0b1308250111010e,0x081b061001110803,0x084087012e020000,0x13490b3b0b3a0803\n"
".quad 0x0a40011201110c3f,0xc000150005030000,0x160400000b330a02,0x0460031d03134900\n"
".quad 0xe0001c0b0b080301,0x1c08030028060000,0x003b010b0700000d,0x170039340800004b\n"
".quad 0x0a00000660001109,0x00780b160046003b,0x00780c1000780518,0x2200050d34003b04\n"
".quad 0x00050e3400110200,0x080300240f90000b,0x92040b0f0b0b0b3e,0x0d13110000055700\n"
".quad 0x0d00004e0d122100,0x00000b320b38b000,0xa4003f1349000f13,0x0000001349002614\n"
".quad 0x010804f90bf322ea,0x203a65666e65676c,0x00312e3520474445,0x710f00010864e704\n"
".quad 0x0b03abd2021b1165,0x01054db24b023300,0x0100f2a373040001,0x22b94b020070039c\n"
".quad 0x00019678910b0000,0x02007303060023a0,0xc8b493718d213a4b,0xb6c304000202abc9\n"
".quad 0x25654d7272453003,0x02040050000d0500,0x63755330000e06bd,0x001c0600003502e3\n"
".quad 0x64696c61766e49e6,0x18010065756c6156,0x79726f6d654d7000,0x2816ed6110009c41\n"
".quad 0x6974697100340200,0x1301001a7a696c61,0x02f6001f03002600,0x6e55747261647543\n"
".quad 0x00676e6964616f6c,0x6f725003fe001b04,0x73694472656c6966,0x1c050064656c6261\n"
".quad 0x4e0061746f4e3500,0x01ff002206006465,0x5379646165726c41,0x0700646574726174\n"
".quad 0x6570706f7d080022,0x6f438200f0080064,0x00dc72756769666e,0x746950520020091d\n"
".quad 0x001d0c1d01156863,0x006c6f626d79538d,0x74736f48dd00190d,0x007265746e696f50\n"
".quad 0x200404e002001e10,0x65549f0020111d00,0x1a12006572757478,0x01516e6942310500\n"
".quad 0x684304fd0021131d,0x7365446c656e6e61,0x00726f7470697263,0x636d654d92002414\n"
".quad 0x2618be7269447970,0xea64412100221500,0x26ab8143664f3361,0x80007d03001d1600\n"
".quad 0x6961466863746546,0x4ea6001e171d01b6,0x00646e756f42746f,0x636e795389001b18\n"
".quad 0x191d02286e6f7268,0x65746c6946a00098,0x1d00dc7474655372,0x6d726f4e4400201a\n"
".quad 0x4d43001e1b16001e,0x7845510153657869,0x201c1600d4756365,0x746559746f4ee000\n"
".quad 0x6e656d656c706d49,0xe00102de1f1c0228,0x62fe4c6f6f544001,0x7573b1007a200028\n"
".quad 0x746e656963696666,0x009823002d687744,0x6563616672755396,0x707544d0001a2500\n"
".quad 0x615665746163696c,0x6d614e6f02e16972,0x37040000212b0065,0x00202c1f00200001\n"
".quad 0x17002001005d0300,0x738000f00100202d,0x626c696176616e55,0x636000992e002800\n"
".quad 0x0245217461706d6f,0xac1b6e6f433000b7,0x73735a0138310028,0x4334180301676e69\n"
".quad 0x75614c726f697203,0x351602910001ee6e,0x4d04fc001902001e,0x4568747065447861\n"
".quad 0x0064656465656378,0x6c694690002300c1,0x02bd65706f635365,0xeb00080020c2002f\n"
".quad 0x0a023800c3003a00,0x5021003fc41d005f,0x6e756f4355030b65,0x01a500c52d002774\n"
".quad 0xe80108ac00013102,0x1d02022100e22800,0x27050037e4002f00,0x24046e01001ae500\n"
".quad 0x0035ff1e01147075,0x496c656e72654bec,0xb501c8006567616d,0x1804d6696e553701\n"
".quad 0x75427061e3056bc9,0x656a624f72656666,0x0022cd1703647463,0x170200246d6e553f\n"
".quad 0x61727241910024ce,0x1804f2614d734979,0x1a03050d02001acf,0xb407010301d02800\n"
".quad 0xd1002e00aa461400,0x6975716341a9003d,0x13003fd200646572,0x010016d31f005574\n"
".quad 0xd4002f0094734121,0xd51705090403001d,0x6e5543434580001f,0x02de0104a4726f63\n"
".quad 0x707573500106d619,0x6d694c7d05f46f70,0xae030165d7007469,0x006573556e497700\n"
".quad 0x7265655051001fd8,0x003d00460706fd41,0x7874505e01c501d9,0x617247840017da00\n"
".quad 0x1703517363696870,0x696c764e6a0023db,0x0020dc1700b86b6e,0x706d6f4374694a73\n"
".quad 0xdd1e0527461106dc,0x6372756f53870063,0x0325000769ac0065,0x200019ad17003305\n"
".quad 0x02020c02b17c6853,0x27ae1f002705068e,0x0237030754000400,0x65704f300023af17\n"
".quad 0x737953676ebd0404,0x009902b0006d6574,0x4887009b73655231,0xeb9000656c646e61\n"
".quad 0x6167656c6c49e707,0x910065746174536c,0x03f42900950b0019,0x003d01a552100203\n"
".quad 0x3c068903004904d8,0x754f5403f105bc00,0x00733d008a664f74,0x656d69549d0021bd\n"
".quad 0x08001abe0074756f,0x064b000578020517,0x410301021e05bf2f,0x643f01c66e452002\n"
".quad 0x6f4e34020025c000,0x520021c117002174,0x120c5a6e4f746553,0x001fc400270c4c50\n"
".quad 0x6544734990023b03,0x0704b1796f727473,0x726573734187001f,0x6f54700013c60074\n"
".quad 0x370079796e614d6f,0x0868000019c70073,0x527000b00109a104,0x03a1657473696765\n"
".quad 0x017501020028c81f,0xc10024c918002405,0x5365726177647261,0x003d00126b636174\n"
".quad 0x736e4962019905ca,0x05cb290545757274,0x64652466e701064e,0x03024105cc2d01bb\n"
".quad 0x1e072f705320001b,0xce0063504d0020cd,0x43cf1905690401b6,0x65762202af6f1201\n"
".quad 0x05d02907bd050025,0x696d726550600261,0x001906a02903e374,0xae06a12803fc5315\n"
".quad 0x1d02990502fc0002,0x4db8074e02001ba2,0x00686374616d7369,0x08025c00009406a3\n"
".quad 0x1805486e4f23005b,0x61657274800048a4,0x840807f77061436d,0x0304002507842f04\n"
".quad 0x851f00256110011d,0x00fe654d20050025,0xb10106006907862f,0x0023870064654f00\n"
".quad 0x64656e696f6a8f07,0x7349420500228800,0x0023891f08276c6f,0x74694708b86d1006\n"
".quad 0x8f001c0300228a00,0x8b00746e65764564,0x6e6f72577005003c,0x8c002702c8685467\n"
".quad 0x148d1703d1040025,0x2009ef0005870100,0x2802240401197055,0x6f6e6b770626078e\n"
".quad 0x41330014e7006e77,0x736142af002b6970,0x13020000ce900065,0x50022f00160e07b9\n"
".quad 0x0d9dc5502d080d9d,0x0022cb5002006376,0x0f08b92e021f0d9d,0x0070550229000017\n"
".quad 0x9376110070020025,0x11005ed65502310c,0x7474610390005eb2,0x72123d8155020072\n"
".quad 0x7003020195e4b390,0x1522735502005073,0x02002e0015b41000,0x5a022a000e06b972\n"
".quad 0x4c040074011500e4,0x02002f00755a1800,0x8d5f02003a48b946,0x756e32008d021500\n"
".quad 0xd65f02004700356d,0x00324db803100163,0x0014b51201775f02,0x7a695372002a6210\n"
".quad 0xb21201095f020065,0x696d616e79a3011e,0x02001e6d656d5363,0x2f0f00dbb61e0f4b\n"
".quad 0x02292400370f1912,0x00ed03002f00ed64,0x641f0000ed641f01,0x1f00ed64120500ed\n"
".quad 0x00ed64180300edb3,0x29a9180084660323,0x001b0f0812dd0f02,0x01b6123d32014808\n"
".quad 0x010043c300c90a15,0x600309000022dc32,0x0014410325b6bc01,0x0014421500146815\n"
".quad 0x6c20001401209101,0x1700167311002a64,0x1700164215001678,0x040016431500167c\n"
".quad 0x13156107072402c2,0x0100690852009110,0x5700f0b110012334,0xa102203501006a08\n"
".quad 0x3901004172747009,0x740900d0002f22dc,0x0013783b0100706d,0x1704802100518a00\n"
".quad 0x00010008b0004100,0x00403c01006b0950,0x6f760a7f00730d12,0x0205ba750b006469\n"
".quad 0x123d01a5035f000f,0x6968740c710202e4,0xa5b11101dce20073,0x16003078760d3001\n"
".quad 0x1279760d3600b3a9,0x127a160012b21200,0x880b002d01fa0100,0x1f00820d000f09ba\n"
".quad 0x007c0e1707008203,0x00700e1800760e17,0x044f000b07146508,0x5f330304ce1378f1\n"
".quad 0xb1900540001e785f,0x6f6c660fa60043cc,0x051b100404007461,0x2b02040063051800\n"
".quad 0x3200110000140606,0x035a00063478614d,0x002001002b032501,0x586d69446f001601\n"
".quad 0x00593f04001a0200,0x04005a3b04001a03,0x4d6469724741001a,0x593f030019051f00\n"
".quad 0x005a3b0300190600,0x99020b9702001907,0x0025081800be0509,0x118d6c61746f5454\n"
".quad 0x0021090028002902,0x1b03ce7072615741,0x002b12b901005c0a,0x2d0609fc0400160b\n"
".quad 0x13431000220c1801,0x170d180bc4521000,0x09e1411010a50300,0x001e0e0028114400\n"
".quad 0x7265764f757047c8,0x0100180f0070616c,0xa20104e7501404e7,0x740200211000280f\n"
".quad 0x07e363657845440e,0x746e4961001f1118,0x18121808f7726765,0x70614d6e61436600\n"
".quad 0xcb001e1300290ad3,0x6f4d657475706d6f,0xc80300fe14006564,0x7464695744319f00\n"
".quad 0x321303001f150068,0x8f05001f161f001f,0x1700746867696548,0x1f003f3313030020\n"
".quad 0x1f003f0305001f18,0x2f10f30105002019,0x614c3005007e1a00,0x1b1f0065020bbc79\n"
".quad 0x1c1f006c030c0026,0x73380007010c0027,0x06127a0300271d00,0x6e8201671e1a01f5\n"
".quad 0xc6746e6572727563,0x34001f1f00733801,0x1820180d1a636345,0x737542696350ab00\n"
".quad 0xe202001621006449,0x0019220064494803,0x1703137763635439,0x1802ad0601f80200\n"
".quad 0x626f6c4762001d24,0x7375423200236c61,0x4c71002225180128,0x0337656863614332\n"
".quad 0x02bf0a06044c261f,0x7341b20029270028,0x6e69676e45636e79,0x55001e281802d865\n"
".quad 0x000d096966696e55,0x09050289291f0e64,0x030c00262a1f01cd,0x620501cd2b1f01a6\n"
".quad 0x004c726568746147,0x0218030b00252d1f,0x6c415f0a02c32e1f,0x23000b02c62f0074\n"
".quad 0x22000a02c9301f00,0x6d6f9f01ed311c00,0x66320064496e6961,0x02760604b9010004\n"
".quad 0x75437203005e331f,0x1f00ed70616d6562,0x1f0164090a002434,0x1b01690311002b35\n"
".quad 0x3f04031203002c36,0x3f0403001f371f04,0x7c0305001f381f04,0x3f04030020391f01\n"
".quad 0x3f0305001f3a1f04,0x3f020500203b1f00,0x0e090500bc3c1f04,0x09030c00263d1f01\n"
".quad 0x4d090500cb3e1f01,0xb9030c00263f1f00,0x030c002800c02f00,0x09040028c11f0075\n"
".quad 0x090b0025c21f01fa,0x0312002cc31f00a0,0x06038d00c42f007e,0x2f036661656e6943\n"
".quad 0x00260805036700c5,0x0119030c0026c61f,0x07d2020c0027c71f,0x694d41060026c81f\n"
".quad 0x1f00ef0215536d70,0x2f0079030f0029c9,0x4302ff00074100ca,0x74696c6962617061\n"
".quad 0xcb00726f6a614d79,0x726f6e696f0c0025,0x009d0b060136cc00,0x011007020029cd19\n"
".quad 0x7365697469561871,0x2205db00ce2e119d,0x1900250505bf314c,0x24636f4c3f0025cf\n"
".quad 0x0b097c00d02f0000,0x3900d1003f0d940a,0x2cd21b002c0b0809,0x29005967616e3400\n"
".quad 0x3b734921001cd300,0x616f42757047aa00,0x1c08003ad4006472,0x4970756f72479900\n"
".quad 0x08f5000023d50044,0x6f74413012f74e11,0x28d61900e0050dcc,0x6c676e695304f000\n"
".quad 0x6c62756f446f5465,0x7de4696365725065,0x11c2526672655050,0x615040002fd70029\n"
".quad 0x0200b50207ad6567,0x07f000d8003f14dc,0x1f00260300e10303,0x65725060010217d9\n"
".quad 0x00a106122a706d65,0x73553709c400da2b,0x63726f46341d2765,0x30db0029013d0101\n"
".quad 0x767265736552cf00,0x0019dc0032396465,0x030019dd00333f03,0x3c0b00a4de00343b\n"
".quad 0x01050020df002f14,0x002b03089a020192,0x4f7c100bc300e02f,0x00f1e1006e697470\n"
".quad 0x65526873756c46e0,0x4f69725765746f6d,0x01040401dde21d12,0x0e01aae31f013806\n"
".quad 0x40014a7365735541,0x7339001d54656761,0x061deb020035e400,0x7246400037020299\n"
".quad 0x0b1add0400376d6f,0xea1f0002e90f0cfc,0x900101020f020161,0x002146040000ef00\n"
".quad 0x0f00745f71803300,0x6c2094161d736e75,0x6011080700676e6f,0x80053852135d0713\n"
".quad 0x20099b00ca4d1205,0x6110a30000157942,0xca14120100058705,0x0108058d5100190b\n"
".quad 0x926e00190b7cc112,0x010e1b6d12011005,0x1201180599c20fdd,0x12736765526d756e\n"
".quad 0x7012011c059e9200,0x001502810c567874,0x696212012005a5b9,0xac6000187972616e\n"
".quad 0x00045f6312012405,0x3100164143220c65,0x66441300722805b2,0xb000240100c80b11\n"
".quad 0x65727012012c05b9,0x6d6881001d726566,0x0cfb767261436d65,0x802696c210002101\n"
".quad 0x11040500746e690f,0xa0030c0004f40fa2,0x000022a900781201,0x0c7912010001a203\n"
".quad 0x000c7a1201044400,0x003f000163010827,0x000022bf13040790,0x57130c400fd9130c\n"
".quad 0x100012d110000621,0x20006e1310001214,0xe7140c800f71130c,0x3100117a13000022\n"
".quad 0xf000310f92045000,0xf700060032000704,0x942b84040010040f,0xd79390e2b5900007\n"
".quad 0x0031065025001104,0x6025001004141701,0x2a0018b310001808,0x1b00490512130860\n"
".quad 0x1021115b04000100,0x1900050031000705,0x0574008905102a10,0x000f0195ccb89000\n"
".quad 0xb244001702263401,0x088025000f0195cc,0x800195ccb35b0017,0x01001b0045030085\n"
".quad 0x0026000704b02100,0x0f0410df00006c30,0xe43100db05c02600,0x250001002ca4abc8\n"
".quad 0x95e4b444002f0820,0x5b0017a016000f01,0x040084a00195e4b5,0x25240001001b0046\n"
".quad 0xe13f248a00112801,0x0e4f4304169a000d,0x3f16450c039d0000,0xc53c0913ad0010fe\n"
".quad 0x000f223f12110012,0x282a5a002a1b14c4,0xd53b11f900133538,0x000ffd3f1665000e\n"
".quad 0x430012433c251526,0x082f043000010013,0x11042227ab012949,0x182117001804000c\n"
".quad 0x00181013000c0000,0x13000c0000182017,0x0000181f17001820,0x1e5300180013000c\n"
".quad 0x1e26001814000000,0x1700300b17001800,0x260018091700180b,0x0018071700180009\n"
".quad 0x1805170018000726,0x1803170018051700,0x00180002e6031300,0x0018011400180117\n"
".quad 0x08083121009c1217,0x2603000c02b02f00,0x0c00d026000c0150,0x6000000c02102600\n"
".quad 0xb329d5080a04312b,0x0024190300240160,0x00205502cc0c1704,0x000439001011f000\n"
".quad 0x101800033900101c,0x1021f0258a450100,0x0100100703a10100,0x1b030021f0610001\n"
".quad 0x31037e10130080ff,0x0a20202c8b081c04,0x0008060397020098,0x001401302600a802\n"
".quad 0x0401d43913342c04,0x088d4f0230262b59,0x1000070250210018,0x002cba1f1767db00\n"
".quad 0x240500a921120040,0x2600180800400401,0x13027800004000f0,0x170028021300503b\n"
".quad 0x138b362017001022,0x001001bc26007002,0x0000109d1701bc00,0x480000107e170250\n"
".quad 0x02400000105f1702,0x1702380000104017,0x0217023000001021,0x03b435000b000010\n"
".quad 0x7e2a00b000001000,0x00b003372a00b003,0xba1b00b00002f439,0x221b00b06b1b00b0\n"
".quad 0x2a00b000f42a00b0,0xb0006b2a00b000ba,0x5b3900b0001d2a00,0x0000090100a00013\n"
".quad 0x00010300180c8ef0,0xd10400c00012f439,0x0001040018ec1b16,0x17830400e012722a\n"
".quad 0x2a00010400186a1b,0x1b18000401001225,0xd52a06c50400181d,0xcd1b188004001811\n"
".quad 0x18441b0290040018,0x00183c1b06250400,0x8000107b39000104,0x0000004c01450101\n"
".quad 0x8e2a000103001873,0x861b00f00401a00f,0x0f012a0001040018,0x0ef93919520401c0\n"
".quad 0x8d2a000104001800,0x004c00090101e00e,0x0001030018850000,0x01b00402000e1d2a\n"
".quad 0x2a0001040018151b,0x08d1500802200080,0x01722a0001040018,0x5b1b00106a1b0138\n"
".quad 0x441b0010531b0010,0x2c1b00103c1b0010,0x152a0010241b0010,0x2a00100d1b001001\n"
".quad 0x1000e62a001000ee,0xcf2a001000d72a00,0x001000c02a001000,0x00a92a001000b82a\n"
".quad 0x2a001000a12a0010,0x10008a2a00100092,0x612a001000692a00,0x001000512a001000\n"
".quad 0x00392a001000492a,0x1b001000312a0010,0x001000182a001020,0x0010001b0010081b\n"
".quad 0x982a04d0003d6039,0x03f831782a041837,0x26702a03d82bf82a,0x2a033820f02a0050\n"
".quad 0xf815f02a03181b70,0xe82a02d810702a02,0x029805602e02b80a,0xf38950baff00010f\n"
".quad 0x09260095b4ffffff,0x01017810000fe280,0x0007ffe0a295c1f8,0x0001007919003fde\n"
".quad 0x0c00321e52029800,0x03f0607040020772,0x010a940947220020,0x795c003fde5104ca\n"
".quad 0x3fde623396023ab5,0x0f23056007720200,0x1006000639001000,0x3700100500053900\n"
".quad 0x7210410010040004,0x2000900300a00100,0x1b003008001d7202,0x10ff000939001008\n"
".quad 0x0300100503840300,0x20001b002007d4be,0x087a220010031800,0x4909120010050041\n"
".quad 0x001f081100a00503,0x20001007f1e0ff42,0x7fe4ff3f001f0309,0x100c00f00e0200d0\n"
".quad 0x0b00400f01300c01,0x0080ff080872105a,0x00a0020080ff0929,0x0f00100919012008\n"
".quad 0x001f7385220b0020,0x33de000010eb047f,0x5f1900d00f0300b0,0x09018003e7000078\n"
".quad 0x472400100500ee01,0xb00c026002056479,0x02b0e01000200002,0x000101300383ff3f\n"
".quad 0x50210240081702d0,0x03140b2a000d0779,0xc0660040f0140060,0x0f0001007918000f\n"
".quad 0xffff03800f500010,0x400c0f01500f80ff,0x0d02300f00f00a02,0x1a00400f0201a00f\n"
".quad 0xb00c01c007dbc603,0x60051f0060041b00,0x04782b02000c0903,0x00a00a00900c0220\n"
".quad 0x7102100f2902200f,0xf00a01700c02000e,0x0e00300e00100e00,0xa00f0c01a00f0020\n"
".quad 0x0b39ca01800f1d03,0x0a000a3f00f00b00,0x0c6d03b00f0a0530,0x00d00c00b00c0090\n"
".quad 0x1b07500c2d00700f,0x1b01400d1d00700c,0x1b00c0051b009004,0x001b00700c068006\n"
".quad 0x0d1800a00c1b00a0,0x0c02c003782b0010,0x02c00300043f0470,0x0c07a00c02800e78\n"
".quad 0x02000c01e00c01c0,0xf00c01d00c01b00c,0x2f02200f00d00a01,0x0f00f00c1d02900f\n"
".quad 0xa00fe236032d0080,0x02b00f09cf021602,0x0c02b00f02700c37,0x017a0231ec07000f\n"
".quad 0x03100d00900418a0,0x0020044ffc027821,0x3402120b027b8230,0x022a00d0321e2100\n"
".quad 0xf00a01800a01f000,0x5f01900900200000,0x0900600168000278,0x031f01d0094c0201\n"
".quad 0x50a70278221902a0,0x0060081b0600600f,0x0c03e00f0060071b,0x0060017800007858\n"
".quad 0x0060081400010012,0x7c1f00300902de01,0x0030000a2c0a0030,0x034e010a0030801f\n"
".quad 0x0a0180041b02d009,0x500c0130021b0300,0x04300d0c00f00f0a,0x6e7919210b04600f\n"
".quad 0x0b00b02600003400,0x10034b457a1301b0,0x0090050572244000,0x9e010f008e02ff33\n"
".quad 0x0a004003e7c50000,0x06050672105901e0,0x00800502000a0ea0,0x011e01050080251f\n"
".quad 0x0300800c19008009,0xc00a004021140020,0x0080050030101402,0x060c72245700600a\n"
".quad 0x194000400c005009,0x140c4f15c8ff0d78,0x0530000100200001,0x0c00200e1f04b00a\n"
".quad 0x400800200e000c3f,0x0d3302f00c0d7819,0x00100c1100100102,0x02023900d006ff24\n"
".quad 0xd00d0303390ed00c,0x300f1b03300e1b0e,0x30021b1000021b03,0x1a003002000c3f00\n"
".quad 0x02c00e1b00300d1d,0x06100c0a00400f1f,0x1b00500e1b01400c,0x0f000b0f1101500f\n"
".quad 0x00200d7229060520,0x1b0090101b01800e,0x0c1b01c00e009011,0x11170040101b0040\n"
".quad 0x600d00720c400010,0x1c13121012702202,0x100070f070491a49,0x2612300070e13400\n"
".quad 0x1706700006b002c0,0x250f78103006c000,0xff110e3902000750,0x01400e01100a0200\n"
".quad 0x1800200f0000b00f,0x0031001f10798031,0x02722457037010e9,0x1902a00003400a0d\n"
".quad 0x193101f00a037005,0x0114023f03500378,0x0c0020111f140210,0x310800201100023f\n"
".quad 0x5003150651037819,0x082a035002022b03,0x0150030729035002,0x1200200f0600a00f\n"
".quad 0x013005006f798021,0x1100111002722042,0xe002722131002040,0x5903029000002f03\n"
".quad 0x0e0150010d027810,0x0e1f00800f1b03a0,0x101f0220001800a0,0x0a02f00f1b0a0230\n"
".quad 0xfc802302a0000260,0x21fc0002b0030910,0x060b72245902b004,0x02200b0b2a02200b\n"
".quad 0x022008782106900b,0x06f0091b02200b15,0x200c070007230803,0x400a0a0720091f00\n"
".quad 0x0120070878194007,0x2007072a02200815,0x3902200611c00202,0x004e010220080404\n"
".quad 0x080070041f078009,0x40e90c220514400f,0x003f3b0a00ab1914,0x0e0d00200f04004d\n"
".quad 0xd00c3d10000f0aa0,0x0d07e00a01600c00,0x00100c0a700d07d0,0x0e0300500c00300c\n"
".quad 0x0c00900c095007ef,0x0c0a00b0051f0180,0x70031b00700e00d0,0xc0071b00c0001b00\n"
".quad 0x85221800e0081f00,0x5902000613001f73,0x3902700407047810,0x00500f0270ff0805\n"
".quad 0x7003171a00700f06,0x0013500070081f00,0x02b00802a0000070,0x0119800f0d0d600f\n"
".quad 0x738750381980e01f,0xa008122015150100,0x1014142a00100000,0x122a001010132a00\n"
".quad 0x001008112a00100c,0x000226001004102a,0x02900c02700e0010,0x0f0001700f02800c\n"
".quad 0x000f00500e180180,0x20051b0ca00e0d03,0xb0041b0660001b08,0x0070121b00700a07\n"
".quad 0x0680041b0070131b,0x1b00c0111b0e500e,0x130010131800c012,0x1519001005024f78\n"
".quad 0xc012028043150010,0x79832002900b0220,0x321e9801b0030609,0x1004000110798300\n"
".quad 0x0010052bfc111200,0x4800100c00011248,0x1448001010000113,0x0115460010140001\n"
".quad 0x1702d00000101800,0x0f6d10700f071020,0x3d0380f01f111d00,0x4004142f0330081a\n"
".quad 0x900c0d03000f0903,0x1c02200f11000b02,0x600f00100a00600e,0x70141b02700e2d12\n"
".quad 0x0230000270151902,0x08000e580230101f,0x047221400b01100f,0x002001002e5cb0ff\n"
".quad 0x8103000f0d00100f,0x2860011b2d00010f,0x0001004022000104,0x010800300004cd3d\n"
".quad 0x0d110400400b1f00,0x0005c900005f0595,0x0c261c13130a0040,0x00201215d8110001\n"
".quad 0x23002400135f0400,0x13272807094c001d,0x081100010c5ecd9b,0x003d8000005f0c25\n"
".quad 0x040040a81f0a0080,0x57003f00074b8821,0x0040c41f0b004002,0x031f2b004ddf2304\n"
".quad 0x040040f01f0b0040,0x004a693f2859b413,0x0c5f6402130a0040,0x4f00079c1d210001\n"
".quad 0x1f0a00400011bf00,0x07addc210400401a,0x0b004001ee003f00,0xafca21040040321f\n"
".quad 0x0c004011002f0007,0xb0db23040040401f,0x0a004000222f0080,0xd3c9210400405d1f\n"
".quad 0x00400192003f0007,0x5b21040040851f0b,0x004029002f0007d5,0x0100702bfc02030c\n"
".quad 0x00100007d6842100,0x4f0b032400170e27,0x1304004043002f03,0x0040000c262af998\n"
".quad 0x40731f000040241f,0x1f000040a41f0400,0x0040af1f00004025,0x261f000040b01f04\n"
".quad 0x040040e31f000040,0x40271f000040bc1f,0x030040012e2f0000,0x40281f000040c81f\n"
".quad 0x1f040040a21f0000,0x0040291f000040d4,0x1303004002032f00,0x0040008c260040e0\n"
".quad 0x40b81f0000402a1f,0x270007d86c210400,0x00402b1f00400800,0x741f040040e51f00\n"
".quad 0x0000402c1f000040,0x7c2e03004003412f,0x0000402d1f0100d8,0x80881f0400408c1f\n"
".quad 0x2e0000402e1f0000,0x34000100001402fd,0x00c009094600d890,0x1bc003172305c00c\n"
".quad 0x020040c01300010b,0x00180400400e099c,0xd01f03008003552f,0x0000802d1f000080\n"
".quad 0xd9002e0400806b1f,0x1f0000802d1f0080,0x00401013040040b4,0x40051f004000a026\n"
".quad 0x13040040d81f0000,0x1f0040b0170040b0,0x00044c2f00004006,0x341000da60390301\n"
".quad 0x2f01000c07ac0313,0xdca0230300800468,0x0c1f004001160308,0x040040aa1f000080\n"
".quad 0x40041f00c0de802e,0x040c60364d130000,0x80030548df302308,0x3213040380001f00\n"
".quad 0x24000100062b0594,0x6a0080050af600e1,0x0001008018000001,0xe4802e040040571f\n"
".quad 0x1f000040031f0040,0x06e8002104004092,0x0d02000300d90500,0x13050040cf1f0040\n"
".quad 0x00400001250080ea,0x40fa1f000040071f,0x070008eb80230400,0x2f000040091f0080\n"
".quad 0x00ee260400800165,0x00400d0240060039,0xe5f11a050040e21f,0x40141c3060031239\n"
".quad 0x2e03004002a32f00,0x00801f1f00c0fb80,0x002e040040d01f00,0x000040201f01c0fe\n"
".quad 0x012d04008003302f,0x000040211f01c001,0x400313050040801f,0x40121f0040011600\n"
".quad 0x4010041806130000,0xa82a0000010f6d0c,0x0057000826000800,0x3900010c03280517\n"
".quad 0x0038080008002550, 0x088017000100062f, 0x0000000000000000\n"
".text\n");
#ifdef __cplusplus
extern "C" {
#endif
extern const unsigned long long fatbinData[3511];
#ifdef __cplusplus
}
#endif
#ifdef __cplusplus
extern "C" {
#endif
static const __fatBinC_Wrapper_t __fatDeviceText __attribute__ ((aligned (8))) __attribute__ ((section (__CUDAFATBINSECTION)))=
{ 0x466243b1, 1, fatbinData, 0 };
#ifdef __cplusplus
}
#endif
codes/mmElem/matrixMul_dlink.fatbin.c 0 → 100644
View file @ e4ae2422
#ifndef __SKIP_INTERNAL_FATBINARY_HEADERS
#include "fatbinary_section.h"
#endif
#define __CUDAFATBINSECTION ".nvFatBinSegment"
#define __CUDAFATBINDATASECTION ".nv_fatbin"
asm(
".section .nv_fatbin, \"a\"\n"
".align 8\n"
"fatbinData:\n"
".quad 0x00100001ba55ed50,0x0000000000000110,0x0000005001010002,0x00000000000000c0\n"
".quad 0x00000000000000be,0x0000004600010007,0x0000000c00000040,0x0000000000002013\n"
".quad 0x0000000000000000,0x0000000000000268,0x754d78697274616d,0x00000000206f2e6c\n"
".quad 0x010102464c457fa2,0x0002660001000733,0xc0230001006e00be,0xf500010012000801\n"
".quad 0x380040004605460d,0x0100040040000300,0x72747368732e0000,0x2700082e00626174\n"
".quad 0x735f00ff00086d79,0x766e2e0078646e68,0x2100326f666e692e,0x2e00df004800010f\n"
".quad 0x0100402200010003,0x0108003000322e00,0x722f0400400b1f00,0x0174131113004000\n"
".quad 0x000100a82200010e,0x2a00240600061811,0x0000065700180008,0x0500480f01a80500\n"
".quad 0x003801130040a81b,0x2f0038081500010f,0x0008801700010006,0x0000000000000000\n"
".text\n");
#ifdef __cplusplus
extern "C" {
#endif
extern const unsigned long long fatbinData[36];
#ifdef __cplusplus
}
#endif
#ifdef __cplusplus
extern "C" {
#endif
static const __fatBinC_Wrapper_t __fatDeviceText __attribute__ ((aligned (8))) __attribute__ ((section (__CUDAFATBINSECTION)))=
{ 0x466243b1, 2, fatbinData, (void**)__cudaPrelinkedFatbins };
#ifdef __cplusplus
}
#endif
codes/mmElem/matrixMul_dlink.reg.c 0 → 100644
View file @ e4ae2422
#define NUM_PRELINKED_OBJECTS 0
codes/mmElem/matrixmul.conf 0 → 100644
View file @ e4ae2422
[DEFAULT]
debug =True
# Name of the gdb executable
gdbExecName = /usr/local/cuda-10.1.243/bin/cuda-gdb
# Which fault model to use, 0 -> single; 1 -> double;
# 2 -> random; 3 -> zeros; 4 -> least 16 significant bits (LSB);
# 5 -> least 8 significant bits (LSB)
# If you want multiple fault models, place them separated by ','
# faultModel = 0,2,3
faultModel = 0
# Injection site
# Can be:
# RF -> Register File
# INST_OUT -> Instruction Output (NOT IMPLEMENTED YET)
# INST_composed -> Instruction Adress (NOT IMPLEMENTED YET)
injectionSite = RF
# Max time factor to finish the app, this will be multiplied by the application running time
# For example if your app spend 2s, and the maxWaitTimes is 5, the max running time before it is
# Considered as a crash is 10s
maxWaitTimes = 5
# binary file of the application
# Must be full path
benchmarkBinary = /home/badia/carol-fi-carol-fi_cuda-parallel/codes/mmElem/matrixMul
# Commands to set the session inside GDB environment
benchmarkArgs = -device=1 -wA=16384 -hA=16384 -hB=16384 -wB=16384
# CSV output file. It will be overwrite at each injection
csvFile = codes/mmElem/fi_matrix_mul_single_bit.csv
# You should create a script on the benchmark source folder to verify GOLD_OUTPUT x INJ_OUTPUT
goldenCheckScript = codes/mmElem/sdc_check.sh
# Number of signals that will be sent to the application
seqSignals = 20
# Initial sleep time in seconds before start sending signals
# Generally the memory setup time
initSleep = 2.3
codes/mmElem/matrixmul_16K.conf 0 → 100644
View file @ e4ae2422
[DEFAULT]
debug =True
# Name of the gdb executable
gdbExecName = /usr/local/cuda-10.1.243/bin/cuda-gdb
# Which fault model to use, 0 -> single; 1 -> double;
# 2 -> random; 3 -> zeros; 4 -> least 16 significant bits (LSB);
# 5 -> least 8 significant bits (LSB)
# If you want multiple fault models, place them separated by ','
# faultModel = 0,2,3
faultModel = 0
# Injection site
# Can be:
# RF -> Register File
# INST_OUT -> Instruction Output (NOT IMPLEMENTED YET)
# INST_composed -> Instruction Adress (NOT IMPLEMENTED YET)
injectionSite = RF, INST_OUT
# Max time factor to finish the app, this will be multiplied by the application running time
# For example if your app spend 2s, and the maxWaitTimes is 5, the max running time before it is
# Considered as a crash is 10s
maxWaitTimes = 5
# binary file of the application
# Must be full path
benchmarkBinary = /home/badia/mycarol-fi/codes/mmElem/matrixMul
# Commands to set the session inside GDB environment
benchmarkArgs = -wA=16384 -hA=16384 -hB=16384 -wB=16384
# CSV output file. It will be overwrite at each injection
csvFile = codes/matrixMul/fi_matrix_mul_single_bit.csv
# You should create a script on the benchmark source folder to verify GOLD_OUTPUT x INJ_OUTPUT
goldenCheckScript = codes/matrixMul/sdc_check.sh
# Number of signals that will be sent to the application
seqSignals = 20
# Initial sleep time in seconds before start sending signals
# Generally the memory setup time
initSleep = 2.1
#kernel = matrixMulCUDA
kernel = matrixMul.cu:208
kernel_end = matrixMul.cu:216
codes/mmElem/matrixmul_24K.conf 0 → 100644
View file @ e4ae2422
[DEFAULT]
debug =True
# Name of the gdb executable
gdbExecName = /usr/local/cuda-10.1.243/bin/cuda-gdb
# Which fault model to use, 0 -> single; 1 -> double;
# 2 -> random; 3 -> zeros; 4 -> least 16 significant bits (LSB);
# 5 -> least 8 significant bits (LSB)
# If you want multiple fault models, place them separated by ','
# faultModel = 0,2,3
faultModel = 0
# Injection site
# Can be:
# RF -> Register File
# INST_OUT -> Instruction Output (NOT IMPLEMENTED YET)
# INST_composed -> Instruction Adress (NOT IMPLEMENTED YET)
injectionSite = RF
# Max time factor to finish the app, this will be multiplied by the application running time
# For example if your app spend 2s, and the maxWaitTimes is 5, the max running time before it is
# Considered as a crash is 10s
maxWaitTimes = 5
# binary file of the application
# Must be full path
benchmarkBinary = /home/badia/carol-fi-carol-fi_cuda-parallel/codes/matrixMul/matrixMul
# Commands to set the session inside GDB environment
#benchmarkArgs = -device=1 -wA=16384 -hA=16384 -hB=16384 -wB=16384
benchmarkArgs = -device=1 -wA=24576 -hA=24576 -hB=24576 -wB=24576
# CSV output file. It will be overwrite at each injection
csvFile = codes/matrixMul/fi_matrix_mul_single_bit.csv
# You should create a script on the benchmark source folder to verify GOLD_OUTPUT x INJ_OUTPUT
goldenCheckScript = codes/matrixMul/sdc_check.sh
# Number of signals that will be sent to the application
seqSignals = 20
# Initial sleep time in seconds before start sending signals
# Generally the memory setup time
initSleep = 4.9
codes/mmElem/sdc_check.sh 0 → 100644
View file @ e4ae2422
#!/usr/bin/sh
# SDC checking diff
# Must compare all things here
# Any particular output comparison must be made here
# To be considered as an SDC or CRASH the
# DIFF_LOG and DIFF_ERR_LOG files must not be empty
# INJ_OUTPUT_PATH, INJ_ERR_PATH, GOLD_OUTPUT_PATH, GOLD_ERR_PATH
# are environment variables defined by the fault_injector.py
# diff stdout
diff -B ${INJ_OUTPUT_PATH} ${GOLD_OUTPUT_PATH} > ${DIFF_LOG}
# Special comparison like the following one can be done in this script
grep -q "Result = FAIL" ${INJ_OUTPUT_PATH} >> ${DIFF_LOG}
# diff stderr
diff -B ${INJ_ERR_PATH} ${GOLD_ERR_PATH} > ${DIFF_ERR_LOG}
# Must exit 0
exit 0
\ No newline at end of file
common_functions.py 0 → 100644
View file @ e4ae2422
import os
import pickle
import re
import sys
import common_parameters as cp
if sys.version_info >= (3, 0):
import configparser # python 3
else:
import ConfigParser # python 2
"""
Support function to execute a command
and return the output.
If the command contains NEWLINE character
it will result in a list.
"""
def execute_command(gdb, to_execute):
ret = gdb.execute(to_execute, to_string=True)
return ret.splitlines()
"""
Serialize a dictionary into a
file path using pickle.
"""
def save_file(file_path, data):
with open(file_path, "wb") as f_out:
pickle.dump(data, f_out)
f_out.close()
"""
Serialize a dictionary into a
file path using pickle.
"""
def append_file(file_path, data):
with open(file_path, "ab") as f_out:
pickle.dump(data, f_out)
f_out.close()
"""
Load a dictionary from a file path using pickle.
return a dictionary
"""
def load_file(file_path):
with open(file_path, "rb") as f_in:
data = pickle.load(f_in)
return data
"""
Read configuration file
"""
def load_config_file(flip_config_file):
# Read configuration file
if sys.version_info >= (3, 0):
conf = configparser.ConfigParser()
else:
conf = ConfigParser.ConfigParser()
conf.read(flip_config_file)
return conf
"""
Kill all remaining processes
"""
def kill_all(kill_string, logging=None):
for cmd in kill_string.split(";"):
os.system(cmd + " > /dev/null 2>&1")
if logging:
logging.debug("kill cmd: {}".format(cmd))
"""
GDB python cannot find common_functions.py, so I added this directory to PYTHONPATH
"""
def set_python_env():
current_path = os.path.dirname(os.path.realpath(__file__))
os.environ['PYTHONPATH'] = "$PYTHONPATH:" + current_path + ":" + current_path + "/classes"
os.environ['OMP_NUM_THREADS'] = '1'
return current_path
"""
Remove all useless information produced by CUDA-GDB on the output files
before they got to the SDC check script
"""
def remove_useless_information_from_output(output_file_path):
ok_output_lines = []
with open(output_file_path, 'r') as ifp:
lines = ifp.readlines()
for line in lines:
is_line_addable = True
for pattern in cp.POSSIBLE_USELESS_GDB_OUTPUT_PATTERNS:
# It is addable or not
search_result = re.search(pattern=pattern, string=line)
if search_result:
is_line_addable = False
if is_line_addable:
ok_output_lines.append(line)
# Overwrite the output file
with open(output_file_path, 'w') as ofp:
ofp.writelines(ok_output_lines)
"""
Show output function
to allow pretty printing
"""
def printf(*args):
string_to_print = "" # ""\r"
for i in args:
string_to_print += "{0} ".format(i)
print(string_to_print)
common_parameters.py 0 → 100644
View file @ e4ae2422
# Max size of register
SINGLE_MAX_SIZE_REGISTER = 32
# Times to profile
# this will be the max number of executions
# to profiler application
MAX_TIMES_TO_PROFILE = 2
# Log path to store all injections info
LOGS_PATH = 'logs'
# Temporary file to store kernel information
KERNEL_INFO_DIR = LOGS_PATH + '/tmp/carol-fi-kernel-info.txt'
# For golden generation
GOLD_ERR_PATH = LOGS_PATH + '/tmp/carol_fi_golden_bench_err.txt'
GOLD_OUTPUT_PATH = LOGS_PATH + '/tmp/carol_fi_golden_bench_output.txt'
# Files that will be compared to golden ones
INJ_OUTPUT_PATH = LOGS_PATH + '/tmp/carol_fi_inj_bench_output_{}.txt'
INJ_ERR_PATH = LOGS_PATH + '/tmp/carol_fi_inj_bench_err_{}.txt'
# Internal python scripts
FLIP_SCRIPT = 'flip_value.py'
PROFILER_SCRIPT = 'profiler_new.py'
# Temporary difference logs
DIFF_LOG = LOGS_PATH + '/tmp/diff_{}.log'
DIFF_ERR_LOG = LOGS_PATH + '/tmp/diff_err_{}.log'
# Debug env vars
# Debug FI process
DEBUG = True
# Debug profiler process
DEBUG_PROFILER = True
# Log file for SignalApp thread
SIGNAL_APP_LOG = LOGS_PATH + '/tmp/signal_app_thread_{}.txt'
# Num of sleep time divisor
NUM_DIVISION_TIMES = 100.0
# Common body of log filename
LOG_DEFAULT_NAME = LOGS_PATH + '/tmp/carolfi-flipvalue-{}.log'
# MAX INT 32 bits
MAX_INT_32 = 4294967295
# Most of the benchmarks we cannot wait until the end of the processing
# Considering most of 90% of the time
MAX_SIGNAL_BEFORE_ENDING = 0.9
# termination, program, alarm, asynchronous, job, operation error, miscellaneous, signal interruption
# 'SIGINT' must not be here, since I used it to send an interruption to app
SIGNALS = ['SIGKILL', 'SIGTERM', 'SIGQUIT', 'SIGHUP', # termination codes
'SIGFPE', 'SIGILL', 'SIGSEGV', 'SIGBUS', 'SIGABRT', 'SIGIOT', 'SIGTRAP', 'SIGEMT', 'SIGSYS', # program codes
'SIGALRM', 'SIGVTALRM', 'SIGPROF', # alarm codes
'SIGIO', 'SIGURG', 'SIGPOLL', # asynchronous codes
'SIGCHLD', 'SIGCLD', 'SIGCONT', 'SIGSTOP', 'SIGTSTP', 'SIGTTIN', 'SIGTTOU', # job control
'SIGPIPE', 'SIGLOST', 'SIGXCPU', 'SIGXFSZ', # operation codes
'SIGUSR1', 'SIGUSR2', 'SIGWINCH', 'SIGINFO', # miscellaneous codes
'strsignal', 'psignal', # signal messages
# cuda signals
'CUDA_EXCEPTION_0', 'CUDA_EXCEPTION_1', 'CUDA_EXCEPTION_2', 'CUDA_EXCEPTION_3', 'CUDA_EXCEPTION_4',
'CUDA_EXCEPTION_5',
'CUDA_EXCEPTION_6', 'CUDA_EXCEPTION_7', 'CUDA_EXCEPTION_8', 'CUDA_EXCEPTION_9', 'CUDA_EXCEPTION_10',
'CUDA_EXCEPTION_11',
'CUDA_EXCEPTION_12', 'CUDA_EXCEPTION_13', 'CUDA_EXCEPTION_14', 'CUDA_EXCEPTION_15']
# All trash produced by GDB must be add here in this list
# Using the Regular Expression format (python re)
POSSIBLE_USELESS_GDB_OUTPUT_PATTERNS = [
r'.*Thread.*received signal SIGINT, Interrupt.*', # Thread SIGINT message
r'.*New Thread.*', # New GDB Thread creation
r'.*Thread debugging using.*enabled.*', # Lib thread enabled
r'.*Using host.*library.*', # Using host library
r'.*Switching focus to CUDA kernel.*', # Switching focus to CUDA kernel message
r'.*0x.*in.*<<<.*>>>.*', # Kernel interruption message
r'.*Inferior.*\(process.*\) exited normally.*', # GDB exited normally message
r'.*Thread 0x.*exited.*', # Thread exited
r'.*0x.* in cu.* () from /usr/lib/.*libcuda.*', # Cuda lib calls
r'.*0x.*in.*\[clone.*\].*\(\).*', # OMP calls
r'.*0x.*in.*', # General API call
r'.*Inferior.*\(process.*\).*', # General inferior process
]
# Injection sites
RF = 0
INST_OUT = 1
INST_ADD = 2
INJECTION_SITES = {
'RF': RF,
'INST_OUT': INST_OUT,
'INST_ADD': INST_ADD
}
# Which fault model to use, 0 -> single; 1 -> double;
# 2 -> random; 3 -> zeros; 4 -> least 16 significant bits (LSB);
# 5 -> least 8 significant bits (LSB)
FLIP_SINGLE_BIT = 0
FLIP_TWO_BITS = 1
RANDOM_VALUE = 2
ZERO_VALUE = 3
LEAST_16_BITS = 4
LEAST_8_BITS = 5
# Focus error string
FOCUS_ERROR_STRING = "Focus not set on any active CUDA kernel."
fault_injector.py 0 → 100755
View file @ e4ae2422
#!/usr/bin/env python3
import argparse
import os
import random
import re
import shutil
import time
import datetime
import signal
import common_functions as cf # All common functions will be at common_functions module
import common_parameters as cp # All common parameters will be at common_parameters module
import sys
import signal
import threading
from threading import Thread, Lock
from classes.RunGDB import RunGDB
from classes.SummaryFile import SummaryFile
from classes.Logging import Logging
from classes.SignalApp import SignalApp
"""
[THIS FUNCTION CAN BE EDITED IF DESIRED]
User defined function
this function must return an empty or not string.
The string will be appended in the last collum of summary CSV file
the column will have 'user_defined' as header
if the string is always empty the column will be empty, otherwise it
will contain the returned values for each injection
"""
def receiveSignal(signalNumber, frame):
syncro.wait()
print ("Alcanzado el breakpoint, y recibida la señal",signalNumber);
def receiveEnd(signalNumber, frame):
print ("Esperando sincronismo del final");
wait_finish.wait()
print ("Recibida la señal de final del programa",signalNumber);
def user_defined_function(injection_output_path):
# This is a temporary example for carol-fi-codes suite
# it will search for a LOGFILENAME int the benchmark output if it finds
# then the desired pattern will be returned
with open(injection_output_path, "r") as fp:
for l in fp.readlines():
m = re.match(r"LOGFILENAME:.*/(\S+).*", l)
if m:
return m.group(1)
return ""
"""
CTRL + C event
"""
def signal_handler(sig, frame):
global kill_strings, exit_injector
exit_injector = True
for cmd in kill_strings.split(";"):
os.system(cmd + " > /dev/null 2>&1")
# os.system("rm -f {}/bin/*".format(current_path))
# for th in gpus_threads:
# th.join()
# sys.exit(0)
"""
Check if app stops execution (otherwise kill it after a time)
"""
def check_finish(max_wait_time, logging, timestamp_start, process, thread, kill_string):
is_hang = False
# Wait maxWaitTimes the normal duration of the program before killing it
# max_wait_time = int(conf.get(section, "maxWaitTimes")) * end_time
sleep_time = max_wait_time / cp.NUM_DIVISION_TIMES
if cp.DEBUG:
cf.printf("THREAD: {} MAX_WAIT_TIME {} CHECK FINISH SLEEP_TIME {}".format(thread, max_wait_time, sleep_time))
# Watchdog to avoid hangs
p_is_alive = process.is_alive()
now = int(time.time())
diff_time = now - timestamp_start
while diff_time < max_wait_time and p_is_alive:
time.sleep(sleep_time)
p_is_alive = process.is_alive()
now = int(time.time())
diff_time = now - timestamp_start
# Process finished ok
if not p_is_alive:
logging.debug("PROCESS NOT RUNNING")
if cp.DEBUG:
cf.printf("THREAD {} PROCESS NOT RUNNING".format(thread))
# check execution finished before or after waitTime
if diff_time < max_wait_time:
logging.info("Execution on thread {} finished before waitTime. {} seconds.".format(thread, diff_time))
else:
logging.info("Execution on thread {} finished after waitTime. {} seconds.".format(thread, diff_time))
is_hang = True
logging.debug("now: {}".format(now))
logging.debug("timestampStart: {}".format(timestamp_start))
# Kill all the processes to make sure the machine is clean for another test
cf.kill_all(kill_string=kill_string, logging=logging)
# Also kill the subprocess
process.kill_subprocess()
return is_hang
"""
Copy the logs and output(if fault not masked) to a selected folder
"""
def save_output(is_sdc, is_hang, is_crash, is_masked, logging, unique_id, flip_log_file, inj_output_path,
inj_err_path, diff_log_path, diff_err_path, signal_app_log_path, thread):
# FI successful
fi_injected = False
if os.path.isfile(flip_log_file):
with open(flip_log_file, "r") as fp:
content = fp.read()
if re.search('Fault Injection Successful', content):
fi_injected = True
fp.close()
dt = datetime.datetime.fromtimestamp(time.time())
ymd = dt.strftime('%Y_%m_%d')
y_m_d_h_m_s = dt.strftime('%Y_%m_%d_%H_%M_%S')
y_m_d_h_m_s = unique_id + "-" + y_m_d_h_m_s
dir_d_t = os.path.join(ymd, y_m_d_h_m_s)
# Log and create the paths
if not fi_injected:
cp_dir = os.path.join(cp.LOGS_PATH, 'failed-injection', dir_d_t)
logging.summary("Fault Injection Failed")
elif is_hang:
cp_dir = os.path.join(cp.LOGS_PATH, 'hangs', dir_d_t)
logging.summary("Hang")
elif is_crash:
cp_dir = os.path.join(cp.LOGS_PATH, 'crashs', dir_d_t)
logging.summary("Crash")
elif is_sdc:
cp_dir = os.path.join(cp.LOGS_PATH, 'sdcs', dir_d_t)
logging.summary("SDC")
elif is_masked:
cp_dir = os.path.join(cp.LOGS_PATH, 'masked', dir_d_t)
logging.summary("Masked")
elif not os.path.isfile(inj_output_path):
cp_dir = os.path.join(cp.LOGS_PATH, 'no_output_generated', dir_d_t)
logging.summary("no_output_generated")
else:
cp_dir = os.path.join(cp.LOGS_PATH, 'unknown', dir_d_t)
logging.summary("Unknown")
if not os.path.isdir(cp_dir):
os.makedirs(cp_dir)
# Moving all necessary files
for file_to_move in [flip_log_file, inj_output_path,
inj_err_path, diff_log_path, diff_err_path, signal_app_log_path]:
try:
shutil.move(file_to_move, cp_dir)
except Exception as err:
if cp.DEBUG:
cf.printf("THREAD {} ERROR ON MOVING {} -- {}".format(thread, file_to_move, str(err)))
"""
Check output files for SDCs
"""
def check_sdcs_and_app_crash(logging, sdc_check_script, inj_output_path, inj_err_path, diff_log_path, diff_err_path):
is_sdc = False
is_masked = False
is_app_crash = [False]
if not os.path.isfile(inj_output_path):
logging.error("outputFile not found: " + inj_output_path)
is_app_crash = True
elif not os.path.isfile(cp.GOLD_OUTPUT_PATH):
logging.error("gold_file not found: " + cp.GOLD_OUTPUT_PATH)
raise ValueError("GOLD FILE NOT FOUND")
elif not os.path.isfile(sdc_check_script):
logging.error("sdc check script file not found: " + sdc_check_script)
raise ValueError("SDC CHECK SCRIPT NOT FOUND: " + sdc_check_script)
elif not os.path.isfile(inj_err_path):
logging.error("possible crash, stderr not found: " + inj_output_path)
is_app_crash[0] = True
elif not os.path.isfile(cp.GOLD_ERR_PATH):
logging.error("gold_err_file not found: " + cp.GOLD_ERR_PATH)
raise ValueError("GOLD ERR FILE NOT FOUND: " + cp.GOLD_ERR_PATH)
# Removing the output trash info
# It automatically overwrite the file in the output path
cf.remove_useless_information_from_output(output_file_path=inj_output_path)
cf.remove_useless_information_from_output(output_file_path=inj_err_path)
if os.path.isfile(cp.GOLD_OUTPUT_PATH) and os.path.isfile(inj_output_path) and os.path.isfile(
cp.GOLD_ERR_PATH) and os.path.isfile(inj_err_path):
# Set environ variables for sdc_check_script
os.environ['GOLD_OUTPUT_PATH'] = cp.GOLD_OUTPUT_PATH
os.environ['INJ_OUTPUT_PATH'] = inj_output_path
os.environ['GOLD_ERR_PATH'] = cp.GOLD_ERR_PATH
os.environ['INJ_ERR_PATH'] = inj_err_path
os.environ['DIFF_LOG'] = diff_log_path
os.environ['DIFF_ERR_LOG'] = diff_err_path
compare_script_result = os.system("sh " + sdc_check_script)
if compare_script_result != 0:
raise ValueError("SDC/Crash script returned a value different from 0. Cannot proceed")
# Test if files are ok
with open(diff_log_path, 'r') as fi:
out_lines = fi.readlines()
if len(out_lines) != 0:
# Check if NVIDIA signals on output
for line in out_lines:
if 'PASS' in line:
is_masked= True
break
if 'FAIL' in line:
is_sdc= True
break
for carol_fi_signal in cp.SIGNALS:
for line in out_lines:
if carol_fi_signal in line:
is_app_crash[0] = True
if len (is_app_crash) == 1:
is_app_crash.append(carol_fi_signal)
else:
is_app_crash[1]=is_app_crash[1]+" "+carol_fi_signal
break
# if is_app_crash[0]:
# break
#if not is_app_crash:
# is_sdc = True
# with open(diff_err_path, 'r') as fi_err:
# err_lines = fi_err.readlines()
# if len(err_lines) != 0:
# is_app_crash[0] = True
# if len (is_app_crash) == 1:
# is_app_crash.append('Unknown')
return is_sdc, is_app_crash, is_masked
"""
Check the carolfi-xxxx logfile
the status of the injected fault
"""
def check_injection_outcome(host_thread, logging, injection_site):
# Search for set values for register
# Must be done before save output
# Check THREAD FOCUS. Check if could change the block and the thread
register = block = thread = "___"
# Search for block
block_focus = logging.search("CUDA_BLOCK_FOCUS")
if block_focus:
m = re.search(r"CUDA_BLOCK_FOCUS:.*block[ ]+\((\d+),(\d+),(\d+)\).*", block_focus)
if m:
block = "{}_{}_{}".format(m.group(1), m.group(2), m.group(3))
thread_focus = logging.search("CUDA_THREAD_FOCUS")
# Search for thread
if thread_focus:
m = re.search(r"CUDA_THREAD_FOCUS:.*thread[ ]+\((\d+),(\d+),(\d+)\).*", thread_focus)
if m:
thread = "{}_{}_{}".format(m.group(1), m.group(2), m.group(3))
register_selected = logging.search("SELECTED_REGISTER")
# Search for the register
if register_selected:
m = re.search(r"SELECTED_REGISTER:(\S+).*", register_selected)
if m:
register = m.group(1)
# Was the fault injected?
try:
old_value = re.findall(r'old_value:(\S+)', logging.search("old_value"))[0]
new_value = re.findall(r'new_value:(\S+)', logging.search("new_value"))[0]
fi_successful = True
# Check specific outcomes
# No need to process for RF
instruction = 'register'
if cp.INJECTION_SITES[injection_site] in [cp.INST_OUT, cp.INST_ADD]:
# if fault was injected ASSM_LINE MUST be in the logfile
assm_line = logging.search("ASSM_LINE")
instruction = re.match(r".*:\t(\S+) .*", assm_line).group(1)
except TypeError as te:
instruction = new_value = old_value = None
fi_successful = False
if cp.DEBUG:
cf.printf("THREAD {} FAULT WAS NOT INJECTED. ERROR {}".format(host_thread, te))
return block, fi_successful, new_value, old_value, register, thread, instruction
"""
Function to run one execution of the fault injector
return old register value, new register value
"""
def gdb_inject_fault(**kwargs):
global kill_strings
# These are the mandatory parameters
bits_to_flip = kwargs.get('bits_to_flip')
fault_model = kwargs.get('fault_model')
unique_id = kwargs.get('unique_id')
max_time = kwargs.get('max_time')
end_time = kwargs.get('end_time')
current_path_local = kwargs.get('current_path')
# injection site
injection_site = kwargs.get('injection_site')
benchmark_args = kwargs.get('benchmark_args')
benchmark_binary = kwargs.get('benchmark_binary')
host_thread = kwargs.get('host_thread')
seq_signals = kwargs.get('seq_signals')
init_sleep = kwargs.get('init_sleep')
sdc_check_script = kwargs.get('gold_check_script')
# signalCmd
signal_cmd = kwargs.get("signal_cmd")
gdb_exec_name = kwargs.get('gdb_path')
gdb_kernel = kwargs.get('kernel')
# Define all path to current thread execution
# Logging file
flip_log_file = cp.LOG_DEFAULT_NAME.format(unique_id)
inj_output_path = cp.INJ_OUTPUT_PATH.format(unique_id)
inj_err_path = cp.INJ_ERR_PATH.format(unique_id)
signal_app_log = cp.SIGNAL_APP_LOG.format(unique_id)
diff_log_path = cp.DIFF_LOG.format(unique_id)
diff_err_path = cp.DIFF_ERR_LOG.format(unique_id)
# Starting FI process
if cp.DEBUG:
cf.printf("THREAD {} STARTING GDB SCRIPT".format(host_thread))
logging = Logging(log_file=flip_log_file, unique_id=unique_id)
logging.info("Starting GDB script")
# Generate configuration file for specific test
gdb_env_string = "{}|{}|{}|{}|{}|file {}; set args {}|{}".format(gdb_kernel,os.getpid(),",".join(str(i) for i in bits_to_flip), fault_model,
flip_log_file, benchmark_binary, benchmark_args,
injection_site)
if cp.DEBUG:
cf.printf("THREAD {} ENV GENERATE FINISHED".format(host_thread))
# First we have to start the SignalApp thread
signal_app_thread = SignalApp(max_wait_time=end_time, signal_cmd=signal_cmd,
log_path=signal_app_log, unique_id=unique_id,
signals_to_send=seq_signals,
init_sleep=init_sleep,syncro=syncro,waitfinish=wait_finish)
# Create one thread to start gdb script
# Start fault injection process
fi_process = RunGDB(unique_id=unique_id, gdb_exec_name=gdb_exec_name, flip_script=cp.FLIP_SCRIPT,
carol_fi_base_path=current_path_local, gdb_env_string=gdb_env_string,
gpu_to_execute=host_thread,
inj_output_path=inj_output_path, inj_err_path=inj_err_path)
if cp.DEBUG:
cf.printf("THREAD {} STARTING PROCESS".format(host_thread))
# Starting both threads
fi_process.start()
print ("Esperando breakpoint.....")
#syncro.wait()
#syncro.reset()
"syncro = threading.Barrier(2, timeout=5) "
signal_app_thread.start()
if cp.DEBUG:
cf.printf("THREAD {} PROCESSES SPAWNED".format(host_thread))
# Start counting time
timestamp_start = int(time.time())
# Check if app stops execution (otherwise kill it after a time)
# max_wait_time, logging, timestamp_start, thread, kill_string
is_hang = check_finish(max_wait_time=max_time, logging=logging, timestamp_start=timestamp_start,
process=fi_process, thread=host_thread,
kill_string=kill_strings,init_hang=signal_app_thread.ishang)
if cp.DEBUG:
cf.printf("THREAD {} FINISH CHECK OK".format(host_thread))
# finishing and removing thrash
fi_process.join()
# fi_process.terminate()
signal_app_thread.join()
# Get the signal init wait time before destroy the thread
signal_init_wait_time = signal_app_thread.get_int_wait_time()
del fi_process, signal_app_thread
if cp.DEBUG:
cf.printf("THREAD {} PROCESS JOINED".format(host_thread))
# Change the behavior of this function if any other information
# needs to be added in the final summary
user_defined_string = user_defined_function(injection_output_path=inj_output_path)
# # Check output files for SDCs
is_sdc, is_crash, is_masked = check_sdcs_and_app_crash(logging=logging, sdc_check_script=sdc_check_script,
inj_output_path=inj_output_path, inj_err_path=inj_err_path,
diff_log_path=diff_log_path, diff_err_path=diff_err_path)
if cp.DEBUG:
cf.printf("THREAD {} CHECK SDCs OK".format(host_thread))
# Check if the carolfi logfile contains the information
# to confirm the fault injection outcome
block, fi_successful, new_value, old_value, register, thread, instruction = check_injection_outcome(
host_thread=host_thread,
logging=logging,
injection_site=injection_site
)
# Copy output files to a folder
save_output(is_sdc=is_sdc, is_hang=is_hang,is_crash=is_crash[0],is_masked=is_masked, logging=logging, unique_id=unique_id,
flip_log_file=flip_log_file, inj_output_path=inj_output_path, inj_err_path=inj_err_path,
diff_log_path=diff_log_path, diff_err_path=diff_err_path, signal_app_log_path=signal_app_log,
thread=host_thread)
if cp.DEBUG:
cf.printf("THREAD {} SAVE OUTPUT AND RETURN".format(host_thread))
return_list = [register, old_value, new_value, fi_successful,
is_hang, is_crash, is_sdc, is_masked, signal_init_wait_time, block, thread, instruction, user_defined_string]
return return_list
"""
This function will select the bits that will be flipped
if it is least significant bits it will reduce the starting bit range
"""
def bit_flip_selection(fault_model):
# Randomly select (a) bit(s) to flip
# Max double bit flip
max_size_register_fault_model = cp.SINGLE_MAX_SIZE_REGISTER
# Max size of bits to flip is 2, max double bit flip
bits_to_flip = [0]
# Single bit flip
if fault_model == cp.FLIP_SINGLE_BIT:
bits_to_flip[0] = random.randint(0, max_size_register_fault_model - 1)
# Double bit flip
elif fault_model == cp.FLIP_TWO_BITS:
bits_to_flip = [0] * 2
bits_to_flip[0] = random.randint(0, max_size_register_fault_model - 1)
# Make sure that the same bit is not going to be selected
r = [i for i in range(0, bits_to_flip[0])]
r += [i for i in range(bits_to_flip[0] + 1, max_size_register_fault_model)]
bits_to_flip[1] = random.choice(r)
# Random value
elif fault_model == cp.RANDOM_VALUE:
bits_to_flip[0] = str(hex(random.randint(0, cp.MAX_INT_32)))
# Zero value
elif fault_model == cp.ZERO_VALUE:
bits_to_flip[0] = 0
# Least 16 bits
elif fault_model == cp.LEAST_16_BITS:
bits_to_flip[0] = random.randint(0, 15)
# Least 8 bits
elif fault_model == cp.LEAST_8_BITS:
bits_to_flip[0] = random.randint(0, 7)
return bits_to_flip
"""
print the info for each fault
"""
def pretty_print(header, row):
fault_injected = row[9]
normal_print = "\033[0;37;49m"
failed_print = "\033[1;37;41m"
injected_print = "\033[1;37;42m"
output_str = "fault status: "
output_str += injected_print + "Injected" if fault_injected else failed_print + "Failed"
output_str += normal_print
cf.printf(output_str)
output_str = ""
for name, value in zip(header, row):
if name != "fault_successful":
output_str += "{}: {}\n".format(name, value)
cf.printf(output_str)
cf.printf()
"""
This injector has two injection options
this function performs fault injection
by sending a SIGINT signal to the application
"""
def fault_injection_by_signal(**kwargs):
# Global rows list
global lock, exit_injector
benchmark_binary = kwargs.get('benchmark_binary')
kwargs['signal_cmd'] = "killall -2 {}".format(os.path.basename(benchmark_binary))
fault_models = kwargs.get('fault_models')
iterations = kwargs.get('iterations')
host_thread = kwargs.get('host_thread')
injection_site = kwargs.get('injection_site')
summary_file = kwargs.get('summary_file')
header = kwargs.get('header')
cf.printf("-----------------------------------------------------------------------------------------------")
# Execute the fault injector for each one of the sections(apps) of the configuration file
for fault_model in fault_models:
# Execute iterations number of fault injection for a specific app
num_rounds = 1
while num_rounds <= iterations:
if exit_injector:
return
# Generate an unique id for this fault injection
# Thread is for multi gpu
unique_id = "{}_{}_{}".format(num_rounds, fault_model, host_thread)
bits_to_flip = bit_flip_selection(fault_model=fault_model)
kwargs['unique_id'] = unique_id
kwargs['bits_to_flip'] = bits_to_flip
kwargs['fault_model'] = fault_model
fi_tic = int(time.time())
[register, old_val, new_val, fault_injected,
hang, crash, masked,sdc, signal_init_time, block,
thread, instruction, user_defined_val] = gdb_inject_fault(**kwargs)
# Time toc
fi_toc = int(time.time())
# FI injection time
injection_time = fi_toc - fi_tic
if fault_injected:
output_str = "THREAD:{}, FAULT NUM:{}".format(host_thread, num_rounds)
if len(crash)==1:
tmp="--"
else:
tmp=crash[1]
row = [unique_id, register, num_rounds, fault_model, thread,
block, old_val, new_val, injection_site,
fault_injected, hang, crash[0], sdc, masked ,tmp,injection_time,
signal_init_time, bits_to_flip, instruction, user_defined_val]
for name, value in zip(header, row):
output_str += " {}: {},".format(name, value)
# :-1 to remove the last comma
cf.printf(output_str[:-1])
with lock:
summary_file.write_row(row)
num_rounds += 1
pretty_print(header=header, row=row)
"""
Main function
"""
def main():
global kill_strings, current_path, gpus_threads, lock, syncro, wait_finish
signal.signal(signal.SIGUSR1,receiveSignal);
signal.signal(signal.SIGUSR2,receiveEnd);
parser = argparse.ArgumentParser()
parser.add_argument('-c', '--conf', dest="config_file", help='Configuration file', required=True)
parser.add_argument('-i', '--iter', dest="iterations",
help='How many times to repeat the programs in the configuration file', required=True, type=int)
parser.add_argument('-n', '--n_gpus', dest="n_gpus", help="The number of available GPUs to perform FI."
" Default is 1.", required=False, default=1, type=int)
parser.add_argument('-d', '--device', dest="device", help="The GPU to perform FI."
" Default is 0.", required=False, default=0, type=int)
args = parser.parse_args()
if args.iterations < 1:
parser.error('Iterations must be greater than zero')
# Start with a different seed every time to vary the random numbers generated
# the seed will be the current number of second since 01/01/70
random.seed()
# Read the configuration file with data for all the apps that will be executed
conf = cf.load_config_file(args.config_file)
# Connect signal SIGINT to stop the fault injector
kill_strings = ""
signal.signal(signal.SIGINT, signal_handler)
# First set env vars
current_path = cf.set_python_env()
cf.printf("2 - Starting fault injection")
cf.printf("###################################################")
cf.printf("2 - {} faults will be injected".format(args.iterations))
cf.printf("###################################################")
########################################################################
# Creating a summary csv file
csv_file = conf.get("DEFAULT", "csvFile")
# Csv log
fieldnames = ['unique_id', 'register', 'iteration', 'fault_model', 'thread', 'block', 'old_value',
'new_value', 'inj_site', 'fault_successful', 'hang', 'crash', 'masked', 'sdc', 'Exception','time',
'inj_time_location', 'bits_flipped', 'instruction', 'user_defined']
summary_file = SummaryFile(filename=csv_file, fieldnames=fieldnames, mode='w')
# Lock for summary file parallel
lock = Lock()
# Define the number of threads tha will execute
num_gpus = args.n_gpus
device=args.device
iterations = args.iterations
if args.n_gpus > args.iterations:
num_gpus = args.iterations
bin_path = current_path + '/bin'
if not os.path.exists(bin_path):
os.mkdir(bin_path)
# Create tmp path and clean it if it exists
tmp_path = current_path + "/" + cp.LOGS_PATH + "/tmp"
if not os.path.exists(tmp_path):
raise FileNotFoundError(tmp_path + " path does not exists, run app_profile.py to create it")
# Set binaries for the injection
benchmark_binary_default = conf.get('DEFAULT', 'benchmarkBinary')
gdb_path_default = conf.get('DEFAULT', 'gdbExecName')
each_thread_iterations = iterations / num_gpus
gpus_threads = []
kernel_info_dict = cf.load_file(cp.KERNEL_INFO_DIR)
for thread_id in range(0+device, num_gpus+device):
gdb = "{}/bin/{}_{}".format(current_path, os.path.basename(gdb_path_default), thread_id)
benchmark_binary = "{}/bin/{}_{}".format(current_path, os.path.basename(benchmark_binary_default), thread_id)
os.system("ln -s {} {}".format(gdb_path_default, gdb))
os.system("ln -s {} {}".format(benchmark_binary_default, benchmark_binary))
# These are the mandatory parameters
kwargs = {
'injection_site': conf.get('DEFAULT', 'injectionSite'),
'fault_models': [int(i) for i in str(conf.get('DEFAULT', 'faultModel')).split(',')],
'max_time': float(kernel_info_dict['max_time']) * float(conf.get('DEFAULT', 'maxWaitTimes')),
'end_time': float(kernel_info_dict['max_time_kernel']),
'iterations': each_thread_iterations,
'benchmark_binary': benchmark_binary,
'benchmark_args': conf.get('DEFAULT', 'benchmarkArgs'),
'host_thread':thread_id,
'gdb_path': gdb,
'current_path': current_path,
'seq_signals': int(conf.get('DEFAULT', 'seqSignals')),
'init_sleep': float(conf.get('DEFAULT', 'initSleep')),
'kernel':conf.get('DEFAULT', 'kernel'),
'gold_check_script': "{}/{}".format(current_path, conf.get('DEFAULT', 'goldenCheckScript')),
'summary_file': summary_file,
'header': fieldnames
}
syncro = threading.Barrier(2, timeout=kwargs.get('max_time') )
wait_finish = threading.Barrier(2, timeout=kwargs.get('max_time'))
kill_strings += "killall -9 {};killall -9 {};".format(os.path.basename(benchmark_binary), os.path.basename(gdb))
fi_master_thread = Thread(target=fault_injection_by_signal, kwargs=kwargs)
gpus_threads.append(fi_master_thread)
for thread in gpus_threads:
thread.start()
for thread in gpus_threads:
thread.join()
os.system("rm -f {}/bin/*".format(current_path))
if exit_injector:
cf.printf("\nKeyboardInterrupt detected, exiting gracefully!( at least trying :) )")
else:
cf.printf("Fault injection finished, results can be found in {}".format(csv_file))
########################################################################
########################################################################
# Main #
########################################################################
kill_strings = None
current_path = None
lock = None
exit_injector = False
gpus_threads = []
if __name__ == "__main__":
main()
flip_value.py 0 → 100755
View file @ e4ae2422
import os
import gdb
import time
from classes.BitFlip import BitFlip
from classes.Logging import Logging
import common_parameters as cp
"""
Handler attached to exit event
"""
def exit_handler(event):
global global_logging
global_logging.info(str("event type: exit"))
print ("llego el final")
os.system ("kill -s USR2 " + str(pid))
try:
global_logging.info("exit code: {}".format(str(event.exit_code)))
except Exception as err:
err_str = "ERROR: {}".format(str(err))
global_logging.exception(err_str)
"""
Handler that will put a breakpoint on the kernel after
signal
"""
def set_event(event):
# Accessing global vars
global global_logging, was_hit, bit_lip,bp,t
if (isinstance(event, gdb.BreakpointEvent)):
global_logging.info("Before breakpoint"+ str(time.clock()-t))
global_logging.info ("Enviado senal a "+ str(pid))
os.system ("kill -s USR1 " + str(pid))
bp.enabled=False
gdb.execute('c')
# #os.system ("killall -2 python3")
else:
try:
# Just checking if it was hit
if bit_flip.fault_injected is False:
bit_flip.single_event()
global_logging.info("BIT FLIP SET ON SIGNAL {}".format(event.stop_signal))
except Exception as err:
global_logging.exception("EVENT DIFFERENT FROM STOP SIGNAL: {}".format(str(err)))
"""
Main function
"""
def main():
global global_logging, register, injection_site, bits_to_flip, fault_model, was_hit, bit_flip, arg0
was_hit = False
# Initialize GDB to run the app
gdb.execute("set confirm off")
gdb.execute("set pagination off")
gdb.execute("set target-async off")
gdb.execute("set non-stop off")
# Connecting to a exit handler event
gdb.events.exited.connect(exit_handler)
# Connecting to a stop signal event
gdb.events.stop.connect(set_event)
# Get variables values from environment
# Firsn parse line
[kernel,pid,bits_to_flip, fault_model, flip_log_file,
gdb_init_strings, injection_site] = arg0.split('|')
# Logging
global_logging = Logging(log_file=flip_log_file)
global_logging.info("Starting flip_value script "+" called by " + str(pid) + " for stop kernel " + str(kernel));
try:
for init_str in gdb_init_strings.split(";"):
gdb.execute(init_str)
global_logging.info("initializing setup: " + str(init_str))
except gdb.error as err:
global_logging.exception("ERROR on initializing setup: {}".format(str(err)))
# Set Breakpoint attributes to be use
bits_to_flip = [i for i in bits_to_flip.split(",")]
fault_model = int(fault_model)
bit_flip = BitFlip(bits_to_flip=bits_to_flip, fault_model=fault_model,
logging=global_logging, injection_site=cp.INJECTION_SITES[injection_site])
# Start app execution
t=time.clock();
#gdb.execute("break "+kernel)
bp=gdb.Breakpoint(kernel)
global_logging.info("Put Break "+ str(time.clock()-t))
gdb.execute("r")
i = 0
try:
while 'The program' not in gdb.execute('c', to_string=True):
i += 1
except Exception as err:
global_logging.info("CONTINUED {} times".format(i))
err_str = str(err).rstrip()
global_logging.exception("IGNORED CONTINUE ERROR: {}".format(err_str))
# Make sure that it is going to finish
if 'Failed' in err_str:
gdb.execute('quit')
global_logging.exception("QUIT REQUIRED")
# Call main execution
global_logging = None
register = None
bits_to_flip = None
fault_model = None
was_hit = False
injection_site = None
bit_flip = None
main()
process_start.py 0 → 100644
View file @ e4ae2422
from subprocess import Popen
from sys import argv
with open(argv[1], "w") as fp:
fp.write(str(Popen(argv[2]).pid))
profiler.py 0 → 100644
View file @ e4ae2422
import gdb
"""
Main function
"""
# Initialize GDB to run the app
gdb.execute("set confirm off")
gdb.execute("set pagination off")
gdb.execute("set target-async off")
gdb.execute("set non-stop off")
# gdb_init_strings = str(os.environ["CAROL_FI_INFO"])
gdb_init_strings = arg0
for init_str in gdb_init_strings.split(";"):
gdb.execute(init_str)
gdb.execute("r")
profiler_new.py 0 → 100644
View file @ e4ae2422
import gdb
import time
def exit_handler(event):
global nosalir
nosalir=False
print(str("event type: exit"))
try:
print("exit code: {}".format(str(event.exit_code)))
except Exception as err:
err_str = "ERROR: {}".format(str(err))
print(err_str)
"""
Handler that will put a breakpoint on the kernel after
signal
"""
def set_event(event):
global trun,ocurrencias,t,primera
print ("Es mi primera vez"+ str(primera)+" "+str(ocurrencias))
if (isinstance(event, gdb.BreakpointEvent)):
if (primera):
t=time.clock()
ocurrencias=ocurrencias+1
else:
trun=(time.clock()-t)
primera=not primera
else:
trun=(time.clock()-t)
"""
Main function
"""
def main():
global ocurrencias,t,nosalir,trun,primera
primera=True ;
ocurrencias=0;
# Initialize GDB to run the app
gdb.execute("set confirm off")
gdb.execute("set pagination off")
gdb.execute("set target-async off")
gdb.execute("set non-stop off")
# Connecting to a exit handler event
gdb.events.exited.connect(exit_handler)
# Connecting to a stop signal event
gdb.events.stop.connect(set_event)
# gdb_init_strings = str(os.environ["CAROL_FI_INFO"])
gdb_init_strings = arg0
cadena=gdb_init_strings.split(";",2)
#print (cadena,"-",cadena[0],'-',cadena[1],'-', cadena[2])
section =cadena[0]=="True"
kernel_end=cadena[1]
#print ("B "+section+"ke "+kernel_end+" ....")
#print (cadena[2].split(";"))
for init_str in cadena[2].split(";"):
gdb.execute(init_str)
if (section):
gdb.execute ("break "+kernel_end)
gdb.execute("r")
#nosalir=True
#while nosalir:
if (section):
#print ("Point 1")
gdb.execute("c")
else:
gdb.execute("finish")
#print ("Punto 2")
#print (" Ocurrencias "+str(ocurrencias)+" Tiempo acumulado de ejecucciones "+ str(trun)+ "\n")
gdb.execute("c")
#print (" Ocurrencias "+str(ocurrencias)+" Tiempo acumulado de ejecucciones "+ str(trun))
f=open("tmpxxx_return_profiler.conf","w")
f.write("[DEFAULT] \nOcurrencias = "+str(ocurrencias)+"\nTiempo = "+str(trun)+"\n")
f.close()
#print ("End write file \n")
#sys.stdout.flush()
main()
prueba.py 0 → 100644
View file @ e4ae2422
import os
import gdb
import time
def exit_handler(event):
global nosalir
nosalir=False
print(str("event type: exit"))
try:
print("exit code: {}".format(str(event.exit_code)))
except Exception as err:
err_str = "ERROR: {}".format(str(err))
print(err_str)
"""
Handler that will put a breakpoint on the kernel after
signal
"""
def set_event(event):
global trun,ocurrencias,t
if (isinstance(event, gdb.BreakpointEvent)):
t=time.clock()
ocurrencias=ocurrencias+1
else:
trun=(time.clock()-t)
def main():
global ocurrencias,t,nosalir,trun
was_hit = False
ocurrencias=0
# Initialize GDB to run the appset pagination off
gdb.execute("set confirm off")
gdb.execute("set pagination off")
gdb.execute("set target-async off")
gdb.execute("set non-stop off")
# Connecting to a exit handler event
gdb.events.exited.connect(exit_handler)
# Connecting to a stop signal event
gdb.events.stop.connect(set_event)
gdb.execute("file ~/rodinia_3.1/cuda/lud/cuda/lud_cuda")
gdb.execute("set arg -s 10000")
gdb.execute("break lud_cuda")
gdb.execute('r')
nosalir=True
while nosalir:
gdb.execute("finish")
gdb.execute("c")
print (" Ocurrencias "+str(ocurrencias)+" Tiempo acumulado de ejecucciones "+ str(trun))
f=open("tmpxxx_return_profiler.conf","w")
f.write("Ocurrencias ="+str(ocurrencias)+"\n Tiempo "+str(trun)+"\n")
f.close()
main()
Supports Markdown
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment