gpu-sync Lite

This repository contains code for generating a shared library that intercepts calls to the NVIVIA Cuda 6 runtime library and redirects them to your own locking code.

How to use:

1. Generate cuda function stubs through the GPUSync/Wrapper/make_stubs.py program. See program comments for instructions.
2. Find cuda calls that you would like to trigger locking and insert locking code (described in Wrapper/GPU_Locks.h). We recommend:
3. cudaLaunch(): acquire lock, run kernel, release lock.
4. all variants of cudaMemcpy(): acquire lock, do memcpy, release lock.
5. all async variants of i and ii: acquire lock, do operation.
6. cudaStreamSynchronize, cudaDeviceSynchronize: release last acquired lock.
7. Generate the library via GPUSync/Wrapper/wrap_generate.py.
8. Compile the library:
9. Use the makefile in GPUSync/Locks/Kernel_Locks/ if you would like to use a kernel module.
10. Use the makefile in GPUSync/Locks/POSIX_Locks/ if you would like to use a runtime implementation.
11. Use the locks:
12. Compile your CUDA 6 program to link the cuda library at runtime, not statically. This is achieved with the --cudart shared flag to nvcc.
13. Use the interception library: LD_PRELOAD=/path/to/libcudart_wrapper.so ./cuda_program