垃圾CPU，耗我時光——Jetson Nano 初體驗2

CPU與GPU性能測試

1. CPU性能測試：計算圓周率

bc 命令是任意精度計算器語言，通常在 linux 下當計算器用。它類似基本的計算器, 使用這個計算器可以做基本的數學運算
man 一下 bc 即可知道，a 是 bc 的一個內置函數，代表反正切 arctan ，由於 tan(pi/4) = 1 ，於是 4*arctan(1) = pi

計算圓周率的前一萬位（單線程）並與 Intel(R) Xeon(R) Platinum 8163 CPU 的CPU做對比

# jetson nano CPU 參數 lscpu Architecture: aarch64 Byte Order: Little Endian CPU(s): 4 On-line CPU(s) list: 0-3 Thread(s) per core: 1 Core(s) per socket: 4 Socket(s): 1 Vendor ID: ARM Model: 1 Model name: Cortex-A57 Stepping: r1p1 CPU max MHz: 1428.0000 CPU min MHz: 102.0000 BogoMIPS: 38.40 L1d cache: 32K L1i cache: 48K L2 cache: 2048K Flags: fp asimd evtstrm aes pmull sha1 sha2 crc32 # 計算圓周率的前一萬位（單線程） time echo "scale = 10000; 4*a(1)" | bc -l -q 3.1415926535897... real 5m22.161s user 5m21.496s sys 0m0.020s # Intel(R) Xeon(R) Platinum 8163 CPU 參數 lscpu Architecture: x86_64 CPU op-mode(s): 32-bit, 64-bit Byte Order: Little Endian CPU(s): 1 On-line CPU(s) list: 0 Thread(s) per core: 1 Core(s) per socket: 1 Socket(s): 1 NUMA node(s): 1 Vendor ID: GenuineIntel CPU family: 6 Model: 85 Model name: Intel(R) Xeon(R) Platinum 8163 CPU @ 2.50GHz Stepping: 4 CPU MHz: 2500.008 BogoMIPS: 5000.01 Hypervisor vendor: KVM Virtualization type: full L1d cache: 32K L1i cache: 32K L2 cache: 1024K L3 cache: 33792K NUMA node0 CPU(s): 0 Flags: fpu vme de pse tsc msr pae mce cx8 apic sep mtrr pge mca cmov pat pse36 clflush mmx fxsr sse sse2 ss ht syscall nx pdpe1gb rdtscp lm constant_tsc rep_good nopl cpuid pni pclmulqdq ssse3 fma cx16 pcid sse4_1 sse4_2 x2apic movbe popcnt tsc_deadline_timer aes xsave avx f16c rdrand hypervisor lahf_lm abm 3dnowprefetch invpcid_single pti ibrs ibpb stibp fsgsbase tsc_adjust bmi1 hle avx2 smep bmi2 erms invpcid rtm mpx avx512f avx512dq rdseed adx smap avx512cd avx512bw avx512vl xsaveopt xsavec xgetbv1 # 計算圓周率的前一萬位（單線程） time echo "scale = 10000; 4*a(1)" | bc -l -q 3.1415926535897... real 2m20.695s user 2m19.211s sys 0m0.047s 

單核 CPU 性能大概是 Intel(R) Xeon(R) Platinum 8163 的一半

2. CPU與GPU對比測試

• CUDA在線學習

2.1 四種計算機模型

GPU設計的初衷就是為了減輕CPU計算的負載，將一部分圖形計算的功能設計到一塊獨立的處理器中，將矩陣變換、頂點計算和光照計算等操作從 CPU 中轉移到 GPU中，從而一方面加速圖形處理，另一方面減小了 CPU 的工作負載，讓 CPU 有時間去處理其它的事情。
在GPU上的各個處理器采取異步並行的方式對數據流進行處理，根據費林分類法（Flynn's Taxonomy），可以將信息流（information stream）分成指令（Instruction）和數據（Data）兩種，據此又可分成四種計算機類型：

• 單一指令流單一數據流計算機（SISD）：單核CPU
• 單一指令流多數據流計算機（SIMD）：GPU的計算模型
• 多指令流單一數據流計算機（MISD）：流水線模型
• 多指令流多數據流計算機（MIMD）：多核CPU

2.2 CPU 與 GPU 結構差異

 

CPU 與 GPU 結構差異

(1)CPU設計理念：低延時

 

CPU設計理念：低延時

• ALU：CPU有強大的ALU（算術運算單元）,它可以在很少的時鍾周期內完成算術計算。
  • 當今的CPU可以達到64bit 雙精度。執行雙精度浮點源算的加法和乘法只需要1～3個時鍾周期。
  • CPU的時鍾周期的頻率是非常高的，達到1.532～4gigahertz(千兆HZ, 10的9次方).
• Cache：大的緩存也可以降低延時。保存很多的數據放在緩存里面，當需要訪問的這些數據，只要在之前訪問過的，如今直接在緩存里面取即可。
• Control：復雜的邏輯控制單元。
  • 當程序含有多個分支的時候，它通過提供分支預測的能力來降低延時。
  • 數據轉發。 當一些指令依賴前面的指令結果時，數據轉發的邏輯控制單元決定這些指令在pipeline中的位置並且盡可能快的轉發一個指令的結果給后續的指令。這些動作需要很多的對比電路單元和轉發電路單元。

(2)GPU設計理念：大吞吐量

 

GPU設計理念：大吞吐量

• ALU，Cache：GPU的特點是有很多的ALU和很少的cache. 緩存的目的不是保存后面需要訪問的數據的，這點和CPU不同，而是為thread提高服務的。如果有很多線程需要訪問同一個相同的數據，緩存會合並這些訪問，然后再去訪問dram（因為需要訪問的數據保存在dram中而不是cache里面），獲取數據后cache會轉發這個數據給對應的線程，這個時候是數據轉發的角色。但是由於需要訪問dram，自然會帶來延時的問題。
• Control：控制單元（左邊黃色區域塊）可以把多個的訪問合並成少的訪問。

GPU的雖然有dram延時，卻有非常多的ALU和非常多的thread. 為了平衡內存延時的問題，我們可以中充分利用多的ALU的特性達到一個非常大的吞吐量的效果。盡可能多的分配多的Threads.通常來看GPU ALU會有非常重的pipeline就是因為這樣。

2.3 Nvidia GPU架構

(1)硬件架構

• SP：最基本的處理單元，streaming processor，也稱為CUDA core。最后具體的指令和任務都是在SP上處理的。GPU進行並行計算，也就是很多個SP同時做處理。
• SM：多個SP加上其他的一些資源組成一個streaming multiprocessor。也叫GPU大核，其他資源如：warp scheduler，register，shared memory等。SM可以看做GPU的心臟（對比CPU核心），register和shared memory是SM的稀缺資源。CUDA將這些資源分配給所有駐留在SM中的threads。因此，這些有限的資源就使每個SM中active warps有非常嚴格的限制，也就限制了並行能力。

 

Nvidia GPU硬件架構

(2)軟件架構

CUDA在軟件方面組成有：一個CUDA庫、一個應用程序編程接口（API）及其運行庫(Runtime)、兩個較高級別的通用數學庫，即CUFFT和CUBLAS。CUDA改進了DRAM的讀寫靈活性，使得GPU與CPU的機制相吻合。另一方面，CUDA 提供了片上（on-chip）共享內存，使得線程之間可以共享數據。應用程序可以利用共享內存來減少DRAM的數據傳送，更少的依賴DRAM的內存帶寬。

• thread：一個CUDA的並行程序會被以許多個threads來執行。
• block：數個threads會被群組成一個block，同一個block中的threads可以同步，也可以通過shared memory通信。
• grid：多個blocks則會再構成grid。
• warp：GPU執行程序時的調度單位，目前cuda的warp的大小為32，同在一個warp的線程，以不同數據資源執行相同的指令,這就是所謂 SIMT。

 

Nvidia GPU軟件架構

（3）軟硬件架構對應關系
從軟件上看，SM更像一個獨立的CPU core。SM（Streaming Multiprocessors）是GPU架構中非常重要的部分，GPU硬件的並行性就是由SM決定的。

當一個kernel啟動后，thread會被分配到這些SM中執行。大量的thread可能會被分配到不同的SM，同一個block中的threads必然在同一個SM中並行（SIMT）執行。每個thread擁有它自己的程序計數器和狀態寄存器，並且用該線程自己的數據執行指令，這就是所謂的Single Instruction Multiple Thread。

CUDA是一種典型的SIMT架構（單指令多線程架構），SIMT和SIMD（Single Instruction, Multiple Data）類似，SIMT應該算是SIMD的升級版，更靈活，但效率略低，SIMT是NVIDIA提出的GPU新概念。二者都通過將同樣的指令廣播給多個執行官單元來實現並行。一個主要的不同就是，SIMD要求所有的vector element在一個統一的同步組里同步的執行，而SIMT允許線程們在一個warp中獨立的執行。

2.4 CUDA C編程入門

（1）程序架構

CUDA程序構架分為兩部分：Host和Device。一般而言，Host指的是CPU，Device指的是GPU。在CUDA程序構架中，主程序還是由 CPU 來執行，而當遇到數據並行處理的部分，CUDA 就會將程序編譯成 GPU 能執行的程序，並傳送到GPU。而這個程序在CUDA里稱做核（kernel）。CUDA允許程序員定義稱為核的C語言函數，從而擴展了 C 語言，在調用此類函數時，它將由N個不同的CUDA線程並行執行N次，這與普通的C語言函數只執行一次的方式不同。執行核的每個線程都會被分配一個獨特的線程ID，可通過內置的threadIdx變量在內核中訪問此ID。
在 CUDA 程序中，主程序在調用任何 GPU 內核之前，必須對核進行執行配置，即確定線程塊數和每個線程塊中的線程數以及共享內存大小。

CUDA 設備擁有多個獨立的存儲空間，其中包括：全局存儲器、本地存儲器、共享存儲器、常量存儲器、紋理存儲器和寄存器

CUDA線程可在執行過程中訪問多個存儲器空間的數據，如下圖所示其中：

• 每個線程都有一個私有的本地存儲器。
• 每個線程塊都有一個共享存儲器，該存儲器對於塊內的所有線程都是可見的，並且與塊具有相同的生命周期。
• 所有線程都可訪問相同的全局存儲器。
• 此外還有兩個只讀的存儲器空間，可由所有線程訪問，這兩個空間是常量存儲器空間和紋理存儲器空間。全局、固定和紋理存儲器空間經過優化，適於不同的存儲器用途。紋理存儲器也為某些特殊的數據格式提供了不同的尋址模式以及數據過濾，方便 Host對流數據的快速存取。

CUDA 假設線程可在物理上獨立的設備上執行，此類設備作為運行C語言程序的主機的協處理器操作。內核在GPU上執行，而C語言程序的其他部分在CPU上執行（即串行代碼在主機上執行，而並行代碼在設備上執行）。此外，CUDA還假設主機和設備均維護自己的DRAM，分別稱為主機存儲器和設備存儲器。因而，一個程序通過調用CUDA運行庫來管理對內核可見的全局、固定和紋理存儲器空間。這種管理包括設備存儲器的分配和取消分配，還包括主機和設備存儲器之間的數據傳輸。

（2）CUDA C基礎

CUDA C是對C/C++語言進行拓展后形成的變種，兼容C/C++語法，文件類型為".cu"文件，編譯器為"nvcc"，相比傳統的C/C++，主要添加了以下幾個方面：

• 函數類型限定符：用來確定某個函數是在CPU還是GPU上運行，以及這個函數是從CPU調用還是從GPU調用
  • device表示從GPU調用，在GPU上執行
  • global表示從CPU調用，在GPU上執行，也稱之為kernel函數
  • host表示在CPU上調用，在CPU上執行
• 執行配置運算符：執行配置運算符<<<>>>，用來傳遞內核函數的執行參數。格式如下：
  kernel<<<gridDim, blockDim, memSize, stream>>>(para1, para2, ...);
  • gridDim表示網格的大小，可以是1,2,3維
  • blockDim表示塊的·大小，可以是1,2,3維
  • memSize表示動態分配的共享存儲器大小，默認為0
  • stream表示執行的流，默認為0
  • para1, para2等為核函數參數
• 五個內置變量：這些內置變量用來在運行時獲得Grid和Block的尺寸及線程索引等信息
  • gridDim: 包含三個元素x, y, z的結構體，表示Grid在三個方向上的尺寸，對應於執行配置中的第一個參數
  • blockDim: 包含上元素x, y, z的結構體，表示Block在三個方向上的尺寸，對應於執行配置中的第二個參數
  • blockIdx: 包含三個元素x, y, z的結構體，分別表示當前線程所在塊在網格中x, y, z方向上的索引
  • threadIdx: 包含三個元素x, y, z的結構體，分別表示當前線程在其所在塊中x, y, z方向上的索引
  • warpSize: 表明warp的尺寸
• 變量類型限定符：用來確定某個變量在設備上的內存位置
  • device表示位於全局內存空間，默認類型
  • share表示位於共享內存空間
  • constant表示位於常量內存空間
  • texture表示其綁定的變量可以被紋理緩存加速訪問
• 其他的還有數學函數、原子函數、紋理讀取、綁定函數等

2.5 CPU與GPU的矩陣乘法對比

（1）CPU單線程矩陣乘法

// CPU單線程矩陣乘法 #include <stdio.h> #include <stdlib.h> #include <sys/time.h> #include <unistd.h> #define w 2000 struct Matrix { int width; int height; float *elements; }; void matMul(float * M, float * N, float * P, int width){ for (int i = 0; i < width; i++){ for (int j = 0; j < width; j++){ float sum = 0; for (int k = 0; k < width; k++){ float a = M[i * width + k]; float b = N[k * width + j]; sum += a * b; } P[i * width + j] = sum; } } } int main(){ int width = w; int height = w; float * m = (float *)malloc (width * height * sizeof (float)); float * n = (float *)malloc (width * height * sizeof (float)); float * p = (float *)malloc (width * height * sizeof (float)); for (int i = 0; i < width * height; i++){ m[i] = 9.9; n[i] = 2.5; } struct timeval t1,t2; gettimeofday(&t1,NULL); double timeuse; matMul(m, n, p, w); gettimeofday(&t2,NULL); timeuse = t2.tv_sec - t1.tv_sec + (t2.tv_usec - t1.tv_usec)/1000000.0; printf("Use Time:%f\n",timeuse); return 0; } 

然后編譯運行

gcc  cpu_sigle.c -O3 -o cpu_sigle

./cpu_sigle

Use Time:52.641901 

（2）CPU多線程矩陣乘法

//CPU多線程矩陣乘法 #include <stdio.h> #include <stdlib.h> #include <pthread.h> #include <sys/time.h> #include <string.h> #include <unistd.h> #include <sys/types.h> #include <sys/stat.h> #include <fcntl.h> #define LOG_ #define SIZE 8000 int * A, * B; // 計算矩陣 int * result, * result2, * result3, * result4; // 結果矩陣 /* int A[SIZE][SIZE]; int B[SIZE][SIZE]; int result[SIZE][SIZE]; int result2[SIZE][SIZE]; int result3[SIZE][SIZE]; int result4[SIZE][SIZE]; */ int size; // 矩陣階數 pthread_t tid2[2]; // 雙線程id pthread_t tid3[3]; // 三線程id pthread_t tid4[4]; // 四線程id /* 雙線程函數 */ void twoThread1(){ int i, j, k; for (i = 0; i < size; i++) for (j = 0; j < size; j++) for (k = 0; k < size; k++){ if (i % 2 == 0) result2[i * size + j] += A[i * size + k] * B[k * size + j]; // result2[i][j] += A[i][k] * B[k][j]; } } void twoThread2(){ int i, j, k; for (i = 0; i < size; i++) for (j = 0; j < size; j++) for (k = 0; k < size; k++){ if (i % 2 != 0) result2[i * size + j] += A[i * size + k] * B[k * size + j]; // result2[i][j] += A[i][k] * B[k][j]; } } /* 雙線程函數 end */ /* 三線程函數 */ void threeThread1(){ int i, j, k; for (i = 0; i < size; i++) for (j = 0; j < size; j++) for (k = 0; k < size; k++){ if (i % 3 == 0) result3[i * size + j] += A[i * size + k] * B[k * size + j]; // result3[i][j] += A[i][k] * B[k][j]; } } void threeThread2(){ int i, j, k; for (i = 0; i < size; i++) for (j = 0; j < size; j++) for (k = 0; k < size; k++){ if (i % 3 != 0 && i % 2 != 0) result3[i * size + j] += A[i * size + k] * B[k * size + j]; // result3[i][j] += A[i][k] * B[k][j]; } } void threeThread3(){ int i, j, k; for (i = 0; i < size; i++) for (j = 0; j < size; j++) for (k = 0; k < size; k++){ if (i % 3 != 0 && i % 2 == 0) result3[i * size + j] += A[i * size + k] * B[k * size + j]; // result3[i][j] += A[i][k] * B[k][j]; } } /* 三線程函數 end */ /* 四線程函數 */ void fourThread1(){ int i, j, k; for (i = 0; i < size; i++) for (j = 0; j < size; j++) for (k = 0; k < size; k++){ if (i % 2 == 0 && i % 4 != 0) result4[i * size + j] += A[i * size + k] * B[k * size + j]; // result4[i][j] += A[i][k] * B[k][j]; } } void fourThread2(){ int i, j, k; for (i = 0; i < size; i++) for (j = 0; j < size; j++) for (k = 0; k < size; k++){ if (i % 4 == 0) result4[i * size + j] += A[i * size + k] * B[k * size + j]; // result4[i][j] += A[i][k] * B[k][j]; } } void fourThread3(){ int i, j, k; for (i = 0; i < size; i++) for (j = 0; j < size; j++) for (k = 0; k < size; k++){ if (i % 2 != 0 && i % 3 == 0) result4[i * size + j] += A[i * size + k] * B[k * size + j]; // result4[i][j] += A[i][k] * B[k][j]; } } void fourThread4(){ int i, j, k; for (i = 0; i < size; i++) for (j = 0; j < size; j++) for (k = 0; k < size; k++){ if (i % 2 != 0 && i % 3 != 0) result4[i * size + j] += A[i * size + k] * B[k * size + j]; // result4[i][j] += A[i][k] * B[k][j]; } } /* 四線程函數 end */ int main(){ int i, j, k, m, n; // 循環變量 struct timeval t1, t2; double timeuse; // 計時 char sizeChars[8]; // 階數寫入字符串 char timeChars[16]; // 耗時寫入字符串 // 申請空間, 計算矩陣和結果矩陣 A = (int *)malloc (sizeof (int) * SIZE * SIZE); B = (int *)malloc (sizeof (int) * SIZE * SIZE); result = (int *)malloc (sizeof (int) * SIZE * SIZE); result2 = (int *)malloc (sizeof (int) * SIZE * SIZE); result3 = (int *)malloc (sizeof (int) * SIZE * SIZE); result4 = (int *)malloc (sizeof (int) * SIZE * SIZE); for (i = 0; i < SIZE; i++) for (j = 0; j < SIZE; j++){ /* A[i][j] = 1; B[i][j] = 2; result[i][j] = 0; result2[i][j] = 0; result3[i][j] = 0; result4[i][j] = 0; */ A[i * SIZE + j] = 1; B[i * SIZE + j] = 2; result[i * SIZE + j] = 0; result2[i * SIZE + j] = 0; result3[i * SIZE + j] = 0; result4[i * SIZE + j] = 0; } int fd; fd = open ("./pthreadTime.txt", O_WRONLY | O_CREAT, 0777); lseek(fd, 0, SEEK_SET); for (size = 200; size <= SIZE; size += 200){ printf ("當前階數: %d\n", size); sprintf (sizeChars, "%d, ", size); write (fd, sizeChars, strlen (sizeChars)); #ifdef LOG printf ("A矩陣: \n"); for (i = 0; i < size; i++){ for (j = 0; j < size; j++){ printf ("%d ", A[i * size + j]); // printf ("%d ", A[i][j]); } printf ("\n"); } printf ("B矩陣: \n"); for (i = 0; i < size; i++){ for (j = 0; j < size; j++){ printf ("%d ", B[i * size + j]); // printf ("%d ", B[i][j]); } printf ("\n"); } #endif /* 單線程 */ gettimeofday (&t1, NULL); for (i = 0; i < size; i++) for (j = 0; j < size; j++) for (k = 0; k < size; k++){ result[i * size + j] += A[i * size + k] * B[k * size +j]; // result[i][j] += A[i][k] * B[k][j]; } gettimeofday(&t2, NULL); timeuse = t2.tv_sec - t1.tv_sec + (t2.tv_usec - t1.tv_usec)/1000000.0; #ifdef LOG printf ("單線程結果矩陣: \n"); for (i = 0; i < size; i++){ for (j = 0; j < size; j++){ printf ("%d ", result[i * size + j]); // printf ("%d ", result[i][j]); } printf ("\n"); } #endif printf("單線程耗時: %fs\n", timeuse); sprintf (timeChars, "%lf, ", timeuse); write (fd, timeChars, strlen (timeChars)); for (i = 0; i < size; i++) for (j = 0; j < size; j++){ result[i * size + j] = 0; result2[i * size + j] = 0; result3[i * size + j] = 0; result4[i * size + j] = 0; /* result[i][j] = 0; result2[i][j] = 0; result3[i][j] = 0; result4[i][j] = 0; */ } /* 單線程 end */ /* 雙線程 */ gettimeofday (&t1, NULL); pthread_create (&tid2[0], NULL, (void *)twoThread1, NULL); pthread_join (tid2[0], NULL); pthread_create (&tid2[1], NULL, (void *)twoThread2, NULL); pthread_join (tid2[1], NULL); gettimeofday (&t2, NULL); timeuse = t2.tv_sec - t1.tv_sec + (t2.tv_usec - t1.tv_usec)/1000000.0; #ifdef LOG printf ("雙線程結果矩陣: \n"); for (i = 0; i < size; i++){ for (j = 0; j < size; j++){ printf ("%d ", result2[i * size + j]); // printf ("%d ", result2[i][j]); } printf ("\n"); } #endif printf("雙線程耗時: %fs\n", timeuse); sprintf (timeChars, "%lf, ", timeuse); write (fd, timeChars, strlen (timeChars)); for (i = 0; i < size; i++) for (j = 0; j < size; j++){ result[i * size + j] = 0; result2[i * size + j] = 0; result3[i * size + j] = 0; result4[i * size + j] = 0; /* result[i][j] = 0; result2[i][j] = 0; result3[i][j] = 0; result4[i][j] = 0; */ } /* 雙線程 end */ /* 三線程 */ gettimeofday (&t1, NULL); pthread_create (&tid3[0], NULL, (void *)threeThread1, NULL); pthread_join (tid3[0], NULL); pthread_create (&tid3[1], NULL, (void *)threeThread2, NULL); pthread_join (tid3[1], NULL); pthread_create (&tid3[2], NULL, (void *)threeThread3, NULL); pthread_join (tid3[2], NULL); gettimeofday (&t2, NULL); timeuse = t2.tv_sec - t1.tv_sec + (t2.tv_usec - t1.tv_usec)/1000000.0; #ifdef LOG printf ("三線程結果矩陣: \n"); for (i = 0; i < size; i++){ for (j = 0; j < size; j++){ printf ("%d ", result3[i * size + j]); } printf ("\n"); } #endif printf("三線程耗時: %fs\n", timeuse); sprintf (timeChars, "%lf, ", timeuse); write (fd, timeChars, strlen (timeChars)); for (i = 0; i < size; i++) for (j = 0; j < size; j++){ result[i * size + j] = 0; result2[i * size + j] = 0; result3[i * size + j] = 0; result4[i * size + j] = 0; /* result[i][j] = 0; result2[i][j] = 0; result3[i][j] = 0; result4[i][j] = 0; */ } /* 三線程 end */ /* 四線程 */ gettimeofday (&t1, NULL); pthread_create (&tid4[0], NULL, (void *)fourThread1, NULL); pthread_join (tid4[0], NULL); pthread_create (&tid4[1], NULL, (void *)fourThread2, NULL); pthread_join (tid4[1], NULL); pthread_create (&tid4[2], NULL, (void *)fourThread3, NULL); pthread_join (tid4[2], NULL); pthread_create (&tid4[3], NULL, (void *)fourThread4, NULL); pthread_join (tid4[3], NULL); gettimeofday (&t2, NULL); timeuse = t2.tv_sec - t1.tv_sec + (t2.tv_usec - t1.tv_usec)/1000000.0; #ifdef LOG printf ("四線程結果矩陣: \n"); for (i = 0; i < size; i++){ for (j = 0; j < size; j++){ printf ("%d ", result4[i * size + j]); } printf ("\n"); } #endif printf("四線程耗時: %fs\n", timeuse); sprintf (timeChars, "%lf\n", timeuse); write (fd, timeChars, strlen (timeChars)); for (i = 0; i < size; i++) for (j = 0; j < size; j++){ result[i * size + j] = 0; result2[i * size + j] = 0; result3[i * size + j] = 0; result4[i * size + j] = 0; /* result[i][j] = 0; result2[i][j] = 0; result3[i][j] = 0; result4[i][j] = 0; */ } /* 四線程 end */ } // 釋放空間 free (A); free (B); free (result); free (result2); free (result3); free (result4); A = NULL; B = NULL; result = NULL; result2 = NULL; result3 = NULL; result4 = NULL; // 關閉文件 close (fd); return 0; } 

編譯

gcc cpu_mul.c -O3 -o cpu_mul 

出現問題

undefined reference to `pthread_create' 

問題的原因：pthread不是linux下的默認的庫，也就是在鏈接的時候，無法找到phread庫中哥函數的入口地址，於是鏈接會失敗。

解決：在gcc編譯的時候，附加要加 -lpthread參數即可解決。

再編譯

gcc cpu_mul.c -O3 -lpthread -o cpu_mul 

執行

./cpu_mul
當前階數: 200
單線程耗時: 0.068800s
雙線程耗時: 0.064988s
三線程耗時: 0.065114s
四線程耗時: 0.065234s
當前階數: 400
單線程耗時: 0.616225s
雙線程耗時: 0.592269s
三線程耗時: 0.590107s
四線程耗時: 0.590583s
當前階數: 600
單線程耗時: 2.106992s
雙線程耗時: 2.067765s
三線程耗時: 2.078379s
四線程耗時: 2.079277s
當前階數: 800
單線程耗時: 5.433704s
雙線程耗時: 5.232933s
三線程耗時: 5.241303s
四線程耗時: 5.245513s
當前階數: 1000
單線程耗時: 11.862752s
雙線程耗時: 11.370975s
三線程耗時: 11.380489s
四線程耗時: 11.376795s
當前階數: 1200
單線程耗時: 21.028770s
雙線程耗時: 20.471958s
三線程耗時: 20.481492s
四線程耗時: 20.683800s
當前階數: 1400
單線程耗時: 33.121682s
雙線程耗時: 30.445559s
三線程耗時: 30.390110s
四線程耗時: 30.387628s
當前階數: 1600
單線程耗時: 57.306933s
雙線程耗時: 56.083187s
三線程耗時: 56.102355s
四線程耗時: 56.104007s
當前階數: 1800
單線程耗時: 70.554218s
雙線程耗時: 64.973096s
三線程耗時: 64.993218s
四線程耗時: 64.977718s
當前階數: 2000
單線程耗時: 98.518510s
雙線程耗時: 94.792153s
三線程耗時: 94.810620s
四線程耗時: 94.876091s

然而從htop里面看的多線程時實際只有兩個線程，有一個線程在工作，另一個在sleep，很是奇怪

編譯時去掉優化級別

gcc cpu_mul.c -lpthread -o cpu_mul2 

當前階數: 200
單線程耗時: 0.236366s 雙線程耗時: 0.276663s 三線程耗時: 0.352680s 四線程耗時: 0.365158s 當前階數: 400 單線程耗時: 2.031096s 雙線程耗時: 2.332579s 三線程耗時: 2.967604s 四線程耗時: 3.046181s 當前階數: 600 單線程耗時: 7.005228s 雙線程耗時: 7.990198s 三線程耗時: 10.129884s 四線程耗時: 10.384419s 當前階數: 800 單線程耗時: 16.948757s 雙線程耗時: 19.568093s 三線程耗時: 24.364115s 四線程耗時: 24.274340s 當前階數: 1000 單線程耗時: 38.138426s 雙線程耗時: 42.429860s 三線程耗時: 52.424720s 四線程耗時: 53.720896s 

從htop里面看也是兩個線程，一個CPU占用100%，另一個CPU占用0%，在圍觀......結果未優化的性能更差，同樣只有兩個線程，且一個線程在圍觀。。。真的是醉了。。。有機會再研究多線程的編程吧！
（3）GPU 器件參數輸出

vim device.cu nvcc device.cu -o device 

#include <stdio.h> int main() { int nDevices; cudaGetDeviceCount(&nDevices); for (int i = 0; i < nDevices; i++) { cudaDeviceProp prop; cudaGetDeviceProperties(&prop, i); printf("Device Num: %d\n", i); printf("Device name: %s\n", prop.name); printf("Device SM Num: %d\n", prop.multiProcessorCount); printf("Share Mem Per Block: %.2fKB\n", prop.sharedMemPerBlock / 1024.0); printf("Max Thread Per Block: %d\n", prop.maxThreadsPerBlock); printf("Memory Clock Rate (KHz): %d\n", prop.memoryClockRate); printf("Memory Bus Width (bits): %d\n", prop.memoryBusWidth); printf("Peak Memory Bandwidth (GB/s): %.2f\n\n", 2.0 * prop.memoryClockRate * (prop.memoryBusWidth / 8) / 1.0e6); } return 0; } 

結果：

./device
Device Num: 0
Device name: NVIDIA Tegra X1
Device SM Num: 1
Share Mem Per Block: 48.00KB
Max Thread Per Block: 1024
Memory Clock Rate (KHz): 12750
Memory Bus Width (bits): 64
Peak Memory Bandwidth (GB/s): 0.20

(4)GPU 矩陣乘法

//GPU 矩陣乘法 #include <stdio.h> #include <stdlib.h> #include <sys/time.h> #include <unistd.h> #define w 2000 struct Matrix { int width; int height; float *elements; }; __device__ float getElement(Matrix *A, int row, int col) { return A->elements[row * A->width + col]; } __device__ void setElement(Matrix *A, int row, int col, float value) { A->elements[row * A->width + col] = value; } __global__ void matMulKernel(Matrix *A, Matrix *B, Matrix *C) { float Cvalue = 0.0; int row = threadIdx.y + blockIdx.y * blockDim.y; int col = threadIdx.x + blockIdx.x * blockDim.x; for (int i = 0; i < A->width; ++i) { Cvalue += getElement(A, row, i) * getElement(B, i, col); } setElement(C, row, col, Cvalue); } int main() { int width = w; int height = w; Matrix *A, *B, *C; cudaMallocManaged((void**)&A, sizeof(Matrix)); cudaMallocManaged((void**)&B, sizeof(Matrix)); cudaMallocManaged((void**)&C, sizeof(Matrix)); int nBytes = width * height * sizeof(float); cudaMallocManaged((void**)&A->elements, nBytes); cudaMallocManaged((void**)&B->elements, nBytes); cudaMallocManaged((void**)&C->elements, nBytes); A->height = height; A->width = width; B->height = height; B->width = width; C->height = height; C->width = width; for (int i = 0; i < width * height; ++i) { A->elements[i] = 1.0; B->elements[i] = 2.0; } dim3 blockSize(32, 32); dim3 gridSize((width + blockSize.x - 1) / blockSize.x, (height + blockSize.y - 1) / blockSize.y); struct timeval t1,t2; gettimeofday(&t1,NULL); double timeuse; matMulKernel << < gridSize, blockSize >> >(A, B, C); cudaDeviceSynchronize(); gettimeofday(&t2,NULL); timeuse = t2.tv_sec - t1.tv_sec + (t2.tv_usec - t1.tv_usec)/1000000.0; printf("Use Time:%fs\n", timeuse); return 0; } 

vim gpu_mul.cu
nvcc gpu_mul.cu  -o gpu_mul

Use Time:1.466122s (w=2000) Use Time:267.060169s (w=8000) 

（5）結果對比

通過對矩陣乘法運算的結果對比，同樣2000階的矩陣乘法，單核CPU需要52.64s，而GPU只需要1.466s，矩陣乘法運算速度遠快於CPU，另外編程時也要根據其GPU編程模型來優化

參考資料

• NVIDIA Jetson TX2 查看系統參數狀態
• CUDA編程入門
• CUDA編程(四): CPU與GPU的矩陣乘法對比