x64下安裝OpenCV2.3+CUDA

災難！
一定有人問，Bee為何裝完馬上升級。
因為，Bee要的功能是壞的啊！查了一下，發現OpenCV2.2滿是Bug。
Bee想主因是一口氣加入太多東西，GPU和X64一次加入的結果。
然後在一連串找尋中，Bee發現了OpenCV2.3可以解決問題，然後一看，11小時前更新，還是熱的啊！

OpenCV2.3需要CUDA4.0及VS2010，天啊！一切要重裝。

好! 就重裝。那就先移除CUDA3.2及VS2008。
接下來安裝顯示卡Driver，然後重開機。
再來是裝VS2010、SP1及Nsight2.0。
然後是CUDA SDK4.0，接下來就有問題了。

竟然無法編譯！少了cutil64D.lib。查了一下，好像要自己編出來。Bee邊查邊找就半小時過去了。
去\NVIDIA GPU Computing SDK 4.0\CUDALibraries下打開CUDALibrariesSDK_vs2010.sln編一次。
總算得到cutil64D.lib，然後移去要用的地方。
再來又少了shrUtils64D.lib，而且缺檔案stopwatch.cpp及stopwatch_win.cpp。
這二個檔在\NVIDIA GPU Computing SDK 4.0\C\common\src。
先從專案中移除錯的連接，再加入正確的。
總算可以用了。
最後把x64、W32、Debug及Release全部各編一次，放好備用。

可以裝OpenCV2.3了。
使用OpenCV-2.3.0-win-superpack.exe解開來。
用CMake重做一次，這次就沒有問題了。

沒想到這次是CUDA比較難裝。
不過x64模式OpenCV+CUDA總算全部搞定。

OpenCV2.2+CUDA x64模式編譯成功

使用OpenCV2.2版，用CUDA 3.2版。

因為改用CMake所以很不習慣。

裝了半天，才發現需要NPP函式庫。
NPP函式庫也要用64位元。

用VS2010編也不行，還退回去VS2008。

總算把hog_gpu給編出來了。


再來，要經由CMake重做。變得太快，不知要弄多久。

OpenCV2.1+CUDA 64位元整合：結果失敗

  最近Bee換筆記電腦是具有GeForce GT 420M的顯示卡。
  現在都是安裝Windows 7 x64的版本。故Bee下載CUDA 64位元回來安裝。
  到這裡都沒有問題。只有CUDA 64部分設定要自己手工調整。
 
  之後有許多工具都很不習慣，花了不少時間去找。整個Windows和XP實在差太多了。
  另外有一些其他奇怪的地方有些程式找不到裝置，原來還有UAC的問題。
  好吧！看在64位元可以不受4GB限制，還是去適應好了 。

  弄了數週，才想到回來看看CUDA程式。
  結果，CUDA 64無法和Win32的OpenCV做Link。
  而OpenCV沒有64位元的Library。那只好自己編函式庫了。
 
  奇怪的是OpenCV2.1明明就有寫支援64位元，但一直編不出可以用的函式庫。
  查過在其他平台都是可以用的。但在VS2008及VS2010就有問題。
  沒錯！就是這個問題。但....沒有人解成功。
 
  Bee又安裝了好幾次，沒一次成。查了很久，發現是沒有載入該有的函式庫。
  為何！M$的C++老是玩這種，每次編出來的程式都很難搬。
 
  最後沒辦法，回去Win32。安裝CUDA 32然後Link OpenCV，就過了。
  再等OpenCV下一版再看看。
  不過64位元整合算是失敗了。看來時代還沒有到。

  還有很多應用軟體也都是在Win32模式下，沒有幾套是64位元。
  要換到64位元，看來還是不容易取得優勢。反而是環境大改，真是不習慣。

最近CUDA程式上的進展

最新主要是改寫別人的CUDA程式。Open Source前幾版真是Bug百出，想找要好改的還要各方比較。

改寫除了解Bug外，也有一些問題要解決。

1.資源調度

  最常拿到的是CUDA硬體1.3版的程式，剛好Bee又想用筆電跑，就要改成硬體1.1版。

  這種狀況就要會做資源調度的程式改寫了。


2.改C++程式

  Bee是C的使用者，C++不熟。不過遇到的算是原本CPU程式轉成CUDA之後造成的問題。

  今天抓一個CUDA記憶體取用爆炸問題。追出來問題是長這樣：

  原先使用物件只有建構函式，改寫為CUDA時，要在建構函式中取用CUDA記憶體，將資料轉到CUDA去。

  後面也有其他函式使用CUDA上的資料運算。

  處理單張照片時沒事，放到Webcam執行時，就產生記憶體不足。Bee也感到奇怪，1G的記憶體不夠用。

  拿GPU-Z來看，真的用了1G。主要是每次處理照片，就會多出一些記憶體，是累積到爆的。

  看程式還看不出問題，後來才發現，原來是建構時取用的CUDA記憶體沒釋放。

  因為C++在物件使用完畢後，會自己釋放掉記憶體，所以很多人不寫解構函式。

  可是改寫CUDA時，一樣在建構時取用CUDA的記憶體，自動解構時只會釋放CPU側。

  而且相關的指標會釋放，在CUDA取用內的記憶體就變成沒人管的記憶體。

  Bee加入解構，其內釋放CUDA記憶體，就這樣解決了。

  CUDA還是用C的記憶體操作，不會自動回收。直接使用C++程式移植來的，這點還是要注意。

工作流模式(workflow pattern)相關資料

為了找出並行計算上的各種狀況，需要找到可以表達的模型。

目前已知並行狀態機可以使用Petri Net來表達。

在找Petri Net資料中，找到了工作流模式(workflow pattern)。

也就是所有並行式工作流，都可以使用工作流模式內的模型來表達。


工作流模式基本模型有21種：

1. 順序（Sequence）

2. 平行拆分（Parallel Split）

3. 同步（Synchronization）

4. 排他選擇（Exclusive Choice）

5. 單合併（Single Merge）

6. 多選（Multi-choice）

7. 平行合併（Synchronize Merge）

8. 多合併（Multi-merge）

9. 鑒別器（Discriminator）

10. M中的N模式（N-out-of-M Join）

11. 強制循環（Arbitrary Cycles）

12. 隱式終止（Implicit Termination）

13. 非同步的多實例（Multiple Instances Without Synchronization）

14. 在設計期間預先確定的多實例（Multiple Instances With a Priori Design Time Knowledge）

15. 在運行期預先確定的多實例（Multiple Instances With a Priori Runtime Knowledge）

16. 無法在運行期預先確定的多實例（Multiple Instances Without a Priori Runtime Knowledge）

17. 延遲選擇（Deferred Choice）

18. 交替平行路由（Interleaved Parallel Routing）

19. 里程碑（Milestone）

20. 取消活動（Cancel Activity）

21. 取消實例（Cancel Case）


可以到
http://is.ieis.tue.nl/research/patterns/patterns.htm

上面有動畫實例，可以很容易的了解各模型的差異。

有了這個工作流模式，只要把各狀況找出解決方式，就不會用自己想的怪方法來解，然後遇到不明的問題。

多線程退出算法

取自"多核計算與程序設計"

第二章重點，主要在三張圖。

平行任務分層算法

取自"多核運算與程序設計"

1.先計算任務圖中所有頂點的入度

2.找出所有入度為0的頂點，放入第0層，這樣便得到一個分層

3.假設已得到第K個分層，考慮去除放入0~k層頂點外，其他剩下的頂點所組成的子圖，在子圖中尋找所有入度為0的頂點，放入第K+1層中。

4.令K=K+1，重覆步驟3，直到所有頂點都被放入分層中。

這對RTOS也是很有用的。

"多核計算與程式設計"介紹

1.讀書原由

  在CUDA程式設計上面遇到許多多核心程式計算的問題。個人從單核心作業系統一下子變成多核心程式設計，才發覺對於多核心知識的不足。

  而CPU多核心已流行數年，在新的PC硬體上，要使用雙核心已是基本配備。意指，多核心已經是免費午餐。

  但在軟體設計上，卻很少使用到雙核心，甚至四核心來做為加速。為何會如此?轉換程式是如此慢，其原因為何?

  在收集網路資料後發現，出在使用語言上的問題比作業系統支援多核心來得嚴重。但不管如何，中文資料一樣貧乏。


  a.作業系統支援及其問題
    過去：

    作業系統支援多核心是最早的，在Unix作業系統就已發展出支援多核心的方式。

    但是現代作業系統都有支援的狀況下，為何沒有什麼程式設計師願意使用?原因在於使用的門檻高，且其經濟效應不好。

    因為多核心程式在單核心的CPU運行效率不好，在市場尚未普及前，多數工程師仍未接受多核心程式訓練。

    現在：

    但現在環境已經不同了，多核心時代真的到來。


  b.電腦語言支援多核心問題
    過去：

    其實這才是目前多核心程式發展的最大阻礙。撰寫多核心程式只能使用作業系統支援的狀況下，困難度太高。

    而大部分使用者所使用的語言皆為對於單核心所設計。能使用在多核心的電腦語言則是太少。

    所以大部分程式設計仍然被單核心程式語言所限住。這才是多核心程式真正無法流行的原因。

    因為程式語言就是要讓使用者方便處理設計，若是學習門檻太高，就不易推行。

    現在：

    語言延伸：MPI、OpenMP
    平行語言：Erlang、Scala


  c.使用CUDA之後才發現的問題

    CUDA的推行以資料為主的平行計算。和之前在作業系統所提的工作為主的平行計算上有很大的不同。

    在超過八核心同時運算時，為了要增加產出率，勢必要採行以SIMD為主的語言架構，也就是資料為主的程式概念。

    但個人發現，可以找到的資料不多，所以轉以找多核心程式的資料，所以找到此書。


2.書本結構

  a.基礎知識
    多核計算概述
    多線程編程基礎
    OpenMP程式設計

  b.基礎資料結構及算法
    結構
    Link List
    Hash
    Tree
    AVL search tree
   
  c.平行運算法
    並行程式設計模式
    並行搜尋
    並行排序
    並行數值計算
   
  d.共享資源分散式計算
    分散式計算設計模式
    分散式陣列
    分散式查找
    分散式記憶體管理
   
  e.任務分解及調度
    任務圖分解及調度
    動態任務分解及調度
    Lock-Free編程基礎


3.比對CUDA

  a.有許多問題是一樣的

  b.CUDA中有些問題解法未表明產生原因，在多核心中則有解釋

  c.CUDA可視為多核心的一種實現

  d.可以有效利用CUDA中的原子函式

  e.免除遇到多核心問題，無從看出問題產生原因

GTX275基本測試資料

發現沒有放GTX275的測試資料，補一下。

PC配備：
CPU
AMD Phenom(tm) 9350e
Quad-Core Processor 2.00 GHz

O.S.
Windows XP SP3

RAM
DDR2-800 2G*4(但O.S.只有顯示使用3.25G)

測試圖：

不過GTX275現在買不太到了，可能也要絶版了。現在GTX260以上的卡都不好買。
也許是缺貨加上世代交替的影響。再來玩的CUDA都會換新卡了吧！

使用GPU-Z看GPU狀況

想知道GPU有多操，所以找來這個軟體看看。結果還很好玩。

主要畫面：

在玩Game時，還可以看到不同形態的Game具有不同影響。
也可以觀察溫度的狀況，來了解功耗。

其他檢測：

不過有些資訊也是有這個GPU-Z才得以解釋。
像是在CUDA中測效能時，必須利用GPU內部timer來測。但有一行程式一直無法理解，其註解寫著Warm up。
結果一看到GPU-Z去監視CUDA運行時的利用率，我馬上理解出了什麼事。

原來GPU沒事要進入省電狀況，主要是調整Clock的方式來做。
在玩Game時會動態調整clock及負載狀況。
而在使用CUDA模式時會直接跳到最高的clock。所以需要一段時間做Warm up。

GTX275測得資料：

GTX275感測器狀況：

可以發現多了幾個溫度感知。

GTX275執行CUDA時狀況：

GTX275玩Game時的狀況

沒想到Game沒有操到滿。

GTX275執行完CUDA後很快回復到省電模式。

追加在ION上測到的資料

CUDA新書到

Bee之前都是自己收集資料，看原文去理解。所以差不多是遇到什麼再去查。
終於有本比較集中資料的書了。常用的運算法也都有解釋。雖然Bee也查的差不多了。
不過買了除了自己查外，拿來教人也是真的。
隨著CUDA運算的流行，實驗室影像組的CUDA人員正在擴張中。

除了影像運算外，Bee本身就是控制組的人，也很想用到Robot控制上。
其中最大的原因是GPGPU的價格低到可以和影像處理的DSP差不多了。
一台使用Atom+ION的NetPC也很合適於控制Robot。
看看控制組何時也想用用CUDA了。

2009年參賽獎狀

假共享(False Sharing)問題

在看CUDA Programming Guide時，有一大篇討論Shared Memory Access Patterns with Bank Conflicts問題。這個一直不太了解。

後來找多核心文章，發現假共享(False Sharing)問題。仔細看，這個不就是CUDA裏面提的問題。
由此Bee之前推測的CUDA中的Share memory是軟體管理的cache，這個論點是確認無誤了。因為使用一樣的硬體，才會有一樣的問題，連解法都一樣。

好在之前沒有將share memory做為寫入用，所以問題不大。
之後要小心考量使用share memory的寫入問題了。

CUDA計算積分數列的方法

使用平行運算計算總和，還可以理解。但要算積分，不符合平行運算的方法，這下要如何?
積分數列的數學問題是這樣的：
數列 X=[ a0 a1 a2 ... aN ]
要算出 Y=[ a0 a0+a1 a0+a1+a2 ... a0+a1+a2+...+aN ]

其實有注意到的人會發現，為何這個問題Bee要用平行運算。因為這就是圖像積分法(integral image)的基本運算。因為要移進GPGPU來做，問題就來了。
因為每一個數列元素要等上一個元素的值出現才能算，可是平行運算要沒有相關才能用，這下子有解嗎?
找了許多文章，一張圖解了這個問題。

果然還是因為沒有收集平行運算的演算法，才會不知道往那裡去找。
找到了之後，發現這個動作可以做以下的計算：
Radix sort
Quicksort
String comparison
Lexical analysis
Stream compaction
Sparse matrices
Polynomial evaluation
Solving recurrences
Tree operations
Histograms
看來是很重要的演算法，值得好好去了解。

CPU和GPU的合作，和大腦分工相似。

Bee是搞自動控制的，有一個問題一直是系統不好選擇的。到底是要用分散式處理器來做系統運算，還是用PC做中央控制。
分散控制也是有好處的，系統加元件是很快，因為關係單純，但程式不好改，要改許多處理器程式。
中央控制則是要經過中央管理系統，比較複雜，但程式好改。

這個問題到Bee開始處理影像相關問題，才知道要如何選。
原因在於影像處理無法使用分散運算的方式，若是真的要用，那訊息傳輸系統一定一直在傳影像相關資料，反而是一種浪費。

不過一開始Bee處理影像控制是想用FPGA來做，但它的記憶體不足，必須外掛SDRAM。
規劃到一半GPGPU興起，完全符合影像處理的需求，我就直接轉到GPGPU的領域去。
其實對於影像處理，其瓶頸在於運算相同，但資料量很大，平行計算才能有效解決。
但資料之間有相關，無法分開存放，要使用集中式的存放才能有效運算。

只是理解之後，Bee在想，一樣的工程問題，生物界是如何解的?
果然，大腦的發展，差不多和眼睛的出現是相關的。不是一個感光細胞的眼睛，是一群感光細胞組成的眼睛。

Sensor Array的出現，就是引入大量平行處理的開始。
在multi-sensor的狀況下，腦不需很大，大概就處理環境記憶和時間相關資訊。
但Sensor Array就不同了，記憶體需求量大，且運算力大增，才可做出分析。

在玩CUDA之後，發現並非所有的問題都合適GPU運算。在做統合判定的地方，CPU的運算更合適。
而且CUDA的異質運算結構，可以比OpenMP更容易去選擇運算結構。簡單的說何時要用CPU運算，何時要做GPU運算，是很容易分的。
但OpenMP不是如此明顯，但還是有平行運算的味道在。

最近看了"你腦內的兩個世界"
http://blog.xuite.net/unlimiter1001/unlimiter/29032221
了解兩個半腦的分工，更可以確定人工智能真的需要並行處理器及串行處理器的配合。
腦之所以二個完全分開，正因為腦的運算為成長型。右腦專精於平行處理，左腦專精於串行處理。二種迴路基本型式必然完全不同，所以分成二個腦來成長。
也可以說生物界對於平行運算和串行運算，沒有做出統一的解，所以長出二個處理器來做。
那Intel想用X86來做GPU，想必無法有效和真正的GPU來競爭。果不其然Larrabee的計劃真的失敗。

但並行運算就算是解了影像處理問題，但後繼串行處理的部分也是會增加運算量。原因是資訊也跟著變多，串行處理的資料及儲存也是要跟著增加上去。
所以左右腦的大小一直都是一樣，可見得二者重要性誰也沒有贏過誰。
意指，就算CPU和GPU就算合成一個晶片來做通用運算，二者電路的比例可能仍是很相近。那關於這個預測，要等GPGPU運算普及才會有人探討。

那對於現在，Bee的問題出在沒有收集過平行運算的運算法。因為才開始流行，資料不多。但對於軟體發展，這是無法避免的。
只能開始注意及了解相關的發展了。

使用積分影像法(integral image)取代金字塔圖層

做影像處理常用到金字塔圖層的方式做影像加速處理的方法。
但金字塔圖層會多用掉記憶體，使用量剛好是原圖的一半。
後來發現有人使用積分影像法(integral image)，它可以同時代表原圖及金字塔圖層。
因為可以做不同大小的差分而獲得所需的圖。
只要取點時算A-B-C+D就好了。

如何將影像處理程式改寫為CUDA程式

這裡提出我遇到問題時找到的原則：

1.保留原程式結構，先找出最多層for的地方，先將這個函式GPU化。

2.找出要GPU化的函式所要的輸入及輸出資料。從個別處理改為表格化資料，因為要展成並行處理。

3.找出各for對應的並行層。
  就是將各for展成threadIdx及BlockIdx。
  和圖檔長寬相關的for主要以BlockId.x及BlockId.y做分配。因為這二個系統變數可用範圍最大。

4.先產生CPU可執行函式取代，用以分離函式準備GPU化。但可以驗證及除錯。

5.驗證成功後，GPU化的程式再試。

6.針對GPU結構開始最佳化。在CPU的GPU版本上做除錯及驗證。

7.完成關鍵的CPU轉GPU函式，再找需要做的部分進行合併。

實際上執行仍有許多問題要克服，解決了再整理。

CUDA工作排程的了解

寫了一點CUDA程式，但對於寫出最佳程式並沒有什麼好辦法去了解，難到只有實測一途了嗎?
當然不是，可以找一下CUDA SDK中有一個試算表CUDA_Occupancy_calculator.xls。
玩了試算表，就可以稍為了解CUDA指令排程，其實和發動的執行時所使用資源有關。

可以發現只有三項資源可以調整。
1. 每一個Block使用多少Threads
2. 每一個Thread使用多少Registers
3. 每一個Block使用多少Shared memory

使用CUDA寫程式時對於每一個Block使用多少資源是很明確的，但對於每一個Thraeds使用多少Registers則不是可以直接得到。
不過調整一下數值可以發現CUDA的排程限制。

從文件可以知道一次CUDA指令可以有32個核心動作，這個叫Wraps。在不同核心架構下一次可以發動的Wraps數目不同。
因為工作上是以Block為主要規劃，所以Threads的數目除以32就是Wraps數。
像Bee用的G9600M最大的Wraps數目是24，24是不太好的數字，因為Bee寫的程式動作大部分和2的倍數有關，所以常常無法有效使用。

但不是可以使用的Threads多，Wraps數目就會多，因為因為每一個Threads都是一個core task會用掉Registers。對每一個multiprocessor來說，Registers數目是有限的。
所以就會有Registers總數限制，就是第二項要控制的。

通常一個Block會有共有的變數，此時就要使用到Shared memory，這項資源為每一個multiprocessor才有，所以又變成另一個限制。
由此可以看到一件事，雖然一個multiprocessor有八個core，但可以執行的Threads數目不是八。
可能Threads數目是32，八個core要分成四次執行。
另一種可能是因資源用完，32個Threads是分成4個multiprocessor去做。
若是資源使用很多，有可能每一個multiprocessor只能使用4 core，這時1個wraps會佔用更多的multiprocessor。

可以見得使用比較好的顯卡，因為Multiprocessor多，所以可以解比較耗資源的threads task。
以上為玩試算表所得的心得，實際上可能要實驗才能證實。

玩過後，發現程式工作如果分得比較細小的程式段，會有助於執行分配。
但要分到多細小，真的就要用這個試算表下去估了。

如何了解__syncthread()的使用時機：試誤法

在使用CUDA shared memory常會使用__syncthread()來避免資料不同步處理造成的破壞。問題是一般人無法判定何時會產生這個問題，所以會先加以做為保險。

可是加了之後確實會影響執行速度，Bee這邊測到的就差了四倍多。
所以要適當加入__syncthread()才是合理。

剛好之前bee使用CPU及GPU同時開發程式的方法，因為一邊為single-thread另一個為multi-thread，所以可以保證single-thread是正確的。剛好可以比對結果。

若是不加__syncthread()會使結果比對產生錯誤，那就要加。沒有錯誤就照CPU版本的直接移入GPU版本。

這種情況玩個幾次，基本上就可以知道何時要加，何時不要加。

第二個CUDA程式，使用Shared memory

因為有共同讀取的來源資料，所以先行載入shared memory。然後經過運算，之後回存。
因為Windows下的CUDA沒有除錯器，所以先發展CPU下可以執行的GPU架構程式。先將變數之間的關係弄清楚再移進GPU。
至於shared memory也是一種區域性變數，所以用local variable來模擬。

CPU版本程式：
void CPU_mask_rotate(unsigned int *src, unsigned int *dest, int gridDim_x, int gridDim_y)
{
    unsigned int idata[HEAD_LENGTH];

    for(int blockIdx_y=0; blockIdx_y < gridDim_y ; blockIdx_y++){
        for(int blockIdx_x=0; blockIdx_x < gridDim_x; blockIdx_x++){
            // Read to local var.
            for(int i=0; i<HEAD_LENGTH; i++){
                int t = (blockIdx_y*HEAD_LENGTH+i)*gridDim_x + blockIdx_x;
                idata[i] = src[t];
            }
            // convert format
            for(int threadIdx_x=0; threadIdx_x<TRANSPOSE_SIZE; threadIdx_x++){
                for(int j=0; j<HEAD_BLOCK; j++){
                    unsigned int r;
                    r=0;
                    for(int k=0; k<TRANSPOSE_SIZE; k++){
                        unsigned int u,v;
                        u = idata[j*TRANSPOSE_SIZE + k];
                        v = (u>>threadIdx_x & 0x1)<<k;
                        r |= v;
                    }
                    int t = (blockIdx_y * gridDim_x + blockIdx_x)*HEAD_LENGTH + threadIdx_x*HEAD_BLOCK + j;
                    dest[t] = r;
                }
            }
        }
    }
}

GPU版本程式
__global__ void Kernelmask_rotate(unsigned int *src, unsigned int *dest)
{
    __shared__ unsigned int idata[HEAD_LENGTH];

    for(int i=0; i<HEAD_LENGTH; i++){
        int t = (blockIdx.y * HEAD_LENGTH+i)*gridDim.x + blockIdx.x;
        idata[i] = src[t];
    }
    __syncthreads();

    for(int j=0; j<HEAD_BLOCK; j++){
        unsigned int r;
        r=0;
        for(int k=0; k<TRANSPOSE_SIZE; k++){
            unsigned int u,v;
            u = idata[j*TRANSPOSE_SIZE + k];
            v = (u>>threadIdx.x & 0x1)<<k;
            r |= v;
        }
        int t = (blockIdx.y * gridDim.x + blockIdx.x)*HEAD_LENGTH + threadIdx.x*HEAD_BLOCK + j;
        dest[t] = r;
    }
}

CPU寫好時，就改一下變數名( _改為. )就可以copy到GPU的版本。

只是在發展中發現一個奇怪現象。
因為讀入shared memory可以使各thread少去對global memory的讀取，而減少執行時間。
但在寫程式的時候，去算各位址不好判定，所以先用一塊shared memory存結果，再從shared memory存回global memory，因為這樣比較好寫。
後來發現CUDA SDK內的example都是直接回存。這樣可以省去shared memory，不過就是不好寫。