1 / 16

ITCS 6/8010 CUDA Programming, UNC-Charlotte, B. Wilkinson, Jan 15, 2011

Zero-Copy Host Memory. These notes will introduce : These materials comes from Chapter 11 of “CUDA by Example” by Jason Sanders and Edwards Kandrot. ITCS 6/8010 CUDA Programming, UNC-Charlotte, B. Wilkinson, Jan 15, 2011. Zero-copy memory.

summer-sanford
Download Presentation

ITCS 6/8010 CUDA Programming, UNC-Charlotte, B. Wilkinson, Jan 15, 2011

An Image/Link below is provided (as is) to download presentation Download Policy: Content on the Website is provided to you AS IS for your information and personal use and may not be sold / licensed / shared on other websites without getting consent from its author. Content is provided to you AS IS for your information and personal use only. Download presentation by click this link. While downloading, if for some reason you are not able to download a presentation, the publisher may have deleted the file from their server. During download, if you can't get a presentation, the file might be deleted by the publisher.

E N D

Presentation Transcript

  1. Zero-Copy Host Memory These notes will introduce: These materials comes from Chapter 11 of “CUDA by Example” by Jason Sanders and Edwards Kandrot. ITCS 6/8010 CUDA Programming, UNC-Charlotte, B. Wilkinson, Jan 15, 2011

  2. Zero-copy memory • Zero-copy refers to the GPU accessing the host memory without explicitly copying the data from the host memory to the GPU memory i.e. zero copying • Depending upon the hardware structure the data may get copied though! • Integrated GPUs that are part of the system chipset and share system memory do not. --- example MacBook Pro • Discrete GPU cards with their own device memory do.

  3. CUDA routines for zero-copy memory Use page-locked memory. Allocate with: cudaHostAlloc (void ** ptr, size_t size, unsigned int flags) Allocates page-locked memory and accessible to the device. One flag: cudaHostAllocMapped-- Maps allocation into CUDA address space. Device pointer to memory obtained by calling cudaHostGetDevicePointer() – needed to account for different address spaces. Reference: NVIDIA CUDA library http://developer.download.nvidia.com/compute/cuda/3_0/toolkit/docs/online/

  4. Flags continued cudaHostAllocWriteCombined -- Allocates memory as “write-combined”, which can be transferred more quickly across PCIe bus on some system configurations, but cannot be read efficiently by most CPUs. Use for memory written by CPU and read by device via mapped pinned memory. Combining flags: cudaHostAllocMapped || cudaHostAllocWriteCombined Reference: NVIDIA CUDA library http://developer.download.nvidia.com/compute/cuda/3_0/toolkit/docs/online/

  5. Code to allocate memory and get pointer for device int *a; // host pointer size = … ; // number of bytes to allocate cudaHostAlloc( (void**)&a, size, cudaHostAllocMapped || cudaHostAllocWriteCombined ); cudaHostGetDevicePointer(&dev_a, a, 0); If needed Always zero currently (for future extensions)

  6. Using pointer to host memory Simply use in kernel call where one would otherwise have used a device memory pointer: MyKernel<<< B,T>>> (*a, … ); without needing to modify the kernel code at all!

  7. #include <stdio.h> #include <cuda.h> #include <stdlib.h> #define N 32 // size of vectors __global__ void add(int *a,int *b, int *c) { inttid = blockIdx.x * blockDim.x + threadIdx.x; if(tid < N) c[tid] = a[tid]+b[tid]; } int main(intargc, char *argv[]) { int T = 32, B = 1; // threads per block and blocks per grid int *a,*b,*c; // host pointers int *dev_a, *dev_b, *dev_c; // device pointers to host memory cudaEvent_t start, stop; // to measure time float elapsed_time_ms; cudaHostAlloc( (void**)&a, size, cudaHostAllocMapped || cudaHostAllocWriteCombined ); cudaHostAlloc( (void**)&b, size, cudaHostAllocMapped || cudaHostAllocWriteCombined ); cudaHostAlloc( (void**)&c, size, cudaHostAllocMapped ); … // load arrays with some numbers cudaHostGetDevicePointer(&dev_a, a, 0); // mem. copy to device not need now, but ptrs needed instead cudaHostGetDevicePointer(&dev_b, b, 0); cudaHostGetDevicePointer(&dev_c ,c, 0); … // start time add<<<B,T>>>(dev_a,dev_b,dev_c); cudaThreadSynchronize(); // copy back not needed but now need thread synchronization … // end time … // print results printf("Time to calculate results: %f ms.\n", elapsed_time_ms); // print out execution time cudaFreeHost(a); // clean up cudaFreeHost(b); cudaFreeHost(c); cudaEventDestroy(start); cudaEventDestroy(stop); return 0; } Example Vector addition without host-device transfers Note flags book seems to miss out this special routine when using cudaHostAlloc

  8. Code to determine whether GPU has the capability of features being used Look at device properties: cudaDeviceProp prop; intmyDevice; cudaGetDevice(&myDevice); cudaGetDeviceProperties(&prop, myDevice); If (prop.property_name != 1) printf(“Feature not available\n”); Returns device executing thread Returns a structure, see next Various property names, see next

  9. structcudaDeviceProp { char name[256]; size_ttotalGlobalMem; size_tsharedMemPerBlock; intregsPerBlock; intwarpSize; size_tmemPitch; intmaxThreadsPerBlock; intmaxThreadsDim[3]; intmaxGridSize[3]; size_ttotalConstMem; int major; int minor; intclockRate; size_ttextureAlignment; intdeviceOverlap; intmultiProcessorCount; intkernelExecTimeoutEnabled; int integrated; intcanMapHostMemory; intcomputeMode; intconcurrentKernels; } Properties

  10. Checking can map page-locked host memory into device address space … cudaDeviceProp prop; intmyDevice; cudaGetDevice(&myDevice); cudaGetDeviceProperties(&prop, myDevice); If (prop.canMapHostMemory != 1) { printf(“Feature not available\n”); return 0; } … Very likely as only needs compute capability > 1.0

  11. Integrated GPU systems Example: My 13” MacBook Pro, 2010 Zero-copy memory particularly interesting with integrated GPU systems where system memory is shared between CPU and GPU. Increased performance will always result when using zero-copy memory (according to the course textbook) CPU GPU 2.4 GHz Intel Core 2 Duo NVIDIA GeForce 320M Shared between CPU and GPU 256 MB DDR3 SDRAM DDR3 SDRAM 4 GB Main memory Intel Graphics Media Accelerator (GMA ) shared bus on 15/17” models

  12. Using multiple GPU on one system Each GPU needs to be controlled by a separate thread: Code GPU 1 Thread 1 Thread 2 GPU 2 So need to write a multi-threaded program using thread APIs/tools such as Pthreads, WinThreads, OpenMP, … .

  13. … #if _WIN32 //Windows threads. #include <windows.h> typedef HANDLE CUTThread; typedef unsigned (WINAPI *CUT_THREADROUTINE)(void *); #define CUT_THREADPROC unsigned WINAPI #define CUT_THREADEND return 0 #else //POSIX threads. #include <pthread.h> typedefpthread_tCUTThread; typedef void *(*CUT_THREADROUTINE)(void *); #define CUT_THREADPROC void #define CUT_THREADEND #endif //Create thread. CUTThreadstart_thread( CUT_THREADROUTINE, void *data ); //Wait for thread to finish. void end_thread( CUTThread thread ); //Destroy thread. void destroy_thread( CUTThread thread ); //Wait for multiple threads. void wait_for_threads( const CUTThread *threads, int num ); #if _WIN32 //Create thread CUTThreadstart_thread(CUT_THREADROUTINE func, void *data){ return CreateThread(NULL, 0, (LPTHREAD_START_ROUTINE)func, data, 0, NULL); } //Wait for thread to finish void end_thread(CUTThread thread){ WaitForSingleObject(thread, INFINITE); CloseHandle(thread); } //Destroy thread void destroy_thread( CUTThread thread ){ TerminateThread(thread, 0); CloseHandle(thread); } //Wait for multiple threads void wait_for_threads(const CUTThread * threads, int num){ WaitForMultipleObjects(num, threads, true, INFINITE); for(int i = 0; i < num; i++) CloseHandle(threads[i]); } #else //Create thread CUTThreadstart_thread(CUT_THREADROUTINE func, void * data){ pthread_t thread; pthread_create(&thread, NULL, func, data); return thread; } //Wait for thread to finish void end_thread(CUTThread thread){ pthread_join(thread, NULL); } //Destroy thread void destroy_thread( CUTThread thread ){ pthread_cancel(thread); } //Wait for multiple threads void wait_for_threads(const CUTThread * threads, int num){ for(int i = 0; i < num; i++) end_thread( threads[i] ); } #endif … Textbook utility routines for multi-threading Found in ../common/book.h Provides for Win32 Threads for Windows or Pthreads for Linux thread = start_thread(funct,ptr) Used to start a new thread Takes as arguments: void* funct (void*) void* ptr Returns CUTThread type thread identifier To terminate thread (join to main thread):end_thread(thread)

  14. Book code Two structures containing data to be sent to each thread: DataStruct data[2]; data[0].a = a; data[0].b = b; … data[1].a = a + N/2; data[1].b = b+N/2; … … CUTThread thread = start_thread(routine, &(data[0]); routine (&(data[1])); … end_thread(thread); Elements of structure previously defined Two threads, each executing routine()

  15. Pinned memory on multiple GPUs Pinned memory only pinned by thread allocating the pinned memory Other threads see it as pageable and access slower. These threads cannot use cudaMemcpyAsync, which requires pinned memory “Portable” pinned memory Memory allowed to move between host threads and any thread to see it as pinned memory Use cudaHostAlloc and include cudaAllocPortable flag

  16. Questions

More Related