<commit_before>---
layout: post
title:  "Cuda 3 - Optimize GPU Programs"
date:   2016-4-5 17:00:00
tags: [cuda, Optimize]
categories: Programming
---

> UDACITY教程 [Intro to Parallel Programming][link] 

> * Basics on GPU, CUDA, Memory Model
 * Parallel Algorithms(Reduce, Scan, Histogram, Sort)
 * *Optimize Parallel GPU Programs*
 * Others(Library, OpenACC, Dynamic parallelism)

[link]: https://www.udacity.com/wiki/cs344

### 1. **APOD**
![apod](http://7xno5y.com1.z0.glb.clouddn.com/apod.png)

### 2. **HotsPots**
Amdahl's Law: 

**max speedup = 1 / (1-p)** (p is portion can be parallelized)

### 3. **Occupancy**
![gpu-occupancy](http://7xno5y.com1.z0.glb.clouddn.com/gpu-occupancy.png)

[Optimize Example: Transpose](https://github.com/wykvictor/cs344-udacity/blob/master/Lesson%20Code%20Snippets/Lesson%205%20Code%20Snippets/transpose.cu)

![gpu-tile](http://7xno5y.com1.z0.glb.clouddn.com/gpu-tiling.png)

### 4. **WARP-Avoid Thread Divergence**
Set of threads that execute the same instructions at a time

Nvidia Card is **32**, so:

{% highlight C++ %}
switch(threadIdx.x / 32) {case 0...31}  // no slowdown
switch(threadIdx.x % 32) {case 0...31}  // 32x slowdown
{% endhighlight %}

### 5. **Streams** 
Sequence of operations execute in order(memory transfers, kernels)
![gpu-stream](http://7xno5y.com1.z0.glb.clouddn.com/gpu-stream.png)

{% highlight C++ %}
cudaStream_t s1, s2;
cudaStreamCreateWithFlags(&s1, cudaStreamNonBlocking);
checkCudaErrors(cudaMemcpyAsync(d_fimTheta, out_t, numbytes, cudaMemcpyHostToDevice, s1));
cudaStreamCreateWithFlags(&s2, cudaStreamNonBlocking);
checkCudaErrors(cudaMemcpyAsync(d_fimMag, out_m, numbytes, cudaMemcpyHostToDevice, s2));

A<<<1, 1024, s1>>>(d_fimTheta);
B<<<1, 1024, s2>>>(d_fimMag);
cudaDeviceSynchronize();
cudaStreamDestroy(s1); cudaStreamDestroy(s2);
{% endhighlight %}

### 6. **Summary**
![gpu-improve-1](http://7xno5y.com1.z0.glb.clouddn.com/gpu-improve-1.png)

![gpu-improve-2](http://7xno5y.com1.z0.glb.clouddn.com/gpu-improve-2.png)

![gpu-improve-3](http://7xno5y.com1.z0.glb.clouddn.com/gpu-improve-3.png)

### 7. Problem 5
Fast histogram:
[Fast histogram](https://github.com/wykvictor/cs344-udacity/blob/master/Problem%20Sets/Problem%20Set%205/student.cu)<commit_msg>something about cudaStreamCreateWithFlags<commit_after>---
layout: post
title:  "Cuda 3 - Optimize GPU Programs"
date:   2016-4-5 17:00:00
tags: [cuda, Optimize]
categories: Programming
---

> UDACITY教程 [Intro to Parallel Programming][link] 

> * Basics on GPU, CUDA, Memory Model
 * Parallel Algorithms(Reduce, Scan, Histogram, Sort)
 * *Optimize Parallel GPU Programs*
 * Others(Library, OpenACC, Dynamic parallelism)

[link]: https://www.udacity.com/wiki/cs344

### 1. **APOD**
![apod](http://7xno5y.com1.z0.glb.clouddn.com/apod.png)

### 2. **HotsPots**
Amdahl's Law: 

**max speedup = 1 / (1-p)** (p is portion can be parallelized)

### 3. **Occupancy**
![gpu-occupancy](http://7xno5y.com1.z0.glb.clouddn.com/gpu-occupancy.png)

[Optimize Example: Transpose](https://github.com/wykvictor/cs344-udacity/blob/master/Lesson%20Code%20Snippets/Lesson%205%20Code%20Snippets/transpose.cu)

![gpu-tile](http://7xno5y.com1.z0.glb.clouddn.com/gpu-tiling.png)

### 4. **WARP-Avoid Thread Divergence**
Set of threads that execute the same instructions at a time

Nvidia Card is **32**, so:

{% highlight C++ %}
switch(threadIdx.x / 32) {case 0...31}  // no slowdown
switch(threadIdx.x % 32) {case 0...31}  // 32x slowdown
{% endhighlight %}

### 5. **Streams** 
Sequence of operations execute in order(memory transfers, kernels)
![gpu-stream](http://7xno5y.com1.z0.glb.clouddn.com/gpu-stream.png)

{% highlight C++ %}
cudaStream_t s1, s2;
// behaviour is full concurrency: default stream also maps to a single Stream
// https://devblogs.nvidia.com/parallelforall/gpu-pro-tip-cuda-7-streams-simplify-concurrency/
// cudaStreamCreateWithFlags(&s1, cudaStreamNonBlocking);
cudaStreamCreate(&s1);
checkCudaErrors(cudaMemcpyAsync(d_fimTheta, out_t, numbytes, cudaMemcpyHostToDevice, s1));
cudaStreamCreate(&s2);
checkCudaErrors(cudaMemcpyAsync(d_fimMag, out_m, numbytes, cudaMemcpyHostToDevice, s2));

A<<<1, 1024, s1>>>(d_fimTheta);
B<<<1, 1024, s2>>>(d_fimMag);
cudaDeviceSynchronize();
cudaStreamDestroy(s1); cudaStreamDestroy(s2);
{% endhighlight %}

### 6. **Summary**
![gpu-improve-1](http://7xno5y.com1.z0.glb.clouddn.com/gpu-improve-1.png)

![gpu-improve-2](http://7xno5y.com1.z0.glb.clouddn.com/gpu-improve-2.png)

![gpu-improve-3](http://7xno5y.com1.z0.glb.clouddn.com/gpu-improve-3.png)

### 7. Problem 5
Fast histogram:
[Fast histogram](https://github.com/wykvictor/cs344-udacity/blob/master/Problem%20Sets/Problem%20Set%205/student.cu)
{'fraction_non_alphanumeric': 0.17725037417147743, 'fraction_numerical': 0.05088732093222151, 'mean_word_length': 5.2456608811749, 'pattern_counts': {'":': 0, '<': 15, '<?xml version=': 0, '>': 19, 'https://': 7, 'lorem ipsum': 0, 'www.': 2, 'xml': 0}, 'pii_count': 0, 'substrings_counts': 0, 'word_list_counts': {'cursed_substrings.json': 0, 'profanity_word_list.json': 0, 'sexual_word_list.json': 0, 'zh_pornsignals.json': 0}}
{'dir': 'git-commits-cleaned', 'id': '5158979', 'n_tokens_mistral': 2067, 'n_tokens_neox': 1864, 'n_words': 322}