University:
Boston University
Course:
GRS CS 910 | Directed Study: Systems
Academic year:

2022
Views:

104

Pages:

17
Author:

maselarmwpa

CUDA - Vector add example __global__ void vec_add (float *A, float *B, float *C, int N) { // Using 1D threadID and block ID Thread int i = threadIdx.x + blockDim.x*blockIdx.x; if (i>=N) {return;} C[i] = A[i] + B[i]; } int main() { .... // Launch kernel with N/256 blocks of 256 threads int blocks = int(N‐0.5)/256 + 1; vec_add<<>> (A_d, B_d, C_d, N); } indexing CUDA - Vector add example __global__ void vec_add (float *A, float *B, float *C, int N) { // Using 1D threadID and block ID int i = threadIdx.x + blockDim.x*blockIdx.x; if (i>=N) {return;} C[i] = A[i] + B[i]; } Perform operation int main() { .... // Launch kernel with N/256 blocks of 256 threads int blocks = int(N‐0.5)/256 + 1; vec_add<<>> (A_d, B_d, C_d, N); } CUDA - Vector add example int N = 200; // Allocate and initialize host CPU float *A_h = new float[N]; float *B_h = new float[N]; float *C_h = new float[N]; for(int i=0; i>> (A_d, B_d, C_d, N); // Copy back from device to host cudaMemcpy(C_h, C_d, N*sizeof(float), cudaMemcpyDeviceToHost); // Free device cudaFree(A_d); cudaFree(B_d); cudaFree(C_d); CUDA - Vector add example Copy to device // Copy host memory to device cudaMemcpy(A_d, A_h, N*sizeof(float), cudaMemcpyHostToDevice); cudaMemcpy(B_d, B_h, N*sizeof(float), cudaMemcpyHostToDevice); // Run grid of N/256 blocks of 256 threads each int blocks = int(N‐0.5)/256 + 1; vec_add<<>> (A_d, B_d, C_d, N); // Copy back from device to host cudaMemcpy(C_h, C_d, N*sizeof(float), cudaMemcpyDeviceToHost); // Free device cudaFree(A_d); cudaFree(B_d); cudaFree(C_d); CUDA - Vector add example // Copy host memory to device cudaMemcpy(A_d, A_h, N*sizeof(float), cudaMemcpyHostToDevice); cudaMemcpy(B_d, B_h, N*sizeof(float), cudaMemcpyHostToDevice); // Run grid of N/256 blocks of 256 threads each int blocks = int(N‐0.5)/256 + 1; vec_add<<>> (A_d, B_d, C_d, N); Launch kernel // Copy back from device to host cudaMemcpy(C_h, C_d, N*sizeof(float), cudaMemcpyDeviceToHost); // Free device cudaFree(A_d); cudaFree(B_d); cudaFree(C_d); CUDA - Vector add example // Copy host memory to device cudaMemcpy(A_d, A_h, N*sizeof(float), cudaMemcpyHostToDevice); cudaMemcpy(B_d, B_h, N*sizeof(float), cudaMemcpyHostToDevice); // Run grid of N/256 blocks of 256 threads each int blocks = int(N‐0.5)/256 + 1; vec_add<<>> (A_d, B_d, C_d, N); // Copy back from device to host cudaMemcpy(C_h, C_d, N*sizeof(float), cudaMemcpyDeviceToHost); // Free device cudaFree(A_d); cudaFree(B_d); cudaFree(C_d); Copy back to host CUDA - Vector add example // Copy host memory to device cudaMemcpy(A_d, A_h, N*sizeof(float), cudaMemcpyHostToDevice); cudaMemcpy(B_d, B_h, N*sizeof(float), cudaMemcpyHostToDevice); // Run grid of N/256 blocks of 256 threads each int blocks = int(N‐0.5)/256 + 1; vec_add<<>> (A_d, B_d, C_d, N); // Copy back from device to host cudaMemcpy(C_h, C_d, N*sizeof(float), cudaMemcpyDeviceToHost); // Free device cudaFree(A_d); cudaFree(B_d); cudaFree(C_d); Free device CUDA - Vector add example __global__ void vec_add (float *A, float *B, float *C, int N) { // Using 1D threadID and block ID int i = threadIdx.x + blockDim.x*blockIdx.x; if (i>=N) {return;} C[i] = A[i] + B[i]; } Thread 0 Thread i Thread i+1 Thread N-1 Fetch A[0], B[0] Fetch A[i], B[i] Fetch A[i+1], B[i+1] Fetch A[N-1], B[N-1] Operate Operate Operate Operate Write C[0] Write C[i] Write C[i+1] Write C[N-1] ... ... CUDA - To keep in mind ‣ Keep executions simple - Leave tough things to the CPU ‣ Performance: - Load balance - Data parallelism - Avoid conflicts - Keep the GPU busy! ‣ How fast can we go? CUDA - To keep in mind ‣ Load balance - All threads should do the same amount of work ‣ Data Parallelism (Massive) - Arrange your algorithm so is data parallel friendly - Look for regularity ‣ Avoid conflicts - Data access conflicts will serialize your applications or give you wrong answers! CUDA - To keep in mind ‣ Keep the GPU busy! - High peak compute throughput compared to bandwidth - Fermi: ‣ 1TFLOP peak throughput, 144 GB/s peak off chip memory access (36 Gfloats per second) ‣ 4*1TFLOP = 4000GB/s for peak throughput! ‣ 1000/36 ≈ 28 operations per fetched value - Need to hide latency! CUDA - To keep in mind ‣ How fast can we go? - Keep Amdahl’s law in mind - Know how much parallelization can be done in your application - Measure independent parts of your algorithm before going to the GPU 1 speedup = (1 − P ) + P S P: parallel portion S: sequential portion CUDA - Applications ‣ Ultrasound imaging ‣ Molecular dynamics ‣ Seismic migration ‣ Astrophysics simulations ‣ Graphics ‣ ....