1 / 12

Enhanced CUDA Techniques for Fast Histogram Computation with Parallel Programming

This document discusses advanced CUDA techniques for efficiently computing histograms using parallel processing. It highlights the implementation of a computationally intense evaluation function and improved crossover and mutation operators. The CUDA version employs atomic operations for thread safety during histogram updates and utilizes CUDA events for precise timing measurements. We outline the use of shared memory and texture memory for enhancing performance and illustrate kernel implementations designed for large-scale data processing on the GPU. Sample code is provided for practical understanding.

odessa-spears
Télécharger la présentation

Enhanced CUDA Techniques for Fast Histogram Computation with Parallel Programming

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. CUDA More GA, Events, Atomics

  2. GA Revisited • Speedup with a more computationally intense evaluation function • Parallel version of the crossover and mutation operators • In this version of CUDA, no random in GPU • Used rand.cu, linear congruent; fairly poor random number generator • New kernel invocation based on POP_SIZE • <<<128,32>>> • See file online for code

  3. Events • Can use CUDA events to time how long something takes cudaEvent_t start, stop; cudaEventCreate(&start); cudaEventCreate(&stop); … cudaEventRecord(start, 0); // Do some work on the GPU, stream 0 cudaEventRecord(stop, 0); cudaEventSynchronize(stop); float elapsedTime; cudaEventElapsedTime(&elapsedTime, start, stop)); printf(“Time: %3.2f milliseconds\n”,elapsedTime);

  4. Texture Memory • Just Mentioning – requires use of some graphics-like functions for non-graphics use • E.g. tex1Dfetch(textureIn, offset); • Designed for OpenGL rendering pipeline • Read-Only memory but cached on-chip • Good for memory access patterns using spatial locality (e.g. access other items near a 2D coordinate)

  5. Atomics • Requires Compute Capability 1.2 or higher • Our Tesla supports 1.3 • GeForce GTX TX 580 and Tesla S2070 support 2.0 • Atomics are used to manage contention from multiple processors or threads to shared resources • Primitives discussed earlier; e.g. Exchange instruction to implement barrier synchronization, keep memory consistent

  6. Example: Histogram (CPU version 1) #include <stdio.h> #include <string.h> #define SIZE 26 int main() { char *buffer = "THE QUICK BROWN FOX JUMPED OVER THE LAZY DOGS"; inthisto[SIZE]; // Count of letters A-Z in buffer, where [0]=A's, // [1]=B's, etc. for (int i = 0; i < SIZE; i++) histo[i] = 0; for (int i = 0; i < strlen(buffer); i++) histo[buffer[i]-'A']++; for (int i = 0; i < SIZE; i++) { printf("%c appears %d times.\n", ('A'+i), histo[i]); } return 0; }

  7. Longer Histogram #include <stdio.h> #include <string.h> #define STRSIZE 26*20000 // 520000 chars #define SIZE 26 int main() { char *buffer; inthisto[SIZE]; printf("Creating data buffer...\n"); // Allocate buffer buffer = (char *)malloc(STRSIZE); for (int i = 0; i < 26*20000; i++) buffer[i]= 'A' + (i % 26); for (int i = 0; i < SIZE; i++) histo[i] = 0; printf("Counting...\n"); for (int i = 0; i < strlen(buffer); i++) printf("Counting...\n"); for (int i = 0; i < strlen(buffer); i++) histo[buffer[i]-'A']++; for (int i = 0; i < SIZE; i++) { printf("%c appears %d times.\n", ('A'+i), histo[i]); } free (buffer); return 0; }

  8. First Attempt – GPU Version • Have a thread operate on each character • 520,000 chars • Let’s use 16250 blocks and 32 threads/block • Index into the string is the usual • blockIdx.x * blockDim.x + threadIdx.x __global__ void histo_kernel(char *dev_buffer, int *dev_histo) { inttid = blockIdx.x * blockDim.x + threadIdx.x; char c = dev_buffer[tid]; dev_histo[c-'A']++; } histo_kernel<<<16250,32>>>(dev_buffer, dev_histo);

  9. FAIL – Why?

  10. Atomics to the Rescue • Guarantee only one thread can perform the operation at a time • Must compile with nvcc –arch=sm_11 or nvcc –arch=sm_12 atomicAdd, atomicSub, atomicMin, atomicMax, atomicExch, atomicInc See http://developer.download.nvidia.com/compute/cuda/2_0/docs/NVIDIA_CUDA_Programming_Guide_2.0.pdf

  11. Modified Kernel __global__ void histo_kernel(char *dev_buffer, int *dev_histo) { inttid = blockIdx.x * blockDim.x + threadIdx.x; char c = dev_buffer[tid]; atomicAdd(&dev_histo[c-'A'],1); }

  12. Shared Memory Atomics • Compute Capability 1.2+ allows atomics in shared memory; using 16250 blocks, 32 threads __global__ void histo_kernel(char *dev_buffer, int *dev_histo) { __shared__ intblock_histo[SIZE]; if (threadIdx.x < 26) block_histo[threadIdx.x] = 0; __syncthreads(); inttid = blockIdx.x * blockDim.x + threadIdx.x; char c = dev_buffer[tid]; atomicAdd(&block_histo[c-'A'],1); __syncthreads(); if (threadIdx.x < 26) atomicAdd(&dev_histo[threadIdx.x],block_histo[threadIdx.x]); } Zero histogram per block Add block histogram Add to global histogram

More Related