Cuda 平行運算機制 報告者：林威辰.

Presentation on theme: "Cuda 平行運算機制 報告者：林威辰."— Presentation transcript:

1 Cuda 平行運算機制 報告者：林威辰

2 Slides MPI基本定理 各種平行運算的簡介 CUDA簡介 使用VS2005.net 建置CUDA CUDA基本知識 CUDA硬體架構 CUDA缺點 My Research

3 Slides http://courses.ece.uiuc.edu/ece498/al/Syllabus.html
National Center for High-Performance Computing NCHC教育訓練網

4 MPI基本定理

5 Parallel Computing Introduction
Flynn’s Taxonomy Amdahl’s Law Moore’s Law Finding Concurrency

6 Flynn’s Taxonomy

7 Amdahl’s Law Expected speedup from partial improvement
P：proportion of program that is parallel S：speedup of parallel portion

8 Moore’s Law The number of transistors on ICs doubles every 18 months
In the past, CPUs have been taking advantage of Moore,s Law to： Increase clock frequency Increase exploitation of ILP ILP：Instruction-Level Parallelism Result：CPUs gets faster Increasingly difficult to get faster

9 各種平行運算的簡介

10 平行運算示意圖

11 一般計算 (Serial Computing)

12 平行計算：二處理器

13 平行計算：四處理器

14 N Processors

15 Performance Development

16 Parallel Processing Technology
Shared Memory Multiprocessor Distributed Memory Multiprocessor System Clustering System

17 Shared Memory Multiprocessor
Shared Memory Multiprocessor別稱為Symmetric Multiprocessors，簡稱為SMP。 此架構採用System Bus的方式，將系統的CPU、Local Memory以及I/O裝置相連接，透過相同的作業系統，將不同執行序的工作分發給比較輕鬆的CPU，以達到分工的作用。

18 Shared Memory Multiprocessor(續)

19 Distributed Memory Multiprocessor System
Distributed Memory Multiprocessor System別稱Massive Parallel Processor，簡稱MPP。 這種架構是在同一部計算機中有許多CPU，並且這些CPU擁有屬於自己獨立的Local Memory，而CPU各自之間只能靠Message Passing 作為溝通橋樑 。

20 Distributed Memory Multiprocessor System(續)

21 Clustering System Clustering System架構是數台獨立的計算機，經由高速網路連結在一起，形成一個巨大的系統，而每台獨立的計算機都擁有各自的CPU、Local Memory和作業系統。Clustering System架構因為每部計算機都是獨立的，不需要大量資訊交換，只有在必要時經由高速網路交換資訊 。 目前Clustering System上的平行計算是使用Message Passing的概念，使用TCP/IP的通訊協定作為溝通的橋樑，常見的有兩種，分別如下： Parallel Virtual Machine Message Passing Interface

22 Clustering System(續)

23 Parallel Virtual Machine
Parallel Virtual Machine提供一組Application Program Interface，簡稱API，讓使用者可以直覺並且有效率的開發平行處理程式在現有的硬體上，它將在此系統上的異質作業系統當作單一的平行處理計算機，透明化的處理所有訊息的傳遞、資料的轉換和網路工作的排程。

24 Message Passing Interface
Message Passing Interface定義在一套標準的訊息傳遞介面，而跟PVM不同的是，不是讓不同作業系統在同一虛擬平台上執行，它不包含系統的Task，也不會直接控制I/O裝置的支援，它只是扮演一個溝通介面層的角色。這些特色使得很多PVM平台使用者轉向使用MPI平台。

25 CUDA簡介

26 GPGPU 將GPU用在非傳統的3D圖形顯示卡方面的應用，一般來說，會把這樣的應用叫作GPGPU ( General-pupose computing on graphics processing units ) 。 適用問題：大多是把一個可以用來大量拆解成多個相同、但彼此並不相關的小問題的情況；在這種情況下，用GPGPU的方法，就可以把這些一樣的小問題，給顯示卡的GPU來大量平行化的處理。 缺點：傳統的GPGPU的開發方法，都是透過OpenGL 或 Direct3D這一類現有的圖形函式庫，以編寫shading language 的方法，控制 shader 來想辦法做到自己想要的計算

27 CUDA “Compute Unified Device Architecture”
網址：

28 選擇OS

29 下載套件

30 使用VS2005.net 建置CUDA

31 使用VS2005.NET 加入規則

32 使用VS2005.NET(續) 加入規則(續)

33 使用VS2005.NET(續) 規則檔的選取有兩種方式
在CUDA SDK目錄之中，有提供官方的build rule，位於「C://Program Files\NVIDIA Corporation\NVIDIA CUDA SDK\common\scripts」底下，檔名是「nvcc.rules」

34 使用VS2005.NET(續) 自定建置規則

35 使用VS2005.NET(續) CUDA： CudaCompile： nvcc.rules

36 使用VS2005.NET(續)

37 使用VS2005.NET(續) 有執行GPU程式的副檔名：.cu

38 使用VS2005.NET(續)

39 使用VS2005.NET(續)

40 使用VS2005.NET(續)

41 使用VS2005.NET(續)

42 使用VS2005.NET(續)

43 CUDA基本知識

44 Design Philosophy GPUs use the transistor budget to get wider
Good for data-parallel computations

45 CPU & GPU 比較

46 CUDA 效能測試 CPU GPU GPU Texture GPU without transfer
使用 CPU 來做計算，但是沒有使用 OpenMP 之類的平行化計算，所以應該只有用到一顆 CPU 核心。 GPU 簡單的透過 CUDA global memory 來做，沒有特別最佳化。 GPU Texture 使用 CUDA Linear memory 的 texture 來做存取。 GPU without transfer 簡單的透過 CUDA global memory 來做，沒有特別最佳化；不過在計算時間時，不考慮將記憶體由 host 複製到 device 的時間。 GPU Texture without transfer 使用 CUDA Linear memory 的 texture 來做存取；不過在計算時間時，不考慮將記憶體由 host 複製到 device 的時間。

47 CUDA 效能測試(續)

48 Finding Concurrency At high level, algorithms can be decomposed by tasks and data Task：Groups od instructions that can execute in parallel Data：Partitions in the data that can be used independently Inside tasks, there is parallelism among the instructions Level Tree

49 Example int a[10] = {1,2,3,4,5,6,7,8}; int sum = 0; for (i = 1 ; i <=8 ; i++) { sum += a[i]; } printf(“%d\n”,sum); sum = 36

50 Example(續) main(){ int compute = 8, sum = 0; int a[8] = {1,2,3,4,5,6,7,8}; sum = compute_sum(1,a); printf(“%d\n”,sum) } compute_sum (tid, a){ if (tid >= 8) return a[0]; for(i = 0 ; i < 8 ; i = i+tid*2) a[i] = a[i] + a[i+tid]; return compute_sum (start, tid*2, a); } sum = 36

51 Geforce 8800

52 CUDA的架構 CUDA的程式架構 Host (CPU) Device (GPU)

53 CUDA – C with no shader limitations
Integrated host + device app C program Serial or modestly parallel parts in host C code Highly parallel parts in device SPMD kernel C code

54 CUDA Devices and Threads
A compute device Is a coprocessor to the CPU or Host Has its own DREM Runs many threads in parallel Is typically a GPU but can also be another type of parallel processing Differences Between GPU and CPU threads CPU：software thread GPU：hardware thread (transfer more fast)

55 Arrays of Parallel Threads
A CUDA kernel is executed by an array of threads All threads run the same code (SPMD) Each thread has an ID that it uses to compute memory addresses and make control decisions

56 Threads Blocks：Scalable Cooperation
Divide monolithic thread array into multiple blocks Threads within a block cooperate via shared memory, atomic operations and barrier synchronization. Threads in different blocks cannot cooperate

57 Block IDs and Thread IDs
Each thread uses IDs th decide what data to work on Block ID：1D or 2D Thread ID：1D , 2D , or 3D

58 CUDA Device Memory Space
Each thread can ： R/W per-thread registers R/W per-thread local memory R/W per-thread shared memory R/W per-grid global memory R/W per-grid constant memory R/W per-block texture memory The host can R/W global, constant, and texture memories

59 Parallel Memory Sharing
Local Memory： (per-thread) Private per thread Auto variable, register spill Speed slow Shared Memory： (per-Block) Shared by threads of the same block Inter-thread communication Global Memory: (per-application) Shared by all threads Inter-Grid communication

60 CUDA Device Memory Allocation
cudaMalloc() Allocates object in the device Global Memory Require two parameters Address of a pointer to the allocated object Size of allocated object cudaFree() Frees object from device Global Memory Pointer to freed object

61 CUDA Device Memory Allocation (續)
Example： int width = 32; float* Array; int size = width * width * sizeof(float); cudaMalloc((void**) &Array, size); cudaFree(Array);

62 CUDA Host-Device Data Transfer
cudaMemcpy() Memory data transfer Requires four parameters Pointer to destination Pointer to source Number of bytes copied Type of transfer Host to Host Host to Device Device to Host Device to Device Asynchronous transfer

63 CUDA Host-Device Data Transfer(續)
Example： int width = 32; float* Array; float HostArray[width * width] int size = width * width * sizeof(float); cudaMalloc((void**) &Array, size); cudaMemcpy(&Array, HostArray, size, cudaMemcpyHostToDevice); cudaMemcpy(HostArray, Array, size, cudaMemcpyDeviceToDevice); cudaFree(Array);

64 CUDA Function Declarations
Executed on the： Only callable from the： __device__ float DeviceFunc() device __global__ void KernelFunc() host __host__ float HostFunc() __global__ defines a kernel function Must return void __device__ and __host__ can be used together

65 Language Extension： Built-in Variables
dim3 gridDim; Dimensions of the grid in blocks dim3 blockDim; dim3 blockIdx; Block index within the grid dim3 threadIdx; Thread index within the block

66 Device Runtime Component： Mathematical Functions
Some mathematical functions( e.g. sin(x) ) have a less accurate, but faster device-only version ( e.g. __sin(x) ) __pow __log, __log2, __log10 __exp __sin, __cos, __tan

67 Device Runtime Component： Synchronization Funtion
void __syncthreads(); Synchronizes all threads in a block Once all threads have reached this point, execution resumes normally Used to avoid RAW/WAR/WAW hazards when accessing shared or global memory

68 CUDA硬體架構

69 處理單元 nVidia的GPU裡處理單元 SP ( Streaming Processor ) SM ( Streaming Multiprocessor ) TPC ( Texture Processing Clusters ) 以G80/G92 的架構之下，總共有128個SP，以8個SP為一組，組成16個SM，再以2個SM為一個TPC，共分成8個TPC來運作。

70 處理單元 (續)

71 處理單元 (續)

72 SM 中的 Warp 和 Block device 實際在執行時，會以block為執行單位，把Block分配給SM作計算。
block中的thread，是以「 warp 」為單位，32個thread會組成一個warp來執行。 warp分組的動作是以SM自動進行，會以連續的方式來作分組。 一個SM一次只會執行一個block裡的一個warp。

73 Warp 排程 以下就是一個Warp排程的例子

74 Transparent Scalability
Hardware is free to assign blocks to any processor at any time A kernel scales across any number of parallel processors

75 CUDA缺點

76 CUDA缺點 1. 太新 2. 綁顯示卡 3. 不支援 Double 4. debug麻煩 5. 記憶體配置常常會抓到已使用記憶體
不能使用遞迴

77 My Research

78 Master/Slave架構 主要是從兩個組件結合而成，運算的過程中，由一個Server的架構對應於多組Client，並且從Server動態的分割出多組運算區段，使用TCP/IP通訊協定，傳輸工作分散至各Client，使工作分配類似星狀架構。

79 Model

80 Model(續)

81 整體架構

82 分工 CPU： Message Passing Search GPU Compute arbitrage

83 Example Theorem 8 If C and P is a rationally determined American cll and put price, then C and P is convex function of its exercise price (X) three otherwise identical calls with strike prices Where Remark：The above arguments can also be applied to European options Robert C Merton (1973)

84 Example(續)

85 Thank you