From 0c172b8e640268b507348aab962f7f92893bcce2 Mon Sep 17 00:00:00 2001 From: Xiangyu Gao Date: Wed, 26 Nov 2025 23:09:50 -0800 Subject: [PATCH] add perthreadRuntimeDist --- example/gpu/perthreadRuntimeDist/.gitignore | 2 + example/gpu/perthreadRuntimeDist/Makefile | 145 +++++++++++++++++ .../perthreadRuntimeDist.bpf.c | 62 ++++++++ .../perthreadRuntimeDist.c | 61 ++++++++ .../perthreadRuntimeDist.h | 9 ++ .../thread_time_distribution_demo.cu | 147 ++++++++++++++++++ example/gpu/perthreadRuntimeDist/vec_add | Bin 0 -> 97016 bytes example/gpu/perthreadRuntimeDist/vec_add.cu | 85 ++++++++++ 8 files changed, 511 insertions(+) create mode 100644 example/gpu/perthreadRuntimeDist/.gitignore create mode 100644 example/gpu/perthreadRuntimeDist/Makefile create mode 100644 example/gpu/perthreadRuntimeDist/perthreadRuntimeDist.bpf.c create mode 100644 example/gpu/perthreadRuntimeDist/perthreadRuntimeDist.c create mode 100644 example/gpu/perthreadRuntimeDist/perthreadRuntimeDist.h create mode 100644 example/gpu/perthreadRuntimeDist/thread_time_distribution_demo.cu create mode 100755 example/gpu/perthreadRuntimeDist/vec_add create mode 100644 example/gpu/perthreadRuntimeDist/vec_add.cu diff --git a/example/gpu/perthreadRuntimeDist/.gitignore b/example/gpu/perthreadRuntimeDist/.gitignore new file mode 100644 index 000000000..dda11ef6e --- /dev/null +++ b/example/gpu/perthreadRuntimeDist/.gitignore @@ -0,0 +1,2 @@ +/perthreadRuntimeDist +/.output \ No newline at end of file diff --git a/example/gpu/perthreadRuntimeDist/Makefile b/example/gpu/perthreadRuntimeDist/Makefile new file mode 100644 index 000000000..9d6848501 --- /dev/null +++ b/example/gpu/perthreadRuntimeDist/Makefile @@ -0,0 +1,145 @@ +# SPDX-License-Identifier: (LGPL-2.1 OR BSD-2-Clause) +OUTPUT := .output +CLANG ?= clang +LIBBPF_SRC := $(abspath ../../../third_party/libbpf/src) +BPFTOOL_SRC := $(abspath ../../../third_party/bpftool/src) +LIBBPF_OBJ := $(abspath $(OUTPUT)/libbpf.a) +BPFTOOL_OUTPUT ?= $(abspath $(OUTPUT)/bpftool) +BPFTOOL ?= $(BPFTOOL_OUTPUT)/bootstrap/bpftool +ARCH ?= $(shell uname -m | sed 's/x86_64/x86/' \ + | sed 's/arm.*/arm/' \ + | sed 's/aarch64/arm64/' \ + | sed 's/ppc64le/powerpc/' \ + | sed 's/mips.*/mips/' \ + | sed 's/riscv64/riscv/' \ + | sed 's/loongarch64/loongarch/') +VMLINUX := ../../../third_party/vmlinux/$(ARCH)/vmlinux.h +# Use our own libbpf API headers and Linux UAPI headers distributed with +# libbpf to avoid dependency on system-wide headers, which could be missing or +# outdated +INCLUDES := -I$(OUTPUT) -I../../../third_party/libbpf/include/uapi -I$(dir $(VMLINUX)) +CFLAGS := -g -Wall +ALL_LDFLAGS := $(LDFLAGS) $(EXTRA_LDFLAGS) + +APPS = perthreadRuntimeDist # minimal minimal_legacy uprobe kprobe fentry usdt sockfilter tc ksyscall + +CARGO ?= $(shell which cargo) +ifeq ($(strip $(CARGO)),) +BZS_APPS := +else +BZS_APPS := # profile +APPS += $(BZS_APPS) +# Required by libblazesym +ALL_LDFLAGS += -lrt -ldl -lpthread -lm +endif + +# Get Clang's default includes on this system. We'll explicitly add these dirs +# to the includes list when compiling with `-target bpf` because otherwise some +# architecture-specific dirs will be "missing" on some architectures/distros - +# headers such as asm/types.h, asm/byteorder.h, asm/socket.h, asm/sockios.h, +# sys/cdefs.h etc. might be missing. +# +# Use '-idirafter': Don't interfere with include mechanics except where the +# build would have failed anyways. +CLANG_BPF_SYS_INCLUDES ?= $(shell $(CLANG) -v -E - &1 \ + | sed -n '/<...> search starts here:/,/End of search list./{ s| \(/.*\)|-idirafter \1|p }') + +ifeq ($(V),1) + Q = + msg = +else + Q = @ + msg = @printf ' %-8s %s%s\n' \ + "$(1)" \ + "$(patsubst $(abspath $(OUTPUT))/%,%,$(2))" \ + "$(if $(3), $(3))"; + MAKEFLAGS += --no-print-directory +endif + +define allow-override + $(if $(or $(findstring environment,$(origin $(1))),\ + $(findstring command line,$(origin $(1)))),,\ + $(eval $(1) = $(2))) +endef + +$(call allow-override,CC,$(CROSS_COMPILE)cc) +$(call allow-override,LD,$(CROSS_COMPILE)ld) + +.PHONY: all +all: $(APPS) vec_add + +vec_add: vec_add.cu + @if command -v nvcc >/dev/null 2>&1; then \ + nvcc -arch=sm_61 -cudart shared vec_add.cu -o vec_add -g; \ + else \ + echo "Warning: CUDA not found, skipping vec_add build"; \ + fi + +.PHONY: clean +clean: + $(call msg,CLEAN) + $(Q)rm -rf $(OUTPUT) $(APPS) vec_add + +$(OUTPUT) $(OUTPUT)/libbpf $(BPFTOOL_OUTPUT): + $(call msg,MKDIR,$@) + $(Q)mkdir -p $@ + +# Build libbpf +$(LIBBPF_OBJ): $(wildcard $(LIBBPF_SRC)/*.[ch] $(LIBBPF_SRC)/Makefile) | $(OUTPUT)/libbpf + $(call msg,LIB,$@) + $(Q)$(MAKE) -C $(LIBBPF_SRC) BUILD_STATIC_ONLY=1 \ + OBJDIR=$(dir $@)/libbpf DESTDIR=$(dir $@) \ + INCLUDEDIR= LIBDIR= UAPIDIR= \ + install + +# Build bpftool +$(BPFTOOL): | $(BPFTOOL_OUTPUT) + $(call msg,BPFTOOL,$@) + $(Q)$(MAKE) ARCH= CROSS_COMPILE= OUTPUT=$(BPFTOOL_OUTPUT)/ -C $(BPFTOOL_SRC) bootstrap + + +$(LIBBLAZESYM_SRC)/target/release/libblazesym.a:: + $(Q)cd $(LIBBLAZESYM_SRC) && $(CARGO) build --features=cheader,dont-generate-test-files --release + +$(LIBBLAZESYM_OBJ): $(LIBBLAZESYM_SRC)/target/release/libblazesym.a | $(OUTPUT) + $(call msg,LIB, $@) + $(Q)cp $(LIBBLAZESYM_SRC)/target/release/libblazesym.a $@ + +$(LIBBLAZESYM_HEADER): $(LIBBLAZESYM_SRC)/target/release/libblazesym.a | $(OUTPUT) + $(call msg,LIB,$@) + $(Q)cp $(LIBBLAZESYM_SRC)/target/release/blazesym.h $@ + +# Build BPF code +$(OUTPUT)/%.bpf.o: %.bpf.c $(LIBBPF_OBJ) $(wildcard %.h) $(VMLINUX) | $(OUTPUT) $(BPFTOOL) + $(call msg,BPF,$@) + $(Q)$(CLANG) -Xlinker --export-dynamic -g -O2 -target bpf -D__TARGET_ARCH_$(ARCH) \ + $(INCLUDES) $(CLANG_BPF_SYS_INCLUDES) \ + -c $(filter %.c,$^) -o $(patsubst %.bpf.o,%.tmp.bpf.o,$@) + $(Q)$(BPFTOOL) gen object $@ $(patsubst %.bpf.o,%.tmp.bpf.o,$@) + +# Generate BPF skeletons +$(OUTPUT)/%.skel.h: $(OUTPUT)/%.bpf.o | $(OUTPUT) $(BPFTOOL) + $(call msg,GEN-SKEL,$@) + $(Q)$(BPFTOOL) gen skeleton $< > $@ + +# Build user-space code +$(patsubst %,$(OUTPUT)/%.o,$(APPS)): %.o: %.skel.h + +$(OUTPUT)/%.o: %.c $(wildcard %.h) | $(OUTPUT) + $(call msg,CC,$@) + $(Q)$(CC) $(CFLAGS) $(INCLUDES) -c $(filter %.c,$^) -o $@ + +$(patsubst %,$(OUTPUT)/%.o,$(BZS_APPS)): $(LIBBLAZESYM_HEADER) + +$(BZS_APPS): $(LIBBLAZESYM_OBJ) + +# Build application binary +$(APPS): %: $(OUTPUT)/%.o $(LIBBPF_OBJ) | $(OUTPUT) + $(call msg,BINARY,$@) + $(Q)$(CC) $(CFLAGS) $^ $(ALL_LDFLAGS) -lelf -lz -o $@ + +# delete failed targets +.DELETE_ON_ERROR: + +# keep intermediate (.skel.h, .bpf.o, etc) targets +.SECONDARY: diff --git a/example/gpu/perthreadRuntimeDist/perthreadRuntimeDist.bpf.c b/example/gpu/perthreadRuntimeDist/perthreadRuntimeDist.bpf.c new file mode 100644 index 000000000..a6be3fee6 --- /dev/null +++ b/example/gpu/perthreadRuntimeDist/perthreadRuntimeDist.bpf.c @@ -0,0 +1,62 @@ +#include "vmlinux.h" +#include +#include +#include "perthreadRuntimeDist.h" + + +struct { + __uint(type, BPF_MAP_TYPE_PERF_EVENT_ARRAY); + __uint(key_size, sizeof(__u32)); + __uint(value_size, sizeof(__u32)); +} events SEC(".maps"); + +// 用于记录 start time +struct { + __uint(type, BPF_MAP_TYPE_HASH); + __uint(max_entries, 8192); + __type(key, u32); + __type(value, u64); +} start SEC(".maps"); + +char LICENSE[] SEC("license") = "Dual BSD/GPL"; + +// EXT helper index — 必须存在且可见 +static const u64 (*bpf_get_globaltimer)(void) = (void *)502; + +// GPU kernel entry +SEC("kprobe/cudaLaunchKernel") +int cuda__kernel_entry(struct pt_regs *ctx) +{ + u32 tid = bpf_get_current_pid_tgid(); + + u64 start_cycles = bpf_get_globaltimer(); // bpf_ktime_get_ns()? + + bpf_map_update_elem(&start, &tid, &start_cycles, BPF_ANY); + + return 0; +} + +SEC("kretprobe/cudaLaunchKernel") +int cuda__kernel_exit(struct pt_regs *ctx) +{ + u32 tid = bpf_get_current_pid_tgid(); + u64 *start_cycles = bpf_map_lookup_elem(&start, &tid); + + if (!start_cycles) + return 0; + + u64 end_cycles = bpf_get_globaltimer(); + + struct event_t evt = { + .tid = tid, + .cycles = end_cycles - *start_cycles, + }; + + bpf_perf_event_output(ctx, &events, + BPF_F_CURRENT_CPU, + &evt, + sizeof(evt)); + + bpf_map_delete_elem(&start, &tid); + return 0; +} diff --git a/example/gpu/perthreadRuntimeDist/perthreadRuntimeDist.c b/example/gpu/perthreadRuntimeDist/perthreadRuntimeDist.c new file mode 100644 index 000000000..9d2d8c52c --- /dev/null +++ b/example/gpu/perthreadRuntimeDist/perthreadRuntimeDist.c @@ -0,0 +1,61 @@ +#include +#include +#include +#include "perthreadRuntimeDist.skel.h" +#include "perthreadRuntimeDist.h" + +static void handle_event(void *ctx, int cpu, void *data, __u32 size) +{ + struct event_t *e = data; + printf("[CPU %d] tid=%u cycles=%llu ns\n", + cpu, e->tid, e->cycles); +} + +static void handle_lost(void *ctx, int cpu, __u64 lost) +{ + printf("LOST %llu events on CPU %d\n", lost, cpu); +} + +int main() +{ + struct perthreadRuntimeDist_bpf *skel; + struct perf_buffer *pb; + int events_fd; + + skel = perthreadRuntimeDist_bpf__open(); + if (!skel) { + fprintf(stderr, "Failed to open BPF skeleton\n"); + return 1; + } + + if (perthreadRuntimeDist_bpf__load(skel)) { + fprintf(stderr, "Failed to load BPF skeleton\n"); + return 1; + } + + if (perthreadRuntimeDist_bpf__attach(skel)) { + fprintf(stderr, "Failed to attach BPF skeleton\n"); + return 1; + } + + printf("BPF attached successfully\n"); + + events_fd = bpf_map__fd(skel->maps.events); + pb = perf_buffer__new(events_fd, 16 /*buffer pages*/, + handle_event, handle_lost, NULL, NULL); + + if (!pb) { + fprintf(stderr, "Failed to open perf buffer\n"); + return 1; + } + + printf("Collecting data...\n"); + + while (1) { + int err = perf_buffer__poll(pb, 100 /*ms*/); + if (err < 0) + printf("perf_buffer__poll() error %d\n", err); + } + + return 0; +} diff --git a/example/gpu/perthreadRuntimeDist/perthreadRuntimeDist.h b/example/gpu/perthreadRuntimeDist/perthreadRuntimeDist.h new file mode 100644 index 000000000..fec960828 --- /dev/null +++ b/example/gpu/perthreadRuntimeDist/perthreadRuntimeDist.h @@ -0,0 +1,9 @@ +#ifndef __PERTHREAD_RUNTIME_DIST_H__ +#define __PERTHREAD_RUNTIME_DIST_H__ + +struct event_t { + __u32 tid; + __u64 cycles; +}; + +#endif diff --git a/example/gpu/perthreadRuntimeDist/thread_time_distribution_demo.cu b/example/gpu/perthreadRuntimeDist/thread_time_distribution_demo.cu new file mode 100644 index 000000000..e988a80c8 --- /dev/null +++ b/example/gpu/perthreadRuntimeDist/thread_time_distribution_demo.cu @@ -0,0 +1,147 @@ +// thread_time_distribution_demo.cu +#include +#include +#include + +// Simple CUDA error checking +#define CHECK_CUDA(call) \ + do { \ + cudaError_t err__ = (call); \ + if (err__ != cudaSuccess) { \ + fprintf(stderr, "CUDA error %s (%d) at %s:%d\n", \ + cudaGetErrorString(err__), err__, __FILE__, __LINE__); \ + std::exit(EXIT_FAILURE); \ + } \ + } while (0) + +// Per-thread timing info +struct ThreadTime { + unsigned int blockId; // block index (1D grid) + unsigned int threadId; // global thread ID (1D) + unsigned long long cycles; // cycles spent in work section +}; + +// Kernel: each thread runs a synthetic workload, and we measure its time +__global__ void timed_work_kernel(ThreadTime* out, int base_iters) +{ + const unsigned int blockId = blockIdx.x; + const unsigned int threadIdInBlock = threadIdx.x; + const unsigned int globalThreadId = + blockIdx.x * blockDim.x + threadIdx.x; + + // Each thread runs a slightly different number of iterations + // to produce a non-trivial distribution. + int my_iters = base_iters + (globalThreadId & 0x1F); // vary per thread + + // Time the "work" section using clock64() + unsigned long long start = clock64(); + + volatile float acc = 0.0f; + for (int i = 0; i < my_iters; ++i) { + acc += 1.0f; // trivial arithmetic + } + + unsigned long long end = clock64(); + + // Prevent the compiler from optimizing out the loop completely + if (acc == -1.0f) { + printf("This will never be printed\n"); + } + + ThreadTime t; + t.blockId = blockId; + t.threadId = globalThreadId; + t.cycles = end - start; + + out[globalThreadId] = t; +} + +int main() +{ + // Again, keep it small for printing + const int BLOCKS = 4; + const int THREADS_PER_BLOCK = 64; + const int TOTAL_THREADS = BLOCKS * THREADS_PER_BLOCK; + + const int BASE_ITERS = 100000; // base workload per thread + + printf("=== Per-thread runtime distribution demo ===\n"); + printf("Grid: %d blocks, Block: %d threads (total %d threads)\n", + BLOCKS, THREADS_PER_BLOCK, TOTAL_THREADS); + + // Allocate device and host buffers + ThreadTime* d_times = nullptr; + CHECK_CUDA(cudaMalloc(&d_times, TOTAL_THREADS * sizeof(ThreadTime))); + + ThreadTime* h_times = (ThreadTime*)std::malloc(TOTAL_THREADS * sizeof(ThreadTime)); + if (!h_times) { + fprintf(stderr, "Host malloc failed\n"); + std::exit(EXIT_FAILURE); + } + + // Launch kernel + dim3 grid(BLOCKS); + dim3 block(THREADS_PER_BLOCK); + + timed_work_kernel<<>>(d_times, BASE_ITERS); + CHECK_CUDA(cudaGetLastError()); + CHECK_CUDA(cudaDeviceSynchronize()); + + // Copy results back + CHECK_CUDA(cudaMemcpy(h_times, d_times, + TOTAL_THREADS * sizeof(ThreadTime), + cudaMemcpyDeviceToHost)); + + // Get device clock rate to convert cycles -> time + cudaDeviceProp prop; + CHECK_CUDA(cudaGetDeviceProperties(&prop, 0)); + + // clockRate: kHz (cycles per millisecond) + const double sm_freq_hz = (double)prop.clockRate * 1000.0; + + printf("\nDevice: %s\n", prop.name); + printf("SM clock rate: %.3f MHz\n\n", sm_freq_hz / 1.0e6); + + // Print per-thread timing + printf("Per-thread timing:\n"); + printf("tid block cycles time_us\n"); + printf("--------------------------------------\n"); + + unsigned long long min_cycles = ~0ULL; + unsigned long long max_cycles = 0; + unsigned long long sum_cycles = 0; + + for (int i = 0; i < TOTAL_THREADS; ++i) { + const ThreadTime& t = h_times[i]; + + double time_sec = (double)t.cycles / sm_freq_hz; + double time_us = time_sec * 1.0e6; + + printf("%3u %5u %10llu %10.3f\n", + t.threadId, t.blockId, + (unsigned long long)t.cycles, time_us); + + if (t.cycles < min_cycles) min_cycles = t.cycles; + if (t.cycles > max_cycles) max_cycles = t.cycles; + sum_cycles += t.cycles; + } + + double avg_cycles = (double)sum_cycles / TOTAL_THREADS; + double min_us = (double)min_cycles / sm_freq_hz * 1.0e6; + double max_us = (double)max_cycles / sm_freq_hz * 1.0e6; + double avg_us = avg_cycles / sm_freq_hz * 1.0e6; + + printf("\nSummary over all threads:\n"); + printf(" min cycles = %llu (%.3f us)\n", + (unsigned long long)min_cycles, min_us); + printf(" max cycles = %llu (%.3f us)\n", + (unsigned long long)max_cycles, max_us); + printf(" avg cycles = %.1f (%.3f us)\n", + avg_cycles, avg_us); + + // Cleanup + CHECK_CUDA(cudaFree(d_times)); + std::free(h_times); + + return 0; +} \ No newline at end of file diff --git a/example/gpu/perthreadRuntimeDist/vec_add b/example/gpu/perthreadRuntimeDist/vec_add new file mode 100755 index 0000000000000000000000000000000000000000..681c868ffb1ea9f7a42d410e97dc7f44769973f1 GIT binary patch literal 97016 zcmeFad3;nw);C^ttJB>#-ATHege(YY_ALn=k^mtP!V+jO>>$V%LJ}Yvk{ALaii0dl zCZM=Gu8gAMj*9EJprWAT0xq~Kf=f_Fac5k}`#q;_FUgqqdEe)K-`_t!FQ4?OQ{Qt= zRh_CjRrmI-<}bx1V+|=q=u)Ck5Vx>{&MC5hRdo@WA~8h7i6YTY^c8mCHvEO;+R}Fn zfVe``sac^wNiT&Ge?U*s*&;+kr4uE+_AXC?MnY70G*5all(hJ^^!_9*7os6NTwXxq zq76MKP3MPbv>y3pcZtS#&udLLS?%_7rGQ!WB6pGAJC@!%mR^YZEq{h+$e-jB`Ws@| zZA;r!TeQ_%YU#D5V_=i65T!Ecn*Lcs#eb)~DVAQ{NNuMreS#d0D?~%(O@$u$<$v|E z+$pA6<+axjMOM2*)GA&KsIFQxD0e`0MP_wX&5Bi-tA-599F&{2yf$kfSApu~#~*dy zxJfewx@9-zS>}YXR`Vfyge6)y$^SAnXF~DeK22v=j*s1P`2MnCAFWwSGE{HUAsM>J zp4H{dQ`(O|lF|OSl+%<)*Z#^^mnB`}9<90m{5%6$zW>Y?o7>5sgo6G^{^oY%FKI`< zw4MBo?c`UZp?_5F%y#nMYsb!(cI5r-$e(N{e^5Ji*0dx4c02x=+m4-C?Z{upy<&vu94VCpXCXTMIt*M+*wy3&Nl$I`8R$EiL zyrHbVp|n(swj)@+qM~ekO;y9h+KLs>Bz9uuvhuofX4FnUXW63KYMoV9U0quayM#}( ziX6M5ro5r5wkD*sys{ybguW#_vnEvV*s_LERW)VxA!JHz-DsE{U0bubYRQWFGB(w^ zgwx9E+m;zJQe0DEK~q*NU;1Aeij++a7*n~js=RXgIW^@=>uYPOR_kV$l%YLKPpz!4 zsjL>IbEY>843w5OK;rDu>dKlW4NFTa>+5Ulr<@8u%&DlJvT_;dr0KQQ%ZnGYeD>gw z{GzgoQuf$NZ792{c6sTdvgMVfh|=|NL1`6S4zP5|s#T?RmG#SOYs#vt8qO(Q8PFO- zTA8Y^FFU8SrZQ9!ZMAy&_~{MV+2u>i>Ps8y%c>fdk1sDSo`$>}6jfDTTD!cVzOroD z`0^R0#nbajb;WDWrfN@b$Stp3(I84^)|_2cQ&Bpta`}p7m122yWn~>|g@IFkW@-7- zGfNkjRlyll?Q&~4Ei0?4A zpZGzty78}f*AfI2PQ6tPxH3Kyp&RxcZ*h(3n7 zN$Z~W(WDu1n&Py4K{QTZOrjstn56S`8O^3fgP%O90*bs1xQqVj3~6}kdZ`JpX3 zVMC(w_5Mq16h-AzokCYIDjx$fd`*eUj}7MvF*_8SocSq%0Q;5-hQTeAt$?uQKr@h7hU$5^Iy8IGcf6N(diuq)+5W$VR z8;q9y!S#Eb&#_Bd^4^;yL`%;P@R^WaBuIWKAx96lw6rYV$~;ZLqpwGJns`T_jqo(# zjy@3KX`&syIl|KfJ9=4!r-^lRU4*9zb+kUhQvf=;D8kdkIyx=F(}X%YD#Fu5Iyx}I z(*!!&EyB~pIqHk>G+~Y!5uPT>(eF-%+e;JV=-~)Y6XWRX5uPT*(Ptw(O@yNlM0lD2 zM{kbs^z7m2Wf7hN;L&vvo+if8`Up=G;^?9<-_kQ4y|XwWeT?A#Y3Vs=VrxFD4d1H` z-=z)Tu?-*JhPSujPnNVU@258W*KPRE+wdQ?;SaXqUv0y`(1w4i4gYW({@ynHZEg4) z+wfPl;V*8(U)Y8}zYV{#6`%dD@r?%-1RDKhBI9m?saT3y}`MY-bM)apx(SK%e{qSc`MKajQwwybX$5=RX%l8rPf% zHWi-;HpbBSj7HWU^EaMY+*o|9BG_c>*@tUTd?I^y!Xw4Uko^~26 zCH6LBqsihMV7)FRzJbuN_+wfe5__8uproyPHxA=!3Gvmrixa+i@%w~+TPWxI1n-s% z&?d%%yW()G~56^k0Yme&k)*R(GeR>(RQ5L(a3F?aK?Eq(WO- z5=6qqD?o!ySCYHn-XbEp+#ZPF5g-@=T&;j95#T$l0cihLs`ca4bo5J;HmM7EjQWkx ze!o#^#i-`q=B;5#E~Aa4AjzhUT#AO^U2WvEU)ab7w)!}kehN1WZ2vIVWaBE-EgHp| zMkVecw5fPQ)5Oie#+QOkPm)~IinSwoNR50{7jPih_-ykhzqYh&+`YmNN9ZbS(Ur6u z^(vXW6iZ|CCI|+bZeSs6fILEeYTS$3lXL0D2<)@L#*do&lkk?9608Aw>^>f3(dGxJ zNsZ6JH8hBHUJ1-JVP&Y<9c=u#xrim6g?2|vyYcPjbJ-Iy9m4j`uE*FIJZM2cR!iGm&S!P z-fq@u9;H|Dc=?Y~&$LQ?1C%UrOpHXn5_lsjByNcrZMkPZ&4s3zT#KzZ(K>dnKjsOB zt?X^Cg!M@9!v?Mn=#F6HifzHhnb$-IK}6BL#k*)gG!^d(UV|tGu^G=FT3@_NhH{|P zCg+-P4tn<(kEWQvj}u~V@jmK{y~QtUvR{)oHF-yq4~PV}Ox%sa9u00;@h052k6d{P z8ntnE!j-!>>`vIY7DDUS>=G-AgH3C82b)$vwD?fd%w56q*MrX-RwY~9o?z3&H^(#; ze}Jyt*HpZJ)@g_@Jm5~J5@7KVio#~q^HpBRNXeFohuT%^W~tp_sbMD5nxdB+tSR^Y z+{Uebc6#HA<8W&;4KSEKPFlf6p1-Rx9JjXWh^TIgIeRQ>$HIG?e?r?SeTLBgLIDHa zj`^JQ4EBV<#sf$_La8lzw5C!EPeql`LPQbPl_MX+xD!^O*^+c)Snmq0cchpM?QMSJ zC)V?{t=*QG{4rYMERtx9d6pEA%jKbASY|N}@788oUPpSHWztVo*%C7X%Q}@Xl@_rV zpW^|({#c-ury3_78q-+(LDL#^9rXt}4o*ZHMmEBCdz)vV@UX7-8F{&BrsZ8;z6Sgl zC1e$_Z%pGeC5;pJW2`qfLb|c|9TZlA(SefQXL!(D0jyF!R>C~~tDbDEqiZv6XCNkx41B#E&k653M znb`-9X$E!(LkQ7gJ`E$#V({EZMPsq?8RWdfAs*xE<7QfUAcw{jLzm?iAZUKoB@~=E ze{b{8tPI7y&5PhHm_N~KG?LS(FqT*d3%c-WPI6YT$qY75gnf472|(ohPa(|RdqC-{_yVPpQmoOGpOW2C?Ex9a+T0aVx?Tij&*v+9f zJJI2^!+DK59{hgLL$o=goW0Grl0OQHA5FN77DCRt0t0*NhW8t~Q`@2^`!nD3a9)JF zH&XQx+f@808q;bSxt8HF zlQ;0xyH&dc3*T)W626WiaODi z&i&m)>j!2G&l>U}jnd-d5h;uS8V_6GTi!wOj49`;k&9d1IK22RMU`m(1yBRHcg21f z2IYciU{OJbxFD;iQH(X02ElgwKI4f5mvA+PhdpFK8(y*I!W7*f(bGr!w{c3Fs7)Ty z^=q8DFSvfqe$f!Hs!CO)svZX=AHh7uZ(gL}Bkt=7RM+Uf!R;3JPh?#q^-z7^Lsv&l z&sKw>Rl`1lV6@q&4h}UEjmv*_xgFMzf5ovhs^_f{65Iw*E;vfq zY8omNJ#N}~5mniexY26BajW+WzWIYKMN|`tbQ2J|pr|{LdSY+$?ZmGyZYH-K3gX^D z_bj!Sd-Xvr&yyjlrg|mT1@MBy4muH;DhQ8<5-!`F@bGRtQQ*B;Ls}ReLFn?B=De?Q z&wu=191#x%pE-mevbT95GTZhydM&C;Lt}*p(+f}2PNn%>YqieYTYDqg>F?mS8@Jxx;@`gVp`>Q8#lMjYLS5QSo?h|kRt%}g(;xjj zHQ1E2Y^V^~KOKGQ7(Kk;HJ_f*hM&r!=s&MNJ@BUo{`A0~9{AG(e|q3g5B%wYKRxi@ z@PNcSN_35$8<_7O?ic<(m8J_5OyXWi|d~Wvfb; z;|;hzeTCGA5~loXCcd8Wh#j=u#$^B~0u8;VPI+`W@D|W;o#fBxcU7XY#b5M!ha7xy z#Fv3DelCa-IHBDkR^gtRuHiAm1)X;=@M)(vjk1<6#e0AaWs8ss(Q}q*o@wdwr8O0+ zL{`nptg4#DwVWu&+lx!;%j%XgS65%#P}^`$T_tn%mDO2g`o+)xx5jmY+W8!ad_`Y@ zwYjM8=#9QVy6E`~<>MJRQ#==Dim%;TmPZu4Jw?o*2>ahlQFmBFpp}6li<{pnxAi+= z1iA3{NBTEddjFL^y~1G)4NJo!);Ph}Hhhr`mBY_5K=HL2U*uni)>&+r1TB5Pl@qFC zs3$^ps2|XV|IlA*>GQLau+>QUA$@!i$h{aE-%+BR3>gm5kRF+*-v8g#3(eaTZe3Dn z-a4aSb_4W(Yu+xTcYtthqlV%N&0B?e>-z7_TfBbTkoAWlEi|{=N0;L*;>wEl@$%Xl zycAs1kfXzc#Rb|?pb&v@5>jeUj70vQ1)BeFyg`da*8fIx{|m?cuY2SV`&kdAKm643 zvdv(IWE1{Wc=ktB|AwTD$`8%&JuGhx-iWyPg7dAx9FxGHyZ9j;W6dEp zt7WH(1RA3|r&#Y8ANe-n=9A)`bf=y4L~w{xe1Q?scZkjIh!JZ&CEqy=!bOZTIYu}= zC5F8>=i6+sS+pk|ukg>JMow~oQay!vHrC+4 z+|L~TlVWGOu-nJ6-B3Lp&RJ~lJ4e?pRDY*24(hY#IqhWsgRc{W^M{s%qQfnYLoH%$ zx=7&JB#tDAi|`Iq(cw@#snGwEi>7O6?T@xoWHIs9d=6o4*mn0fbbgPn-fa}xu%f;F z%_?iY?q(}`4+PJIBUxyF(^uo6?Smy9Igi@I$_ceG6#qhk^a3KTx*{f#k9n3t$S|=S z%OWD2I)YEMDI^v~U$*r9R*uDnce+Fu>F6eyNOUm?8^rMOQ%wcHgEwHn&C zy={5DJ_s=?qRAL2|6={E#aNF-L;I*Uu-IW` zglMSUWR{C)+aODRo8^!YZL9acU!IgxV0)oucRpwlcqjg~y8rVdGwLe0>KBa4qvWvI z0Rxmc-X%6`C60~Z^qqf5){xjNob+2#+2CKktaMOzY*txCMSbP+<$P=EACwzgB#y`G z-IzZMC$^T<_;dYPD{=<tT~@pD8o94$eRhVwXH7#@Mb@g0;=I9N${Wg(#h4*TWr(xNDzd6;vz9BF zlO4#&6X=1i^+T={(lpc(@tu)_LeO-W!}(I%-s)6nR$(PyXG)ezoYeqI#W3$~{l~VB{CXmal9m%WA02T2fuRsH}R8lxKotT?YjOERjkV`h(%vzEl)7SU};+` zQU|9A@w%{`E4+jAq}nBnYU4&UbuK)UKfjak4eTVk4Viz2Fako{Dny#d@3}a8i7?8f z+9PHUMuQrb;^Bn|WC!LC79Da13&%le3=BPxnW|kC}jvs><(Mc81*><`i-uY@v%W5 z&<&tJgKh_%eHi+n&w+jl`WC1S0ZAT#K4>B67|@$QD?kr`ZUFUv3V(pog_rF3lB=f) zxymCuxgE}}=qt*n^Ft?d-P@a0CFPUGd4bDLW1jT=UOo$Ex%t* zk8a3ieLC_Nz?g_$YshW`{`w%F$_ZUm)@J-QqCAXoA&TQY7aF7E>{C%8u!P@@zxDqU z{AJ)LFjn$F!}`?`{bR_E!`Sgf%Ys(8t^ncm8wfW8A4d2{tAEed_Q^M6sC|e5=Ob;~3@3%!SiD{tbz063 ztQvC1AUB8P8lvP3Jy@u$cTiS0^cC3++nF0_3qAjthPw15JXDt`U()ZmP#o$LtzR0^ zr{5zP4E-6=`o?4ypz>z}{~C4V`-MZ;7@~?bPn%{r z4ZIR*TOEW$<&d3Ok+QBQze8>+$<2t8GlJpzQXJh6nK9AzIak}D_E3EN0eE|JhvMny zz}uS}OM!m?{DNrv#;S1nn}8Q$ZWc%5=S9lj0lWe@^?A6RQT@FS_>I7)wZiq}rusJn z-w!-98n1}hcVaqz2po?KBKFH7cpu13$ABz9@|M z0sc1dcKaXri@-CY^~)pmTMGOw;Jl{|wZ9>PZvy@u@b>Dr130bgf8<|3%G?LM82ISu z@~zlSS+oWngG?P{xDAwvtEEMDoe1Boj`_hM?vOFJJH@<0)7kZ^^5HzBk^!;;4r{U!rlA)bB% zBAVx;RZ%$o#spn%P4q-SI?m#^_tXjAb08V|U5C(_^U&Un_hi6BzvDnLi!N&$DSkZJ zQcL_!EH89nR`BJ-zLYLZJ-*(taK4#9&M=F|vkJaK&tK?`#^}ou6DN3c5qT~#+iIWX zHvKdUia2n^*P~X$t!jv+Rzab9Pq7>sdSV?azl)_GlKZtqw-4f+b<*NDSq;JCUv58^ z8NQBN{1^-Wx5bC*VTt`;Bcb0@YilrIbzqT2r&x5MMe8iO)}mW2y3L|HExOyH`z?CN zqQ@%Zci3tLOvKY9`Ubdw1E zj>xt$jZ;1ifza=TgtYK;GeiA-W-eM$)3Cyyos)&1p2-}vg4u!R56sO9xyfqd>Qzm3GMCU=+5VQ%`^5vHN71mfgVa+XG%Au6mS!;%ttq~T_TP;ARSv+l< z>3Y-Rd8-6I%Nol)7{#{N;`>^#zEWq4$1Q%U#k>7F zF~;JjSbV7di!44=|5X-$fhGTo#TQ$AQCFP^wdY!kpJw65ES|Sc$V=^}0p7|XIojfR z>ju8w;sVvFak z67aJuez>K7uf>P#C0HRkRNs*nKiZOSOQuT-zmi9jGaqQz@s`LRjZEo`i2*E~4g-4N zNLG4mi^19Ak?xaFiceErW7Q|7f>>v4iqY9aE!~bkqX(qBI4CKMdZg1iZ8IJohg;sx zcOlzpyo%4RCTV+w%NA2r(;#9V0f(Z02T?om;ik43-PK~scmf}GlK&8Y4()E+WLR&C ze+4#~vi)-sCdA)f29PQSxAyLI{^+JuV@Y;&>68v?$0*@r6j4MLgR6EK*a3xCn*QpRL zLi!B)=;oPEVuMM=R68v(vl{G~&P3__jHGM8B^r*tgF$Q#2TDxTJ>}j6#9=H%oqHy2 zBPnVB0ve7o8?t&OlcM7$%rS58j!zS8#GFJCj(<>ApOn=EPnd*s-^3$GyI;pgweA%+ zWjN{y?B~0H(vvY4yctQlcpfi~GRn$K=33n5M>^Mc1&MxEh;&|u+bR7JnHbdJK}wIn zQw#6l#1|-ie=nqmbf6*~Kg1)QpZE=>O#|tn9e$6;^HUH9G&a8A$()(`FG2nB&OG-X8cG`5;TDzA*vKH1>4*24=EX3F8HBo5(%~ z$s}r~EqgVRDU@`=+qRL2;BJym)UyJfr-AA=P|ZZA?MxtPWYpLo`|buax-5&(J7uT) zAyKkQ=mJX45!!soYN3sHAwlB-+5JiM6+QW1fxnoC@gb@3$cNC956bR?`rvsY1g;EA zCLtx3y+q~jlHF451dFlr`te^S+yz_cnHfef`f2=^%a5Exfmi!m+LC8&@?{W=rAQfFhuFxA3&C`pwf z?8d3uu9#8krqj_UYV%z5flACqDM}84>(tZ3QCam$H~ch~I$Vy;ni>X&q^R9kLQ)mi zQ>7r?NnMBfoT5fzMNC%@dJ#+1LiBSt^&Y}hceO7U57m?nF{7uNG+BuLYD*b@LQySn zVbxTn*vn<9azA=ojfS;cbq_*lo=V2Z9jOk(@+h@$A-YWk5nINn?x zRTA8JKphX@@w9ppvE>c*5o+?L+7J1+)Nc6apc;x!ds{u8jK@{#JNW8dMbFOOQ*UDI zzpu_ic=$l=fPX$zx1hZrsZYA&0i;@s8FNUP@Wv?GQ--S#9PzZ_nu;3jHeB7t3h|8LT0c#QJ%(%l6d|5BTr<(?7Yx^B=#0I{L|gY6 zt~{u}Xt*vyEnhNRcP0w)Ps8O#m0mVnHxGx)4A-Jm_{?yvh1*{00vgs= zx#}?c`zhB2y=Z8-E*ME;!}TD%v_ZL=P{EDL^)D&Jh01jy`sO0#8i|0iNx5nes~VN- zk_yZ@<$4;mZc?uOXv${gN<`^flL82gx zSOqRuu00q~UntkdXqczBsJ#F^Y|XhZZKKCWTFy zYammX*0&7v=b)?q>jq|VHiCGE_g#T&0GT)7W6Zl4j01~c=K*l}6qS`Y%kS~sA)WOY zFCQUy<^*{2PU-p;O?HZz9l+fs;|TG1M512t>l+~K+a(p5@hra`H5lql9Om>+MSsB4 z_FD&YW>&4c)8pq6}zgru(E53vHT^NO^A()H{xbVkSZCbit? zxWS~>7#;s=>ijL{A;!0wzcRhl^kQHe9WOI`Grim#&h$!iDbs7rjZCjKuV;FL`6$yH z%{Q6eWZDZ!|LvAlfHPSExv)E&W=L*&O)L`*`nnld1kkL z2oKxA`!Gi&B`joe@>xJ=7-bMb2@5&ZdoB=KK*th72@4tLI~Tz@9z&VFFPD_Ckm0_f zu|m)gT}=ojEM!RBQ?N#Z=n6u3H1(*cS}8g{l<;7GJ=7iyAsTk3obE7(V=g&U<~Yo3 zBvPmJg>Gu0Gbendx5Y1g+ctuvK8+NCf|i-1+&?+mhcR&;5- zu(2d`9scM!gGT?eG%UM0fF1j>Os39sP&$n^K54mGYD=9AI3L8ToNbePZ zNGm`PB}5M#ybvkxCx1RxD@_P1Y2<7;hB(1$&tygTb;HitbYXR!p8Uj>CM2ROhtbh)@MAux^{_387q_jios3( zX$k2I4ULlCCwgktG^~`mrM4?!DdlvTTx|``Jc3AFXlSPNHYRATDG{wX2!ScpGO^QX zx~#g8mVBPu%J@u^jEF(jT4eJa zx1fy?gRf#IowCv9p)$xIJ*Tw{GPKs_>0?3vAV}Acp;b1|SijafNRX~!X*LXQvw42Q zMxP9R8!^bxbv91`ULnZs4I5%4rHqR5?G;X3B*pZeu8GqmAxwe&U-Jna8&^Q`W! z@$rO{2Q2%H|J&yIw5!Ie+Sq6ON1Nx?o*Hj#WB*fhRZ4P9l^9c}D0-XX@b0e58N zl^5ICXFMUsGYsL1@S|<)@1P)@9pkxdqQ*@;a*3`#FL#xE|pQeeLk} z0Yohx5mFVl9>vhIIgpf%K)d>$qoX-VApM)y18gPh*L zG_Bf|aI&GteR^g6Vv#b&nY;)=kMxIy^o53soZeil%cLP#L7N?JQ2S`1Y%gY0b|!bS7L|`8S`4jp zdV5=|%D*F8J!rgcb|#O-yr!~JaibaS7sfX^y~|Z*zKQ;Asl!tK?QKu|0`$Q+|p|yv*`?AHvBckq@DTzKctkn67W}CX!G% zJXkS;Eqxsqe>G}CCDpeo$!abO?Q$m1u<7bv)~YzGq@7OhgVqdqpp_9}TiO}beae{- z@7BxtDc^p=m&`jf)?BqVGL zMnZq9ncO|7^=sPbGrrmMc0l)&PcIJZYY)FW1p3#SiJ3)O|G`%J*6k1rU1i3e+Qz0=VfRN#h?s)fg3L!R=UnfxIZ4btBp z);|KL3LE1GO|ODiNq(KTri08!nGP|( zMqD;J=bOJU9ctRK2pOFVOdr$X=Ea;}XkN>7gn1j&BJ%;Jqs+xDH`<)UwAft0`D4v9 zn2s}7G7XyRm`*Ueu-vI;Kc*$-Af^+|QA{VBPqEw-lfPZ+Jk>nN{5110(-~$x%gr?B z5bd(a+z$sCUCPZ&ObVk*r8%1EVsj?bCFYq-mzwLCR+)cedWLx)(=*K%m{yx#FkNOk zv2!%KoMrZ4T5nEZ+F+i^bcMN@=}Pk+rmM_7OwTdjXL`Qb%yhl^3)2mz2ZPn`6$y%%(s|sHIFg9)O2F8HqtLM`!c=U zEM|IzS;q8Aa}Cq0%xjrmZT^GlHRkI~|7L#7^jb3pGsQ^1&g{?hdUG7p8_as9H<}kS z-Dcj(^k#D}(_74DrrS;1bke`g?8o#Db28Jr%@s`VHLqp5(|nle1Li@d51KzS-DUb_ zkp4qvCew$_Ak#<8a;A@(=QDlGyo2fE=F?37VSdc?36oxfHqxIo(`OQW%ACOTX|s~) zGvGS3nOkXhFr;*-XGlS_qb2QTz%?hS3nH!n@)4Yx8%jRCDub3Y* zebxNcqG<@0M*3@JzC~v<-EW?4(W{sqFdwq$n@nFfzqe@0Y&R)1i0z1p4fNDap=()21$Wr@g1hRF=-4sT0?#lymhNIet zFYE=L006m-&>b)<{9ItYhi1;>$lXl{6=)$1&Wlmjo5=k*gvfg$@Y(YH!8o7;=wul7 zqk)$8xpGMkz5zfhJL69`EY)U>oNr*>~AqN2;8J6J^w9IDD^MfpzIlwE! zGHl0xD?tC{&f5@1&jGX{1nX~{{0Z5BuZp`6apxL9w}fCxg^F{YjRkyjY#yH0(~X;a z0{F8e!@gO)1^NC*d=k5XC#aHdhaurU1?@HbGu;A`d=>$n@_K^b0C6qTHDP4gdEkU!k1QJ8~h}b*k*pAH@%AI33?? zmdfN;JL7ECec~N#M zpQ#1kB4PHmZB++Ac=X;mL;LQ>kR5I8+{)U;qU-Cj(>nNoZJmiSx*n9B68t)&D|%DQ z##bR(ghGYUEzf)paz?j7rXS8Xx(zl9KppQRPNq(C@OMKF+F!))=&%M_r$2^39DldG zd>+Ve7~J%Pmml-$XT9RQu4QQ1OgU^#-LiTyF3)v->1hK+aGKBPKFyp>MRcETmN1=X zRxzzKH!xjl-o&)Zda4tC8Dj15p z!jOI8M`~OxGQ%y6KZbR16D;Z`Q%YPQyO--Ok)rrD?31K@L>lI!KNXd<&z1^9=sFnr zJN{zaux=sqB4p%8+y93M4dT}p|JOwt{nm`m~#A!t8eo{0wVz_=U|w~#!wTT)7Yah)*$`*r(h z56u(*o!Cp+uea|@#}g1pyb+dX6Qm@%pGJi|Ea~e>wsH))-;p3SLgY**_;dO0I1R_b z{dk4JLiu}09vW`y;o`R(sD@+RT@aX08Pb3D0*sMc4i>$c~%9H1e!_k93f0 zwfb=>>AzJX^#Ld#qArOQX=zm8Q_?~26%3=!aN7)dE@sF4ur%rF5^dBI?gB$z+*b&y z!zvP*h8p3a4WgI(R71XpwU6p@c^FEgFRl)AoFbkxR@>kxJF8O zE_Nixko|KAlDAVq8XqTR*Kw$Je>5C_tSV_PSolOKcMU<*0yeA_u8nZ%(LevhaOPWcI zt(Ed+_~Bg$d>RtceY8=Y4kx-+%J>!NeKC9j;>2oc-i9K;V(c1IAu%s3#ZC;z()gQ4W4qTIWxFYArJkG*8>vms zt0C`*N6(~g=AaZ`fbhq17s?ZUnj&{g$4>NslM>&fRh*#8a)R2)2^wFVm=;QGLt;J^ zX4n?e%seRlr=Xu`pZPn~<>(%$jfOs!$;+@)ASeZXl7Ei{bdl^a7ivdj^7*G|tT#>I z7vx}afNdqBm+P4HxZ#3+P^u(FN<@9!zc5rz8qbkuWFgHDH_2)P-%HPdIBnpz*1Bwp z4RCF2B}FLjg!H_Oo+fO6JK~~rXji;MXYjrJrx5xQe=!fECSoBWU5vLdA$~`uGs^G4clIp@WJlBZ*dpx#*5J(&lY7$xDy6It|CD$5qYz&?NLaOi&B)MO7$(I{u z-$}4A17a1eMY&`xb(y)C~5$}5n%xy0@R_F)GTG6ie0#z35kVaDK^tPYdqeSmH7u?CE**4J6DyG zCgvr{$qVtZIiw(Zcg7n2vK(@*`D#BLRb*Yl4VNjBYOv8r6k3T(Ydix>! zip+jGEPI`jee<-491FP;l4ZLIFJV-TdKCa4cOfgJZ?DfJxYFzxytwpz+R5vBd9H}DmlSp+5a5au?YSm)$bi8 z-^50i%eTXj8-G^$a}eQNhm`DL=m{?=C}c!%+RnK?SMsj;7CshONf=MM7fy2hsO%h_ z79oEbVHABhY6vY5YqJ+kL)LnLm$wp-dQ4?v6>w$P>^Z2pyaN*Vg``A6iUZgnn_Sq- zs=|xFUL#z)(3}Rm$R@v>q@%Vx3hbLO?jML))oYB5?~83v4sO=)7gfRDbZ@gyj^9ip zunT&?wb&-F3u-w@UuTVIqmT)O7vYc|xPbKs0Lvp>>ofkLO%`Aoq&mpS zz-C7DAI7TTI%bpmQ9j?8)dE`)(P#Wyn_QBH+YIu}CBUu>>nG(<`7*}-B$fzy7xEtt z=~_dC1$;5~D?35pWq@yo1)5<`Bo(8}UHxM0866<-FM!{N1Vr~n>_mE9BuB!<9EIF@ zLY#s>?PH!;JP;E|&~;W!EIpWK$w81B+gg%~(~|3BVh<)+k_#bK*IJU}p_aTMCiWuC zb+-9eNZr(0G7?`CvI&<_En*UN_dJElxeZ1aqh5U>pn59hv} zfmcfTAto*vH@X1J3gPTJsr&f|{N^fi#6E?cCrbt)HT#bw-y&x%b+|6?rUR8x52ni}xo-fY_Gt5;AC)0_uKKP6I4(gIpl4L{% z?M-^;u~k3I7j^gv1qC57DI_nt(+b+VIH{Wp4>h1BmjGHu@Xs(Oy3+=z_xQ9Gsk#t( zA)w7+c$1_}P49_mWmv?x=-YtZ9mYMS#q15?rchGQXBa-70Xne(eVKacS4d7eS&q8G zS^t?3spA%)MxNK5-l4ed%8!G`-t0kd!}Ey&&s%ZiX|ZS`ypVS$Y%Q9BX)~mtAMR}X zA+Yo-`M{0$U!kkV?Xd5LV-3Ag$*!32^gkhmAAE+pE%K~a=2H8IT{jyl`;}}){kM_I zcNi(u)$OX-j?99}e4G61OfC5tsnuXAkXrZ=Ya9JdD2$7g)IQ;T0K?w%jhs0a4?xf| zCuT-G{g1r7HchXDi0wJ7BxG zLtS_{d$*YgMmchCi7w)jXgS$$Ayq?7!4u_24f)R$E%OXy!j26$H|mM9w-iH89hDx$ zuIsAs+GEHl`+#yJeTZ?rlVIs#*^GD;hmjLD8tEiqI~9W>&QvlNMySdYqXomI|B?y+?0;J2@xl; zueL_#wAyOqt(gRYmlU1d;xBpIwr{iigdF{g5jT%SsBI$nWiGOY>z>IwO3Vm@PI>WB zVBuRm5;-f`|5s5iSCCS|Bu98}#FCi!cHEU9LX}JO*(>Rbo|OZjp;FyPonW@uXjBI7oZ8~ z=)Y4Tzfx7`L(D#A!zQe!k^ZZxzrMFgQ-6IQBTf4HJ~a`Jl^L1N15_AjY#&Yk*;~_L z7Ee#;b9W-sdou=eMw#r{UoWchP?Vz=3Iqu=8%U)*I1dj#C^v{){t$&rQB2icKD~n~ zoP_=LGI|M~_EUv!ECFZvlPxZ{6E^#KWfnIE^*`qT<}pd~Qj?KaNKgCSwJ3ofi)Rcf zpf3#^boZ4}kTeU0I_SnKV_2toNT=?>c$OJ!oO?;(9Y28cr)-0Ta8wbICyWKvPArRO zLqo*2B%<#=S6#@h)D=-~^<5bRc9x%fOueJWg!fmpDDi9mcOZmNf3+P>-qkdDzVCY$ zh=I^Sec4vVcLOpL&x@xOmHIe>3tw?MpuZ&28&{iKLq=wAHy^^x?C<77n3-8_nsAAp zn3v%@BQxK90*xa|#}$dv;Y6Zz6p`pci_$SeBXg*m&v0fIxcLlc<}f#(;mjQF<};j` zBiwv;5|j=f8lV)!jLea4KGT^w+RbM=Gsn95OlRgeH=pUu47&MDXXXSqpXtmz)lHM$ z$eifrvz(ce+6S7pmKmG906muFV#aCzdOR36Gj2QmckEel%y{f?qAFSO%p}?mVWW|i zz)XsL9O6KhhnZ9C>DU-#d71Isr(r*tmB`0;``C@nV3Oh}8U=)V{U~&RH2w*3H;>>+ z=z;;aBv98nzQK?mu-!u|2!2_?(F??BTbDsyX6pkDa94I=?j_X9q~n<;eJz>jab)%d zF+96UL2RDQXu-rnteTyMY=ajrk)4Kiq*IB)2=vmO4$7Sl%AF3%oes*K4$7S#=4?(V zIr9iNa-h3{GY7giFwJpqx9A?Gx$aLbN@v%MK%Scpsu87AX(ml3BQVH4g&67?H|;Pe zr;f!2ySG|&C(|MB{T8L;VMZX|ZO0-*lumvbfuU|X=S7rGbs2#IH=W@kN(Z)#z%chM z7JZKC2=@_-(w9L-V1%0vQW=3FH=SrAN=KxK(wQD3fC*+%I_X3_9cwZIqug|ciRiZ$ zrPE19V6?lpMJePGPlt?*z!*23C?ZP7hlt+GwAf7tgN#72n@#}{H4(Utz*u)*i;lNw z71MFA`l3aTGM(h6(lABJI z7=g)dIzmE}&WjjK~uK~!Q(WCW(WyIPb^V;F%M zZaR8Fl+IZgfthYPR6&$ZP#A&J+;m)mD4mTk0<+w75Q1nb=A03j?asF-onkNor@QH( z0#Q1PU7{-nFwadd?-Qk$^o_uLH@%Th zlwQ9#0t?*qu02tDvEB%jy6G)?qV&qV5m@N{t3~N$cq35erZ?Y-(rfNUV3C{NX(viA zup5DLH@&S+l-}tx0u^q051lBzbZ!JH-SmbzQF@)+2rPEfyW>RZMR6mr#2ufeX%^F^ z?nxG{VOr(B*rE?IJ;P1!UK@ck+}|@j(@pPE8-X+3^y0J;sCJLBXa&<{?hO{*&a}q8 z$D;JgvJt3t)BDLp>1AYR_UAqhew?7_#|a95oS^u}2?~Impa|&9{=&!8mlHI9o!QMk zp5UCI8SczJ>f@R&B6gWXq!I^!`$6t290NH)LC~4~ zBZLu}IYF`9nf;TGgEJ>6Iy-1jS!x_6Z*^2%Ml5!I}MUABT8OP>gqG|ALl= z5){Rq*}wXDncxJi6VB}4e7sO_f>sJ=_U}Gk3OGS)fiwH0kApoYDB3&d2mJ@8=-*9x z5>sbjYT^qB_>P82NT=!FM=H-mx|9A*q#b^|pp%-CMvLnskjY6)9Z2a1v9#uNN}WJy z7gpPxQ&JaE`U@viTcUN> zQ+l$_FByyUoW#`<8TwHV0&<^GrI&jm!LZbyUoi`M`3?#0hqWr zDh)qwV#@&+uS|$JGT_=FF~0Fgr>-4`x6ksko%GhcJ%GyQ=kV@I*bkfrW{`t!ublSn z@I?M#W<2(cd@w_pNwL3#3gqWA|ad>GtsjezFBK8Sp;TM+%_=l>=)yHFn_ZB2UuHg zFTqAHe`=f!9^YiYjtrdcx)DCuYX3(On7Q%sK(4Xxgx~V#c~Zb^v;SijnE70t9ro_i zz%1Zbu6No?CV(k*(w*s}_J5G87INKo+b5x~^B37W!@wTl*Z^}w9o!}OEhG%n{QmJ{ z?x5Z0XQ-7)T{RHJ3#u0Xq7E`Y05axLr@$ zvT=8RWFKQ2bfCf#-iMh{FhW04pz{@0_LrDz1;e!s8seNSDTV)41D~gP*6+p^dgrK)BJd)5b$8L0h0B z9T5vw5xasAaOq0YS4z5)^k~C+43vYBSuk1)><9}`N|^c7im6>8;8pw>YEMW&#J)&Q zDbh}%!y>x8;D))#9&0^*IUNzSphnTz5w4i<-Ab*xksVAYNvz!d2tozhIG4_qL>eLT z^1A@)sOXr9E+Ow&C*(|$bl^lchUU14tpHyzU9u56gc8BHq52_IWl;}EJ#BRgPDa?}z zq(rA$EdRcUK~Q+R;%qe)+4LdI;TX=bvoMPa_b4wtDx-rh`Uw~h5SEIiw~h)Q(#pjo zK+k}*h1e#-{;4~TU-ICmNLue|1?(-|2s$~#DprNM_5=2T#^^MS#z=V(&S%X5*h#s) zpi~b!W)m$C`w}&0hprKwz0s8k&cQe=yjN$(z5`o@C6b4(fLC;XQWTNI3O)=Jdjt`! zaHN*%Pg1QP7;4ejRf8ZkTzBgzlA;g90P=WcHn20ag=Jw2oQf)prc-P(IqP9vKRN;w z?Z|IB0((MZw}vngn}9H1___8Iof6XZW5brSreZJ^c2n*kC^?JH5b0-mS|IjSJklzB zLtCbUMUn@U?jk>R!`Iq;Pr@i$YZEKX^D(k)&tNqzJX5(J1*IbBI8s~5*m$bxui6xy zQ)*LF$@XloATAfyEB78yGDRnuBDK{*vC9y&3xCq49u5m=0~@gF6xJ&Dxu9g=U~4?~ z6-0}|vAR3xIFt@|JYu=qeGmB|a1M`lIwqy7Nv;w;iD*)|jwdgjoQl*-_?)K$x{e__ zOcf2?h2UPehz-*Ds%V3&v1$~)$wQEiTt#Wc4(SZ2gMJ3PJ**b)SV?)p^p8fhI+D+R zW-D}1i>(OX7=-P@HC!wm+lmycgN*N8>a>kqV>-?iDci!IbOB$&`gF{zt$wVV>RzWE zSxf>HnIxy8#x*(4_dc0=hbve~l9Z`y8T)nzK$l5=)O|r1qm(dds>Oau0#{3ZV0d$D z0UzBm7kGTiW6 zTOpq!7vID!$RIpY617#t1}R47>i|H9*?10#*ym}mb<){%woTjR8fv`^c)l=Qauev# z+kY62qma-;3(#q}{}8yNBLwlYzNL8 z*bONoIPF*APqP%{VhRC_~-`?M-Jv~8J+6V5S{-LzK00P^$|k{y{t{@&t&U5FKI`< zA?`pYwl0_YXGq5f?LVP<>RBxNLeV$!QdKX+AX83smr}JDV>b0^8eYv%-y~pwsJwv) zpXyoM!8uj7AGKDuCn4skR0D6mD0;r-R_7Gq`*GDBt6;o3*b!gmt1&Zi*Q~xr{PL>o z`KYkE4)Pt;=(%{VuU>**lhh*QcT_>_A(EB781qfpkWN+4Ae^SD@6jurR9p$(K~tB* zW@q*4U?IAw7Z5qp)drO7S98;GJWqXv8PrYHVyD$zT>`X+8a@GEtEmrBdN1`8YTjGr zqm6ykCn&eC8ZifNZKwy)y8i0IPWaYLH4LGj5Wd2LS)|^>vNAw*MCk$LMk}(_!vWm# zsY94HIVugKC|8X_;LcM&Lvo=ShrTFN*P$;KsXOsmt{%f@g|Z{QR;s_iON-UF=<6kF zUM?OdsEMdym0IV<85s36+J2^*iaJ!Q?I}VmQ)djq@i_GwMoX=_6@65vp2wc`EY)E= zUYS!%;hE*?N3^9u{W2Ecr>d(^`bzZ{Jb$+8i14vW^+KPXqxzx;S1T7J&sAr`d*`Xy z*f_3HUt)BfuN)}%FX|3_u2s)0z*8l)9Aj^t`gjB$QK+}j;~UfkXyZn8D(ZQmT56}C z488`XY*N=Dr%|P$l)tJ6P`4&^CeoYL2?Xmcie9|ESS4dFT%z((?p8(bs$8m0B-8H% zZ=Z{~sN87R6>1X3PV^Mw8uj{6A^xU@qu;MpDdl*5Ma_cN_3Bdi zQ&665CyU#>UDVFZuK{K>>jlhv+Q2=Cg#F@>M%y~PIXtA5cjKZ(0320 zD>2$1R4&wfmvX_M52^2v^RPM%{&_@gD#242bq8`DQ~#s3cY%-VD(^-2o~wB%WcsSHk&YuB>pLBlGfi>99tI*58>%0=}`n2;k)c^OK9|2>ZcD@R%`;2oEIQ&^> zz_YC1cixN<`Z?$G7-gS#c!1L%ID>%Xe|KIs1OGAS%*~KqokJK||KNNNt@y4pupQ3K z&MhFx|K!|?{{3fX+ih6e=uGxumUq&C_V=BSA#A^smk+`sjzwP3y4y z+4&g4e{9{#6Wr}KA;!(lC(v&Gnz_Jn?b>P3FRn?IC9chUTHDSr+SQtWfJ+-M&2-0G zdw?H4#K?SY<3hf+AM#Q>jLWNVY39ndwJF5w$UNznTwIuKEw=0(WRYy?c@TP zHeQ;!c4zH7@U@w1{o3BmC=6@M?nW5B0N4ID(_Y(&%SB$eP^gn1+CGD)Ha;~oZ0*-r zI2W3!bInv7Hgm#Q`%`@FU9v!;0R~ZYt+UoP;(5nBXjk+@H!{50nI1qC{O$!9 zarCu)%`EQC&NKMP-(6b<)xi<}0RfFXIAZX|rf>2~BTtSPyh&af8tM5w*uI^yqK$uq z&x0M!^SG^t?hR<)gdVU4U*a6Y_qD%*OD{t|iicmbAH}C9u?H{Dhj^>K%LV~G+rdw7 zv#;gLg&pt43l4(lo<>ak#=nF{j(>%hZ{Tt~|1#j^PY}mu-K~v(0!k45ovnP?0EAPU zHymFW2R07#ogb9EBSH@+eKq})`op|_(4cA`482cn+ox_98PwFD|k+QkXxwSDJNBvLkVDzsudXTfc60BA;<8Ryr z?1_GL^I-)1gfW?ub)4s9e5NIB-q$Q+3=gNTlk#sf!HUy8x_lg*!Xu4;GmY-yCuXje z!EK15eD}iECk`w9{LA@S*1xU!%qTMavl}GCM;I&(l8ZD$E8p-J$lG9_Lal<#>C`q1YEUD@@k(*}P zr-2K@zkj_%*pe|3f?7$14%5+dfY^xNkh)p+XqZ-)yFp7OtS%#Oi+O^ydE> zmG`~sJ+YaJSm!kB#}$=)=v75UZYFHydppQ@^dV0g^Q>v z?OA8N%6(3i_N;T0dY@LMJ?r$U(oS1D>1z6$H>+Gze(v=9yn-jGD&`dpx>0jCcwWJR zC+Zy>hmlLVm{&C1t&7=5SJTKPP?;|e19bFLU1SN>ex^I=d-|ZR3uY|&*_VHSm{;&N zRdsnygK}%Qr0epUh8wEt@|uP_bX~emBfmH5QUGK5Ij?WqG!zRrT8GI__vzG_wd7}Co(rc&!V-L4H}0AyQ0A=ehi(O9 zff&!cM8$tV-B9Y+5xR@Z`Zh@1(@0vd8ZTVSwfm7viG9ROAk91MPb1oAf(71-kera4c zR|9_;3Esa$fe@rte)i=X^BRk`Fy@x?xL9q!g3Dnt)kC_nT?pXMzWhx@kdgW*UV$7bxPY%` zdQ=~;Gk9%RuhZ559Y24JZwj_|N7^T*k!mMB#kcC5cOig3`|>^ahtOaf3H!R=zCmNd zHq!NVzr9BzrL~ieCTdkhNYT_rx7UD-+Mw@tQs`(n`c)V+9{CP2-2;$&G!i|OS)kJQhYmK(P(u;t@k6;Ux+05_04Gk*ycaQ zKxp}cri}*m2e>q%jOaee^zl2%L4#EgGHQ|>GI)_Oqdci6xJkiVkTlV*3f`!0DE~bO z4WTurO)jzfG(y+vxWj(jubQ}@oR%sol>e`dw{o2rhLN@Xb9aFLBUa?6xJb`(v~1!Y zPTyfR`gen5}+S3*7%Wp$yt7HhYy%#T_%&jAGKW4q#<|ZWAAmxJ$_#`g=5DHbF zSv-b#U&FojI7j<(w1@&EJ<#H}OW=8R3tM#yk$vmu5RUNxH0IL0+A|{U%TFN_1=w;^ zCjskR`rt){*X5@8(`cu2>B)5S7$K_Jf)_1WIYH%KwK_es6#9TQt6MAHV zDs-{Vf2s5I>x?|Nu+kW~+~(;SNaQA5{LHLD#`-R@?Z$mC{;I!)k8S(ZC|>_mq%%;5ZT*mm*8G`B zE#RxSeGa&^=Fe@)evwO2tL;l*V;;A^it86q;5YFv^|n(&y>i(4YsWtl!cW*-inQTD zj8(I)DEU@U7pywEliFjTZJ96kQf{@n>2uzV$h8NGVroRFj7~k`tjRwaN0sh%1PtGDK`V>1kZ*vy53YX%(MxM{1Igk|D-=y5)b7`2C++L!;Z zrfh|jxrR;tXy=|1UuN##!9^g!Y6j@ApIQ71h-NF$=TD&|;t$Q2pMCilLR%&MrDI*B zgDo0kK+NRV@;@W%S{Lwv^T)K@r(2V92SM4s6@%j>@`yZyi(kB=4|hgxU^NwexHGbs zJk70>vg;=jDFP+WuxtS~7OOj1Vb}lUC5!K_-9q zeh3G+e;5C1&}{EPx}V}T(pV0qvj!RE2cXoO5xgD$TNufjy!Br6(qE|l9n}2`T#}!C z`5lOe{{1FiH~(4(V*aILGYkN_p9}qzK?cgT8&5^Ees~SLU zYV%R)I+~ee_V3o?{fN%)2|v5T6%5or9zbe)52jJp|HMfh%S^I6mOweMB8qXFov);H zVyijEfaAD4Y0rHjiTpRYlQ4bgFP`YQo_XV7LUprl^}HM+eB*V9-|Re#%=p`J z%>M@HdZKn5QTkBNja= zt;w4)J4Lss<5l0Io2y44n)@k4-+a`OD5j6B_uBspDWey~wb+wsnyefj^gAH#@? zzPn>IW-7UQveS44U-KTX+t@qiT)5xqtIT~KeoH1M*)ZM855+X1_eNEhedGe(5xrM7Cj_whj_AGDs_#!?_&+1`*2j@8R8NDXIruS`XKs6_F>84_rD8|WotAu-|byJnqB#ITlA#A+5Vr= z(?q-41HbLdA4FdJlmw;pAMS4+W=t}(xr;ksu&u27J-okKW8#pmgnjwjNO}7?e8*Y0 zjQ{dmOW86u*=l0+$fcZhn9GW5j~urzx1HDQ9o7}A(MVqC1bVOoUn1kUsHu!0%Q#*4 zqxPQHBFm0t9IyLP`?z|>`MMvqJw96r<`@O5T*AHf$aDMhuOhTnvc?>`3onsNnby5> zpKcwL`^&B0ux~-jY&&a?P}gJ$L-%$Z#Ae%X(L@@NZnj!l4z)PzkW;JWth=ita(Byt zmU!EK&~4{dE8Yg6^~(oaeTbj~-D|9tjwru1+IGCn#zRvF;&vfXbc21%y0#YQrWpUg4le?H1PK znoj%PttQ>Ut;}=1D|Oz`>F38<9NzK^$ilMK$4JytMQuR>XD4EAL-iW`%v@G5-lj6d zvCi?fexz;`_~<$jJWbVgj!Ldb%Us!H1)(+JY)z&uQmBh|D-gH3R!Du(R9Nd}_P_IWE0saO*YXPard}BD+1h(ODaaEyfgD$_$k12%&MQn1V8f^AvJ_1~b z5tk}%VO^p4X|tVmrauY!>oqteB1Q$$MZ@W~>5K%u*w*cx_R+RXTYC#r-&kE0)7G0L zSmzUTO1mFhTHf1Xn*<%{2hj$Nk& zcAbN!Z@c`@EoWTj(5+w`u0s#-*snP8xYvx~zUr0|bri<+Z^Iba-+>Xd$B*J3=x9|V zWG||2?c*?EBldTq5ruV`s&->%Iz=7ss-*#UM*skh?}MS7G5l24_vkJOkMdheqqIU)L3ZdXpza@f4swKyCeKmTp#c@S&_uNh3up^F})<)T;%_`xE7Mr~Xw-xC( z(C3a1JBC7!9ze|m2|J|9W77PN+wbn^22j|6+dAxdlr`RO9Ywtprojp!6hvMcMoP$# zVU$rZlEhLomceRvs!|2WVbq#ZE!wDv1&B0Tb+TK?6-n%=1I(vmE53K&^|ZWl2=VPq zE!ZGjhJrd83+o=SF%=O0FH>alMzrkZp*TXer6+#L?ywc2yh7t{FN%1j0Wf=kNON#& z5YY>NL{2sNOou;rtkK`odW58^zHTHs_I5;6WoHbk?qW+#GvNL6$ip_^0S6!jw!Qca z8Vj|rk1G;DT*+yWOF$Upv+|*%nOV1nKg{VG33yRx_**`hqo+1+usKKHzXwPN4Tf9ur*u>D3bl^>cMAv=zR4DMflCfhw6RIpvYHA zS5(bEaFEh?iw5Fk0I-vI*Ax{VQ^AXTyh!4q5aBg;FdGq)im`dEjRAr;4v>fK>WZkz z3Bq1YI0-WxuXJc(T|dSJLtbwSDrq9a(f0=1CtkWg984VO*XQ04hF@K~an+*GBTwHERz{_D?j^g5+;w8)1 z1zz6C%MvfYfQ$18pWeaCd-?idUY_EG^;^rr-7ny^^~<<4^WuIDmxh1DrRm>!ft+DA zgSodX@X1zd7cY15GQ`V4Uce1l&1d-Z5?)Gto#EvHT$&!@6T7pKUc2^h;nGNt=jJcq z((pB0oPWZ_{Wn~i+_$o~jSyuWjYe+2wxQ9z(dmhvZfJ71IA5^Z8*F#G6+y&DF190VtMkit^teQN@?x|S zigg=eMP4vg&&pV~yTqvZ6rP;p?xyIs z+?M;?YmpUWK8GZEcNZU{A9nAM`$l&d_egXjzIFUu-x{MYaa)|y69T*+TaN*tw)6PK zh-WuM-5%$&R^$l!$98^dJ>zbDoxAN@7?}6E*MHBoKiSairkv+r-3d_4Jm)@CaWC8G zyifD;4P02k#zOT*(bMkD$iR*kklo$R%t80i^X|nD zyN?iHBi95;jB&^LR@8QH_!sw-{(QfyK`VcgRXXm?!+x>kv@vu8oan})ezi7bLh&s8`}Wvv|n#(umQ0X?GhPO7P+qfq`TGcY1VJXeEs}2z=Q9(8~m>X;W_j5C+qCz z-L}62nd)=?wiLP5b~ia=Pc*D`uX7%82uYzJ&ga*(yBklqJ6WG-Cwp>J(Az&YnOy?J z8r?n4@koETGu|lU$GN}Vbv_ktfN?r;ehJuqJDR&qXWzZ#Zm+lpUPmb1@?qD0ok>L? z8BZDKwMb=0&$=6yfQvt~Fmks?bL?Q};fw7mMXR8aao=$sZM5B;(Pejkv{RwlI{9uq zgq)9I$c#I!psRMYTjikmfuS85)PIW@%^1di?nckd5^-ekDugypQqKPZ#2{#mCN}N`s}*sz3we+as}|-MjJ=`cB_$U+#Q{NL%aK&zqg{#xz{r~X!jG?`G5^%gBsBsx#z~* z9p7^|pK*6NEvMako_FmIRflV&JB-_<(`&mOpLB25h^zX_cyv1d8-+z35IFldOK5zi zQDcefu?KI~J@&Jx-LT%h*ZHS&0IlP0W@P8BO`H-UjnBDgQ@1Kyr87RX&R<3C=tFLF z$!)*iUH6o`9`of!Og>BD$i|)5`4TAc2E^SR&4tls;1*qp$ML-dy>0)YyS7U~;lI1< zSrrWGCnF6QOY_CTY`a|Yg^J~Q>7`5x+n^6E&EyNmQiaTX*6KL`Yh9``m@f=@r%I{C z#cau|SVzaE+K2AAqi0`xPhV$WXLtMVg}wc_RF(?aduFqRY$>0{YpRrhQ> z`BZt%e9jjO_>j(p-mP4IK3hmFWW8dERYn`ix=mAsQg$|9u4GGSN2+vws*<0#g2pcu z%K6zsHq)-4fV9Ke<$O9jnLSsr7E-hMw59SJPE}GXvOFTtL9Y5sqLiJ_7aquF@ImSz z1!am$Xr{HCEtOGs%S#S)8h<(_(2MHSgt7`4oCLo;Z<9_INQA(eGKF z^j%TPO&*JljL6eyY;=6$jF*@gpNtJn#>dBkHpC0dsrh{7NPZsqC9{=EzA&p&4;>qr z@J0tt*R5M`suEkUDhpmYn-)MX){V9Ybu4GEwguLl8YFWwT8><32eo1e#q zta$LzP

pY^ji)k7JnWw5W2SH+RKW4)rVC>(AsD`tWl9{r6u9D8mxnD=vZ_aNJFv z@x}&5V;W1);BIfISU_u+(v^7bcr3P@nDmAh$0xfLr1nUa*qoK!Whz=T5K*adp!dXy zWVeR`_OTOZOTd2&Xppl)B|bAU(;bT`oQ;hDzb&skS1eU#me}-UrEjr#7Df8tnktl+ z78emAn@Lp8oea^DHJu+GNB1<7KDi~o){R4dGWDhv57cZ zhL9m|WB|n_rxKudlLnJOcfG}8xe(-9MV&N`fZ~zfB`f_v%mUFWUoNJ5`@H4c*>XBn zP?Rni8y!d-8=r_Jt+VNJA-#aXje=t@or<6IhDHXGNolMP(UH^=CZoyxLbkY6vGV05 zkf=8|0&|a-aMF_FQzJoER4z$ z9`U6#i~lU9Fn|S9LL6O`OvHxbN8&@ei8bV)KZhAQhLOK80c-}Do=?r{L8K-xHHRs{ zN^yGf7>S*Xjr4g#14GA9M*Kb?dhHCLs<-RT_Hre2*IjBJQ(-ztsAQ-i?T6YAAses( z6$|sZS57Twb$g_H_SCc=6XH-W=0OZWfp32f=OYZcZ`74An$JrWY}IL5(bY+@{ifpA}6KXHoOg;%aD z&3N8@2UODsG8s;Z-CiPR?xc2N0+Vgb!a z=Q*Vdm@E%d+5@Pu9+`r@Kwcm$*oRdns{@#sX0sJ6jS8?9VC|eo@PLN zAh5MevyxTe9LSC~c0Ymf1GL&rILfCwO9K%bC^Z7K0SaPY9!<;z=Ce^gB{>;!5i=~%mSnDKa z+j-3M9sr&$k^-Rn664^IgvKGjr!AnuNi#+8Yk5)enJ4#p@sUJduQz_=NMCPVAPvh< z0EG7BL4;;lJM%0PF}tvSO)-b1=cS@U-rztI$aIqeHbrv%qYQ0nlkR*e_kxMPjSGtMQV=2PA-QHZ>VE6#2NGq99>fx4P|3y;Im}M@o46X@ zWHdgOj2}%-o=DzjVYCWk?4>Fmk5pkzgpl@=%96;g%9rYS(dy|5#z}}53^6ycbJ_F& zI2H`J`qSN2u~ryBs16v0fPI!n4t}zvO4h-cs90bi=4aGMIT9biY;`6%85^w+$^#*8 z*N5k-V?!zdJXXvQ@>L`kJ%MV($Gx*9;E{$ThMNd}te_^LU`RvCU+^}TX;VZz5;Q(Tx4Nbbpl^~$MHr=GJ` z0qn_Qg$e<5{YoiuU%E5{wy>b7?>_j!@Q}{X5Q3r14{)#L=E{>>ln-_n(+!6 z+3O_+lO9-)p%aO5Om%>1VOfZNtWW`78D^E_pTcrLLqmuNb~gz*ED4Q;0DwX)XC1Fn zq_mu!&6XC}?vR*^xm}nRBTvcb$i;fuCD3aure-nzNXaQFd8sl10^s-*{8*}#ISURg zNll2z>MTHwG2vvqAg1}27Q<}W znq4G&1log%pQ;t9&I+NipA%CoktUxgpxOhPRiJFhSwx=EuZj^sN!fGh>>?*y54h9Q zSM|DAaReG;K+c5_2>1VU0h%=n0Ieqwm8U{AVO0AyqY4pW9GozQ30OgjE93<^lUBe4 zlDO4KdHWZOkc8P+%T`X{OVW}nrP5v|KMRRkAkZI7sq&b!vJZ;BMX_Ub@&)WMWOE9k zOAxU{u3mj4^Z;s_Th2niA-os+nv{<8wxWWh6INx$R%#9e1tc>tb|D>t?#mirV4fNE zz?sfw75PC$>QGAs$6)Nffbv3!JroxH5Q)#vXJ=FM5Nm-~Y0a|b@+EMoLZHxUV3qV? zFJ#Wa_}H-J4QCO%c;26g^{gu@35Zu@3W5?w^JQWGvKg@CWWH@x|nQ6B9(Hed_9&|N>sFkG+(q0MnDE@X5bu8(H6~0pcc3D2C7{{v0Jedx^da*#%M|!|vasky}1;RPV(od2k7|Dgj;`;5w2#Ge(&< zS=lG^`#f|^*5Xp-EY)3C!a+eet+-*ltDx@C@KjTGD%*sD)tM7}I94f=5KXDFg{nz3 zpmoRq5adc7ArpjB$_&ht`kap!rplnYkSy~ae#9#?)eS%vQ#xq&93ezrXJ8Uapuwq0 zs7&h7x1W!7B-^u4%!ITas*>fYLTVY~gZh5yn!eOb8DwhIOQZIf7@_YdQ4NR*R%s=T zxH@)UY+_uGY0m5eLqp?JW0M>MAWtwduuGGbUZB@1TsTm2Ato59lzJ;^IkW*GHi@)@ zN^@2Oiv#qOFvj_SBp$L@8GW8jXERt3p`=BK*aU3|<#HhzP;0Uhq1ERaogMJtCP9<&v@hcnqq92(H< z0_Grp3ar&9l9dTK#g(HFGepu`DrW;}74R`rM-``7!#IkLKbM+EQeVp@yd6feC?P2{ zDMch8kTA+%sJI>~LZe&&Q&E9-931{^!MaeRNhFO@RxPiDVa^eft3Yooy(+o^Sk0gs z*xRp=MQaJiCm@83j~%f@slrS|op)t6Y!#gtfl!<;$5ihtiE;`Iga<`ySeI8jQub;o zV+%UH_NR2LWe!Tf`_u$1&eqFl<-bQ@$BI|s^8NX@HkgYeP1 z+)?casnNiRvazq1)-m!7zJ9%mRrEtcTKAw_Gz>*Af`uz5j1rfb)zBbi7^)|_?oIoo zn4+Ud(hI2qW~Yo-s(5GfP$ppjEMX=CDTJ7gc?`5q8NM>K&H~Fw)=mzLOkF`(-w&ja zAsi;)#&mmHsPyT?(O${Uf92YRbmE0EF=(nYBC6$z4t3-b-+G}Kt2vQ#wzri%-U7-X3Jl_3h< zj&eg^}LqIC6J7#kXzDx`}ga3tbR0eY4y&74Q3s3Q8b z-?uc$q^0pw=wpvJ9JbQDs3E5<$vyi@S*2$6DKru31;igrtn65!8_=VcE8bEs1?GB8 zGfmh^rzPqvYH-=E1Aep6A|Du9Y#Yx3!Nk@oal_Jy7EVwBvg(UuNTIN;fHy4%*%-s% zg|J0E_&g-o43zM7?&Aeb{#aS+9mwHG+5I_ix|${8BWHk>y+Q*s=cbUs$QwR7F)%8@ zNpQR1*!9$6iOmn$)ETk8^A*D!`7k-2udu>kzKO`g5OL9F(k$UcU!bR`Mo*thg%mfO zVlg3978dEa5gPR&bL_;xDG%zD6R}AM7ef%cCIjHX3_3JVlk_BMPyzHu(a)em0SryJ z;LW4l6Ap@@%MVqoV_)HWK)ZBHs)ENfm? zJ(I$vi%S?+kT&ZRA<=_`6?`N~=0*eq&t+0{QqklvxZ9CDoVy=tgrFQzgh7 zp{YQadq2Ih<%};?#&Z+da&f7YhFr$xXgW+m7h=G`(20TJ;R*Rt#T^Ekzxen`UrCGQ@Pblc#!mxVtM2c;1ndL?s~+xUz8FPV@a=uagK>$pEa$?yIL4cdK0n4 zh`})9c(aGHRv!?!QiPRYsQ}JB(`)5m(>*Js9L@I`wJP>)A8BNj)DOAUhisx$BwZz!@?m+U-k&RPnv&6{vhrH=x4lL3lA$3qHOPMs$XC zfPk4z4Bk{`a<+mEyxj=Y2nkF{qKV>kR^{arHo^zZtKxIw<)R|W23>ptS7k+r1y*m< z%@DOQ{(BXnlgFH|AUk-lU6D(rND=q0gM22v1n*t}Bej4j0ICgReAjXlCZ@gfC6G8M zaNu6ENCTQsJt&jn&`FLeu9p*7@BS*BCu=1hD99GZ*3jrSq*b#v=nG|J^M(f|2gGX3 z?$jv03dZ<&M2276t2Lt6a=nTh!%*-MDp|}`Xjp+Pp3lGlgh3kvYcK;-zHpfLjVX-o zRjy>D8&C{rLU1)oK(V@2Tk1I+=`&LYtH||Vor|IJkp$jaxg7K_1QEzSJn58c#fm{I zlTQ{9f|~)=8CUDW2(1EjtYT@%iJnxWV~sS>(-X2?poMu+!Jtn4a9A~{$VX1M>>j_{ z!oxqLYZGA)aftV}jUM&Br3IJ=Ng>XN`D4W*%IqGW)Jyrh&p;5pEXbaNXT z0NX^x9jhgw@GGTffjny-IIehLH4XS^qh^)S)VX08tma~{TS4EaIF&v0MDtk#K68*@ ztSJ9k=nK#ZnnFkr#TKon>_c=lYWbugRYQGWq{1HlS#+bVgKN-?%j5v^Dk7~#hZ~>9 zSm>6()l(>h7KrK*4(S~@V8Gus7+ds`hBF(6)F7#b>OzT>v*@2`X&)xb!}JW)0s|t> zhLQuUm~d+wR$5GeOANJ{rG-xazqq>25KRQJO4jk$k9g2b7>m7j?sW*Yc zF+GneIf_+Kh7t?#d0_ZtJW1d2_}Gzg^4D}@2+4Muu^31{xRj^ECpZh3qsbx#;L?+p zMsF+}$cT@5ogvoKSDr&Pp*t^vS)=D1=0$j)${@frTZV!Nbc7!^{@9L%#??o*tq zqFXi7qo+7l(f1vl7@taj8?(g6dPQ}U!ux8t?)ArcH80V(FFaKhAoya%hLTy+#t|A~ zqGl*VnHYo2@$o@wHYX>>M=WClClFV2a?tguI~lCzvhIt}tVY-boG9gsB}JNzBQ5*9 z#-WuktEWrohrW}_%)Qzl9m`wbbq02*9`V#%Gv{$_Mnp8&6rn67tplyVL|(|iNy0Z? zu_{V#@FfPHp!n*MI-U<=cqIF7^o`Io1{%i|{cNP1YTYSeCK_x}4?qZ$3I^hf6jY__Cb-plYyP?Q1AkRu6SThB>ZyNQqK zfkjTCfGI5nU0%t`WS0}l`D5iAve)-+^2?{QP&KdMe;O((dI}?h!d>5MC+9oa2<|Ay z>WT5espO=X{R208@s<_e+gae>BI!{&i=|8yxX=s4Je7Jl1}+v`qXlNPA0j0&3 z*^Fzu!hTvJ5rK6&qF5(FNoMsIe zdRb{Wh{2$7kTf!~+$!krki+0e9_Ew)wKP^J4p)NdL0F?BLcUc_Q$fZ+q$FMKOTZ4` z5EM*`k(Y|Vk}H0rr7R-A$R&D51C{`_FFzzSM(5D6 z_{gwOKu$d3pD)y{K5?GB0Wh~i=^nwwx z>QaF~9}@2LBiVgxr(^Gkaf!} z-=HfCuuw4=uRH&&BB0FF%XGBC@)_qa^ei`FBZBh5b#<0pkLA%B2tprxh#7yz%U zq!s58HH#{Lo>46H!gMMXRG^(KCWU+)L>?h8P^v{`agZ3-#4S&%+0zqRRH$vPzT*`; zPvuH8@MKXTq(c1~t-y;_7lP83M77IqkOF<1y)R^9bZgC8j;obOsWIki3N}NI&dAat z?Pi(HYB$SdB0i=mETliPB0@xajA9gDZDOwW!BHNq;0dvKbzHdG${h9u0H|ZCF%iddA9BJeanJGKPG#1euo0*~H+29vr@$x_HYs|4 z-&x3C93M2h$J1JUA18f0I0EZpmDDu=iXt`wxCi(gwy;v7=V#rjbbK(I8C{qf*s-x= zUc6^N(Iv#<*Oo!7a+lFpXf8+X&{7oC^c$k3 zZ*eWDlDx6XNVCFxQdD(H7M`-#iDZ*!5}1hemLubX$Fb&=LO3g2m(8TGsS?8UJU+$9 zRSx$7bz$c~&npHk6r}KFYAlX&fpzr5m%i%9P&9tid_*M&N?2`AyPQ~#5Sbf;7`GzS ztY**!@~`5~>MggB^#|-IF$zKaMC~+|fUd9b1|$|+Bv`9hcMrlDpLh4 zQx`2gysfzRkgi?b{tJ3lJV^bG0?Wrp)}ZKA)|vU9MZ8Lh1qQ+1B?(` zT^&)PH^rTL1SBQ4n1YzQNh12PeCT2-I%Oi@|=0noUvm!jzJo(jC)>tBFAyAS~$eabd z5*rTW*J>IRY)yqIU>mVw&exP+M;9_G97q^)ib=)F5P^dbMYP&qfXPOPt6=@H)K_Ix z<`73>@^sv?%D9~RqZJe!OXh{M5s68WQQ0vDjGKnXRq8}!&6pGx+AITlLoBPJ&IAAn zTs*~V+xU5^_SE2(@AD*!R;3kUq9aGDDK&RdDU(r=T8EcXK*sf(7)+Pi)EJ&J_X-nR zs8JN1m=tf38n}^<3<WMowH~96aERgKZ*hb=a@=t_JqD*Vj|zUZV}_Vh6j!X6?my5Y5tO3Drx zG}Q2G8wN@paKM!m*rb5PKPY-PMK0LA5U7Say|4XCkbFAz-G zW9;pym}VWr%kx631u<=F??p)HoZxD3A5G||>j`NB;))GRsk5rR7sMHaQ6S1Fr9+Qc zj>0Q~vCb0PB5?IUJ{Lb-#?FdddYOm%3M=JOC0Y3ln~#v7*m7(yWJ5@ev<>X%)^POL z3QV#gf=TdK@0X@MKs2_ar?Bocw50uzUeW>yPWleom4OLat_W420wOFR(3#CIEkHwp z)K}ht3Ylmk%35&`5haFj{~_dpbv)Q&v!Hbhs1(}~BG6ceEfZDHx!CGzBtKJa`2(= z-*zwE7>WyCN(Av*#9m3?#68?-Wvll63EF2=4Z2myN|Ed4WN`Gu#+1t@)JTD%o!c@& z4Z=MNz`|Kz?x3ZG%BaBgSUHhq#sK@;C@TQAr|zk#^hD4 zTO%LGwQH1?>6*60JvT6Zwft1I86~2Kk=FO;4#d2$wih&5h^?#^lhH;6lk7po5%0^p zXa_fqv%+1fCf2K1HZ}m(7aWdg+9wu4irKunZKy~Q#T8pqM|lo;Hz=KY981m2 zQas)5sq zVoijOn{6X~wQao?!IvVKI+Zt*(ign&h8u3K_Xq z2OSwC_d4^Z>R&BC?h#ErbiRI;1B-cWRUH)H0U&Lk&jRs@;w=?|p}noN3dpEs5dnTV zLQJGe11CcUb#z6UYjhMe1~DD0%_aN_XzWSk(bkhnATzVJzy4G&!IF8Q*K8%jH4MtN z82zew9_o>01!o`lCu9RoN03jP5rtn2^JuH=$%V<*_xXkK0qdxk8%;e|wDO=hy?Vf# z!2VFmI4qQ|W`OAJ(^)^YS7@zj*@w;8GPet_UcSp@qf=OB4RwrcI7)i4(%v4d30lT% zSWOOs-XAHJ%jaLZl;tjk$T!$Zu6j~lb$KOgnp_f=gVhKI7}uf1&YcT(L)7aug;ej- zv|1Q!*_LPJUY}A&j5RC)I00tCymF@>A3US^bJ>hCnQ?2tCB5Yn@J%mdRNAzW*wKNZ zGtxk1y{8ZEi69#_a~%QAPC|gmW-z>hj|M#}n0o;sEntWy%xq!Iy%XUV!R^?MuJPbO ziCsmrf^~#ry63nF#Q4|<3`2FsEq2!rNRiK|fl{ujj}bk2wbLs&RLw*9K%-T}XxYdQ z`8(c)IlCNo^blnVy&~A|W8)y>zR^@sdp+CKYDC3*tcaA zo2pEXk6|}^*}T)ImWsR&jhyf%x$tJiA^|md!D*Ranpz%Rjbo|cq)kk^?P}(E(Pk1BO)%I;MGlAA6z-!IlZLcUhlz|~q z7Gmwx3NBhW>BXuojff7toN};uYB+WTfKXOVwa!KDql#5GYDEz}GSr;LvRN{I6=E9uawUWD;&$s#?hUOA6y+$%%G4qj>t~JY5nC@4c!H|I z{4(^IGHWTWN6%0pDbL@1zfYJFBMC#|BGy5m@P`qZ)+?m1vj-n>(?<<|8NaxNDyL!p zBNx^Ce#dMI2D|@Jk?2u6!XqzXT`FF0K_)TnF?c zJeq=YT<07Nc8lk76|cLydrxn=dp{rKw|l>b{YPln^Do#AQBX4<%%NJpK*lTDR9&quwdbLwIIBe11lf!k}N- zjem{^EaU^m&j3EqEX@#7I7X#{ks=;X%7}x`Llv;4J9Fjf4LTjlk@TGKm`}cFrIn37 zm?EJzVrQR;d8x%bk^_Oz z3_Pqvace=$#L(xd{X7s!t==)*2RfO160;IG#FU4m^4DZe*8H^%P3CCBJx%8P!i`Pl ztinxA=3K%HP3E-0`J2xU1M@aEnZxlmHJO9(UTZg*W9(v0=3u(}n#^Hy6HVq2x!q0XxVT+SL;6^_t|oK# zTW^y&;cc|ZoYuCx$sEbn)?`j#+tXxDT1z#V^V1$^GAEyHY%<53-PvRgGP|eA98H#L zx<;QX*4s3opN7;yV|A|R4y(cX=bzPnYDR#KFK@z`to*gh)v|B~Z{7GeB0OTPw;p4J z(61fO41c;NyibQ)`Y2tKB#sB$y3u+pr9xOg{-*G6TU)JXUoNly{$7uG{QV97S~9Ih+BZ|eAi+tm9(9o~*JhME5b9j+0}8bCPV@U+gy)Mr-5e+7;g=TG0P z*Gv3st?2<3evP_W2L8{~Up0@6+v_)ES*dINSS`n)r|E_(yI~5uecUKcd5*)8Vhw;a`$`Hd^o0 zcw*~t9QKF&H(Kwi3E!i`-(3^VVO0-*<|O_m>(1jk#W6rP((bw5#fZx zW4gmP>d%i!{?}Vyd3kkw9{6lqH&{P34b`823*jbhIK}@!_zvXn(|`T(FLbz1?{)a| ztK`$_s(gIftm7X-_(tUayBbeUX+SSXc(^9&`IiyC9^>@bboIBl>im7mt3Q8Chx@aG z4*&Kl`M-0mWwl#Cs+#;aG+NdMLP+;7nkm0_gx7;x4B>=ZZMkP9Tqh3Q-lF6CGm-xG zc^&Rgk~;i*2q)a0)^y`>U7;W9_%rP);wBycmL^pnf8N#MlL)U@{|gA;5ZPu$r*%Yq z`z4)^Ka=b5|ER;1Rr0DarS@aiCW z>s%%L1i}e_e~#B_Q>(-;uM++nIv;=L*ZDlF!+kzLhu?sSndSN%kq$qA@Ot&1(D7^G zxwuOF$AO>h_h}Lq{BBh@>m94a|EpEPzq?Afvqsh5=a+Q;ck6JUE7IX}I=mJ?--dA3 z^IA<;dv*K|>G-vr&)ODMuFo6k_`48Z4{jr?gwHZO0wAr|8R*->D)Apf_=ZR@@0(PY z5zc<}nNI!TOFH~r9V%c#hyS|{ukH73(qAUm;Jy#x8!%73W?F@<*K}_b;VjqZnRSx9 z4)@t%9sYKNudhR|-Y?-dT3=7;i2C*)5nd0T&+7cowX2A))$xC*^D*_*!B#`7s;AFQ z>kr!yUN8S+I=;`?>-cjz+?OSE_@WN?d2b#59v$xU)jIssI=ptC`Xa*X)#o2{{Mz}e zab5lRJ*$MDS|$7(!r70t^zd$Zv(@|a(!t-7n0ZEzVxQU zAJ*Zu^UlX~xG%-&_weRMC_nu+p*-)jQEGZR!rlC9nZ%OFjOqV zTuKTz^@W9={vf8+jdf|@9!t3>^v|K||&?Q*u7~cC_f#^k@X&LJS3J_|0S$24hE>0f~HGk5R4LD$x zMHFtPM`MrI{Ot0T~&SjHEgnz1ZD&|uvBH{^>+4lVtLk! zGd}jvFn)z3Rqz;QY;h7ml9kRXvht8Ua1|#X0VuT&$YyIFS^-x**RvXMUZ!01%;D_P ziaJifdg1@7eLmD-kX! z#r;#p^h^oYx0TfdsKnjA2EA5h=O3a~XEd#TJ{UL{l40$#Dcc@Z4;*c0u2ZpmVmDSFp)w@N1>gmR>V?+B@U!l zo3XEqusUwT*G=z{hn*qmL2YDEj0{SklU&YZc~b?LfHNN4y_XiUhmQ^offb?X$qLpz zIDKUus!__YL~wat85RBT|90^Ua)4M@p>=c227= z<}xV9yrBZ=wDxXa* zSFf^wM?a@#X0RuO`2y_;idL}ze)Gs+;EMlVD$H`JPC2Z9`gxzJp0@E1Go$>r zIMsJ;UesSzT>k|RFTp=d8GicxI=#7TbCI6}W`nBVUc80V&*}8$esfi9+7hdO@tX9; zAHdw(tD~<-e+u7!@Ux#4P5WsN<x6}l4s58}hmLiHbpeSts!u3r86n{v$k0>1xyDE*^4y}5r~ zr)RYYdsBZMY`s~BoBZYo4E)SJ!A9d(3qRBTx8oc09b_c@J^_7u_ZkA8{GtO_;?HG#38!z@>CF8}U4N57P4?CYYSKTwLB%w86MwbzAE`DUx*fA?{nUf)UI!oQm8SQ`)kQ;5(I#xwpzPmicnHQ@8dkEzFYq$i9F znKAxL&+7Chp2{s4e5ad|S9JPz-Laai`Q@8>eF$mUCX?RyW&ChVl^Es6 z(9e`>?!EXPPH#=?$O&~fbi!m|(wX=DHR-oZ>vTGOw~46lwdvzX&oEQ}*G#L3$8`E> zRxb2A9D3&Cm@dDz9FzWzn)oI=bEm(Ec@9S{>+&Bn151CbjhkB~y^~Z4u6tED*Og_J zR!P6*k5sywH9a@E)RzC6Rni~)OO@htdOlh${hM_9TKHMtQYoI!n3U>n$~X77uaf?0 zb}W8Te`2Yh{sTI_`WjrH8CU7gshj_y=2iNxri&!*VLI61*RBchPCT=X2LB9R!;Qhm mJv{oq=~t@suUzj3$wO`VoNf&6hGTxZpkjW!CPB@s^?v}UxjWDR literal 0 HcmV?d00001 diff --git a/example/gpu/perthreadRuntimeDist/vec_add.cu b/example/gpu/perthreadRuntimeDist/vec_add.cu new file mode 100644 index 000000000..4413cb6ff --- /dev/null +++ b/example/gpu/perthreadRuntimeDist/vec_add.cu @@ -0,0 +1,85 @@ +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +/* +nvcc -x cu -cuda vectorAdd.cu -o vectorAdd.cpp +python filter_hashtag.py +g++ vectorAdd-new.cpp -Wall -L /usr/local/cuda-12.6/lib64 -lcudart -o vectorAdd +-g + */ + +__constant__ int d_N; + +// A simple vector addition kernel +__global__ void vectorAdd(const float *A, const float *B, float *C) +{ + int idx = blockIdx.x * blockDim.x + threadIdx.x; + + // Grid-stride loop to handle cases where total threads < N + while (idx < d_N) { + C[idx] = A[idx] + B[idx]; + idx += blockDim.x * gridDim.x; // Grid stride + } +} + +int main() +{ + // Set vector size in constant memory + const int h_N = 1 << 20; // 1M elements + cudaMemcpyToSymbol(d_N, &h_N, sizeof(h_N)); + + size_t bytes = h_N * sizeof(float); + + // Allocate and initialize host memory using vectors + std::vector h_A(h_N), h_B(h_N), h_C(h_N); + + for (int i = 0; i < h_N; ++i) { + h_A[i] = float(i); + h_B[i] = float(2 * i); + } + + // Allocate Device memory + float *d_A, *d_B, *d_C; + cudaMalloc(&d_A, bytes); + cudaMalloc(&d_B, bytes); + cudaMalloc(&d_C, bytes); + + // Copy to device + cudaMemcpy(d_A, h_A.data(), bytes, cudaMemcpyHostToDevice); + cudaMemcpy(d_B, h_B.data(), bytes, cudaMemcpyHostToDevice); + + // Run the kernel in an infinite loop + while (true) { + // Zero output array + cudaMemset(d_C, 0, bytes); + + // Launch kernel + vectorAdd<<>>(d_A, d_B, d_C); + cudaDeviceSynchronize(); + // Copy result back to host + cudaMemcpy(h_C.data(), d_C, bytes, cudaMemcpyDeviceToHost); + + // Print first element as a check + std::cout << "C[0] = " << h_C[0] << " (expected 0)\n"; + std::cout << "C[1] = " << h_C[1] << " (expected 3)\n"; + std::cout << "C[2] = " << h_C[2] << " (expected 6)\n"; + + // Sleep for 1 second + sleep(1); + } + + // Note: This code will never reach cleanup due to infinite loop + cudaFree(d_A); + cudaFree(d_B); + cudaFree(d_C); + + return 0; +}