GPU 的基本介紹

GPU 的硬體架構
 

這裡我們會簡單介紹，NVIDIA 目前支援 CUDA 的 GPU，其在執行 CUDA 程式的部份（基本上就是其 shader 單元）的架構。這裡的資料是綜合 NVIDIA 所公布的資訊，以及 NVIDIA 在各個研討會、學校課程等所提供的資料，因此有可能會有不正確的地方。主要的資料來源包括 NVIDIA 的 CUDA Programming Guide 1.1、NVIDIA 在 Supercomputing '07 介紹 CUDA 的 session，以及 UIUC 的 CUDA 課程。 GPU 的基本介紹 目前 NVIDIA 推出的顯示晶片，支援 CUDA 的是 G80 系列的顯示晶片。其中 G80 顯示晶片支援 CUDA 1.0 版，而 G84、G86、G92、G94、G96 則支援 CUDA 1.1 版。基本上，除了最早的 GeForce 8800 Ultra/GTX 及 320MB/640MB 版本的 GeForce 8800GTS、Tesla 等顯示卡是 CUDA 1.0 版之外，其它 GeForce 8 系列及 9 系列顯示卡都支援 CUDA 1.1。詳細情形可以參考 CUDA Programming Guide 1.1 的 Appendix A。 所有目前支援 CUDA 的 NVIDIA 顯示晶片，其 shader 部份都是由多個 multiprocessors 組成。每個 multiprocessor 裡包含了八個 stream processors，其組成是四個四個一組，也就是說實際上可以看成是有兩組 4D 的 SIMD 處理器。此外，每個 multiprocessor 還具有 8192 個暫存器，16KB 的 share memory，以及 texture cache 和 constant cache。大致上如下圖所示： 詳細的 multiprocessor 資訊，都可以透過 CUDA 的 cudaGetDeviceProperties() 函式或 cuDeviceGetProperties() 函式取得。不過，目前還沒有辦法直接取得一個顯示晶片中有多少 multiprocessor 的資訊。 在 CUDA 中，大部份基本的運算動作，都可以由 stream processor 進行。每個 stream processor 都包含一個 FMA（fused-multiply-add）單元，可以進行一個乘法和一個加法。比較複雜的運算則會需要比較長的時間。 執行過程 在執行 CUDA 程式的時候，每個 stream processor 就是對應一個 thread。每個 multiprocessor 則對應一個 block。從之前的文章中，可以注意到一個 block 經常有很多個 thread（例如 256 個），遠超過一個 multiprocessor 所有的 stream processor 數目。這又是怎麼回事呢？ 實際上，雖然一個 multiprocessor 只有八個 stream processor，但是由於 stream processor 進行各種運算都有 latency，更不用提記憶體存取的 latency，因此 CUDA 在執行程式的時候，是以 warp 為單位。目前的 CUDA 裝置，一個 warp 裡面有 32 個 threads，分成兩組 16 threads 的 half-warp。由於 stream processor 的運算至少有 4 cycles 的 latency，因此對一個 4D 的 stream processors 來說，一次至少執行 16 個 threads（即 half-warp）才能有效隱藏各種運算的 latency。 由於 multiprocessor 中並沒有太多別的記憶體，因此每個 thread 的狀態都是直接保存在 multiprocessor 的暫存器中。所以，如果一個 multiprocessor 同時有愈多的 thread 要執行，就會需要愈多的暫存器空間。例如，假設一個 block 裡面有 256 個 threads，每個 thread 用到 20 個暫存器，那麼總共就需要 256x20 = 5,120 個暫存器才能保存每個 thread 的狀態。目前 CUDA 裝置中每個 multiprocessor 有 8,192 個暫存器，因此，如果每個 thread 使用到 16 個暫存器，那就表示一個 multiprocessor 同時最多只能維持 512 個 thread 的執行。如果同時進行的 thread 數目超過這個數字，那麼就會需要把一部份的資料儲存在顯示記憶體中，就會降低執行的效率了。 Shared memory 目前 CUDA 裝置中，每個 multiprocessor 有 16KB 的 shared memory。Shared memory 分成 16 個 bank。如果同時每個 thread 是存取不同的 bank，就不會產生任何問題，存取 shared memory 的速度和存取暫存器相同。不過，如果同時有兩個（或更多個） threads 存取同一個 bank 的資料，就會發生 bank conflict，這些 threads 就必須照順序去存取，而無法同時存取 shared memory 了。 Shared memory 是以 4 bytes 為單位分成 banks。因此，假設以下的資料：     __shared__ int data[128]; 那麼，data[0] 是 bank 0、data[1] 是 bank 1、data[2] 是 bank 2、…、data[15] 是 bank 15，而 data[16] 又回到 bank 0。由於 warp 在執行時是以 half-warp 的方式執行，因此分屬於不同的 half warp 的 threads，不會造成 bank conflict。 因此，如果程式在存取 shared memory 的時候，使用以下的方式：     int number = data[base + tid]; 那就不會有任何 bank conflict，可以達到最高的效率。但是，如果是以下的方式：     int number = data[base + 4 * tid]; 那麼，thread 0 和 thread 4 就會存取到同一個 bank，thread 1 和 thread 5 也是同樣，這樣就會造成 bank conflict。在這個例子中，一個 half warp 的 16 個 threads 會有四個 threads 存取同一個 bank，因此存取 share memory 的速度會變成原來的 1/4。 一個重要的例外是，當多個 thread 存取到同一個 shared memory 的位址時，shared memory 可以將這個位址的 32 bits 資料「廣播」到所有讀取的 threads，因此不會造成 bank conflict。例如：     int number = data[3]; 這樣不會造成 bank conflict，因為所有的 thread 都讀取同一個位址的資料。 很多時候 shared memory 的 bank conflict 可以透過修改資料存放的方式來解決。例如，以下的程式：     data[tid] = global_data[tid];     ...     int number = data[16 * tid]; 會造成嚴重的 bank conflict，為了避免這個問題，可以把資料的排列方式稍加修改，把存取方式改成：     int row = tid / 16;     int column = tid % 16;     data[row * 17 + column] = global_data[tid];     ...     int number = data[17 * tid]; 這樣就不會造成 bank conflict 了。 Global memory 由於 multiprocessor 並沒有對 global memory 做 cache（如果每個 multiprocessor 都有自己的 global memory cache，將會需要 cache coherence protocol，會大幅增加 cache 的複雜度），所以 global memory 存取的 latency 非常的長。除此之外，前面的文章中也提到過 global memory 的存取，要儘可能的連續。這是因為 DRAM 存取的特性所造成的結果。 更精確的說，global memory 的存取，需要是 "coalesced"。所謂的 coalesced，是表示除了連續之外，而且它開始的位址，必須是每個 thread 所存取的大小的 16 倍。例如，如果每個 thread 都讀取 32 bits 的資料，那麼第一個 thread 讀取的位址，必須是 16*4 = 64 bytes 的倍數。 如果有一部份的 thread 沒有讀取記憶體，並不會影響到其它的 thread 速行 coalesced 的存取。例如：     if(tid != 3) {         int number = data[tid];     } 雖然 thread 3 並沒有讀取資料，但是由於其它的 thread 仍符合 coalesced 的條件（假設 data 的位址是 64 bytes 的倍數），這樣的記憶體讀取仍會符合 coalesced 的條件。 在目前的 CUDA 1.1 裝置中，每個 thread 一次讀取的記憶體資料量，可以是 32 bits、64 bits、或 128 bits。不過，32 bits 的效率是最好的。64 bits 的效率會稍差，而一次讀取 128 bits 的效率則比一次讀取 32 bits 要顯著來得低（但仍比 non-coalesced 的存取要好）。 如果每個 thread 一次存取的資料並不是 32 bits、64 bits、或 128 bits，那就無法符合 coalesced 的條件。例如，以下的程式：     struct vec3d { float x, y, z; };     ...     __global__ void func(struct vec3d* data, float* output)     {         output[tid] = data[tid].x * data[tid].x +             data[tid].y * data[tid].y +             data[tid].z * data[tid].z;     } 並不是 coalesced 的讀取，因為 vec3d 的大小是 12 bytes，而非 4 bytes、8 bytes、或 16 bytes。要解決這個問題，可以使用 __align(n)__ 的指示，例如：     struct __align__(16) vec3d { float x, y, z; }; 這會讓 compiler 在 vec3d 後面加上一個空的 4 bytes，以補齊 16 bytes。另一個方法，是把資料結構轉換成三個連續的陣列，例如：     __global__ void func(float* x, float* y, float* z, float* output)     {         output[tid] = x[tid] * x[tid] + y[tid] * y[tid] +             z[tid] * z[tid];     } 如果因為其它原因使資料結構無法這樣調整，也可以考慮利用 shared memory 在 GPU 上做結構的調整。例如：     __global__ void func(struct vec3d* data, float* output)     {         __shared__ float temp[THREAD_NUM * 3];         const float* fdata = (float*) data;         temp[tid] = fdata[tid];         temp[tid + THREAD_NUM] = fdata[tid + THREAD_NUM];         temp[tid + THREAD_NUM*2] = fdata[tid + THREAD_NUM*2];         __syncthreads();         output[tid] = temp[tid*3] * temp[tid*3] +             temp[tid*3+1] * temp[tid*3+1] +             temp[tid*3+2] * temp[tid*3+2];     } 在上面的例子中，我們先用連續的方式，把資料從 global memory 讀到 shared memory。由於 shared memory 不需要擔心存取順序（但要注意 bank conflict 問題，參照前一節），所以可以避開 non-coalesced 讀取的問題。 Texture CUDA 支援 texture。在 CUDA 的 kernel 程式中，可以利用顯示晶片的 texture 單元，讀取 texture 的資料。使用 texture 和 global memory 最大的差別在於 texture 只能讀取，不能寫入，而且顯示晶片上有一定大小的 texture cache。因此，讀取 texture 的時候，不需要符合 coalesced 的規則，也可以達到不錯的效率。此外，讀取 texture 時，也可以利用顯示晶片中的 texture filtering 功能（例如 bilinear filtering），也可以快速轉換資料型態，例如可以直接將 32 bits RGBA 的資料轉換成四個 32 bits 浮點數。 顯示晶片上的 texture cache 是針對一般繪圖應用所設計，因此它仍最適合有區塊性質的存取動作，而非隨機的存取。因此，同一個 warp 中的各個 thread 最好是讀取位址相近的資料，才能達到最高的效率。 對於已經能符合 coalesced 規則的資料，使用 global memory 通常會比使用 texture 要來得快。 運算單元 Stream processor 裡的運算單元，基本上是一個浮點數的 fused multiply-add 單元，也就是說它可以進行一次乘法和一次加法，如下所示：     a = b * c + d; compiler 會自動把適當的加法和乘法運算，結合成一個 fmad 指令。 除了浮點數的加法及乘法之外，整數的加法、位元運算、比較、取最小值、取最大值、及以型態的轉換（浮點數轉整數或整數轉浮點數）都是可以全速進行的。整數的乘法則無法全速進行，但 24 bits 的乘法則可以。在 CUDA 中可以利用內建的 __mul24 和 __umul24 函式來進行 24 bits 的整數乘法。 浮點數的除法是利用先取倒數，再相乘的方式計算，因此精確度並不能達到 IEEE 754 的規範（最大誤差為 2 ulp）。內建的 __fdividef(x,y) 提供更快速的除法，和一般的除法有相同的精確度，但是在 2216 < y < 2218 時會得到錯誤的結果。 此外 CUDA 還提供了一些精確度較低的內建函式，包括 __expf、__logf、__sinf、__cosf、__powf 等等。這些函式的速度較快，但精確度不如標準的函式。詳細的資料可以參考 CUDA Programming Guide 1.1 的 Appendix B。 和主記憶體間的資料傳輸 在 CUDA 中，GPU 不能直接存取主記憶體，只能存取顯示卡上的顯示記憶體。因此，會需要將資料從主記憶體先複製到顯示記憶體中，進行運算後，再將結果從顯示記憶體中複製到主記憶體中。這些複製的動作會限於 PCI Express 的速度。使用 PCI Express x16 時，PCI Express 1.0 可以提供雙向各 4GB/s 的頻寬，而 PCI Express 2.0 則可提供 8GB/s 的頻寬。當然這都是理論值。 從一般的記憶體複製資料到顯示記憶體的時候，由於一般的記憶體可能隨時會被作業系統搬動，因此 CUDA 會先將資料複製到一塊內部的記憶體中，才能利用 DMA 將資料複製到顯示記憶體中。如果想要避免這個重複的複製動作，可以使用 cudaMallocHost 函式，在主記憶體中取得一塊 page locked 的記憶體。不過，如果要求太大量的 page locked 的記憶體，將會影響到作業系統對記憶體的管理，可能會減低系統的效率。

 

BitCoin 龐式交易 老鼠會

有趣的東西…

先爬幾篇文吧!!!!

 

http://www.koalastothemax.com/

https://en.bitcoin.it/wiki/Mining_hardware_comparison

 

有點算是半個老鼠會，但是自投羅網的

外國叫龐式交易

中華電信MOD 60吋電視機方案大手筆降價5,000元 全新推出鴻海製造InFocus40吋液晶顯示器，專案價只要8,800元！

中華電信MOD 60吋電視機方案大手筆降價5,000元 
全新推出鴻海製造InFocus40吋液晶顯示器，專案價只要8,800元中華電信「MOD 60吋大電視」方案自推出以來，揭開高畫質大螢幕電視收視風潮，廣受消費者好評，申裝量已超過萬件，為感謝客戶支持，並讓全台客戶輕鬆享受大尺寸電視與精采高畫質內容的極致視覺體驗，於明日(12)起，【MOD 60吋大電視】優惠方案調降5,000元，凡申辦本方案之客戶，可享60吋LED液晶電視33,800元特價優惠；於102年3月31日前申辦本方案之客戶亦可享5,000元購物抵用券。此外，中華電信MOD將推出最新優惠方案，鴻海製造InFocus40吋液晶顯示器搭配「光世代 MOD 豪華餐」只要8,800元！本方案於4月22日開放客戶申辦。

 

感恩回饋降價5,000元，60吋電視33,800元輕鬆帶回家！

MOD引領大螢幕電視的收視風潮，【MOD家庭豪華餐 60吋HD電視】方案自推出以來，體驗過的客戶有口皆碑，一致表示透過高畫質60吋大電視觀賞MOD節目，感受到前所未有的鮮豔清晰，無法再適應傳統有線電視的畫質，直呼「回不去了」。

中華電信北區分公司總經理涂元光表示：「為了讓更多客戶體驗極致逼真的視覺震撼，【MOD家庭豪華餐 60吋HD電視】方案大方降價5,000元，申請本方案，鴻海製作60吋大電視優惠價只要33,800元！」，為保障消費者權益，102年4月1日至4月11日以原價申辦本方案並付款之客戶，可享5,000元差價退款優惠(註一)。

 

全國銷售創舉，以客為尊，已申辦老客戶尊榮享有5,000元購物抵用券！

MOD在意老客戶，為感謝客戶支持，凡於101年11月9日至102年3月31日前申辦【MOD家庭豪華餐 60吋HD電視】方案之客戶，即可享5,000元SIO產品購物抵用券(註二)，可用於未來推出之「MOD 鴻海製造系列產品」優惠方案使用。

 

全新鴻海製造40吋液晶顯示器，專案價只要8,800元，小坪數空間、第二台電視首選！

家中擁有二台電視已經逐漸形成趨勢，為滿足民眾的使用需求，全新推出的鴻海製造InFocus40吋顯示器，搭配「光世代 MOD 豪華餐 359 元」方案(註三)，超優惠價只要8,800元！鴻海製造InFocus40吋顯示器，採用世界最先進的SIO超視?第十代液晶面板，時尚鋼琴鏡面窄邊框，完美呈現1080P Full HD非凡出色的超精細畫面，最適合小坪數客廳、房間或第二台電視，搭配MOD家庭豪華餐提供近120個頻道，包含高達55個HD高畫質頻道，以及MOD週週免費強檔電影，是市面上最經濟實惠的聰明選擇，本方案於4月22日開放客戶申辦，想要以超值價格享受非凡視覺饗宴的民眾，機不可失就趁現在！

**************************************************************************

註一：102年4月1日至4月11日以原價申請本方案並付款之客戶，可至中華電信原申辦服務中心辦理5,000元退款，詳情請撥免付費客服專線：0800-080-123

註二：SIO產品購物抵用券將由SIO超視?寄出，詳情請撥免付費客服專線：0809-090-510。

註三：本方案需簽約 2 年，可提供 28 家銀行 12 期刷卡零利率。

2013年台北春季電腦展 中華電信多款熱銷手機平板展場現貨供應

 

2013年台北春季電腦展  中華電信多款熱銷手機平板展場現貨供應

首次下殺！3G無限上網搭購平板電腦月租620元！再享熱門平板超值優惠購機價0元起

智慧型手機搭配指定方案加碼贈500元國內通話金  搭來電答鈴超值包手機現折千元

 

2013年春季最受注目的3C高峰會—台北春季電腦展，即將於4月17日至4月21日在世貿一館熱烈展開。呼應今年「超乎想像的科技浪潮」的展覽主軸，中華電信將帶給消費者意想不到的驚喜，申購平板電腦選搭3G 850上網型方案，中華電信更提供前所未見的加碼優惠，月租首次下殺至620元，各種多元方案搭配展場限定優惠，定能讓消費者享受CP值最高的購機體驗！此外，中華電信攤位不僅備有HTC Butterfly、Samsung GALAXY Note II粉色、SONY Xperia Z等多款熱門智慧型手機現貨，於展場攤位申辦智慧型手機搭配指定方案，加碼再回饋500元國內通話金！申辦一般3G手機購機方案搭配「來電答鈴超值包」，不僅超值包月租享8折優惠，購機價現折1,000元！

 

高中低階智慧型手機一次到位　展場現貨熱賣搭配指定方案再享500元國內通話金及原廠好禮

　　中華電信於春電展不僅展出最夯的HTC Butterfly、Samsung GALAXY Note II、iPhone 5、Sony Xperia Z等多款高階熱銷行動裝置，上述熱門手機在中華電信攤位每日皆有現貨供應^(註1)。此外，為滿足行動上網嘗鮮族，現場將展售最新的HTC Desire P鳳蝶機、LG Optimus Duet美夢機以及Samsung GALAXY SIII mini (La Fleur花漾白)、Nokia Lumia 620、Sharp SH530U等多款中華獨家熱賣機種，搭配新好康方案，月租不到千元，輕鬆加入智慧行動一族現在正是時候！凡於展場申辦新超值一年方案、新好康方案、學生方案且為583型(含)以上資費，除可享原購機價優惠，展場加碼再送國內通信費500元，指定機型再加贈皮套、行動電源或16G記憶卡等原廠好禮！而中華電信日前推出的「iPhone 4S新經濟專案」，主打超低月租949元，即可以優惠購機價6,900元輕鬆購得iPhone 4S (16G)， iPhone 5 (32GB) 亦同步推出限時優惠，選搭超值、輕鬆方案，不僅購機價下殺1,000元，再加贈LANVIN時尚背蓋，想一次購得「果迷套餐」絕不能錯過此次春電展！

　　

3G行動上網 850型月租首次下殺620元　名牌平板0元起帶回家

今年為「行動上網普及年」，中華電信為推廣行動上網普及，首度針對3G 850上網型方案搭配平板電腦提供加碼優惠，申辦指定方案，平板無限上網下殺每月只要620元，且Acer Iconia B1-A71平板0元起；最新獨賣Asus Fonepad及熱銷長紅的Samsung Galaxy Tab2 7.0 搭配無限上網也只要990元起；不僅如此，熱門iPad mini (16G) 及Samsung Galaxy Note8均享限時優惠價6,990元起。此外，無論想選購筆記型電腦、網卡或單純申辦上網門號，亦可享有3G 850上網型方案月租8折優惠，想用最超值的花費遨遊智慧行動生活，盡在中華電信。

 

現場申辦「來電答鈴超值包」搭購機現折千元　下載「Hami Music 」App 再送音樂會門票

　　好康不只如此，凡展期期間於中華電信攤位申辦指定購機方案搭配「來電答鈴超值包」，不僅購機價立即折抵1,000元，「來電答鈴超值包」更可享8折優惠，月租費僅64元，申裝答鈴超值包年約服務選擇綜合、西洋或情調音樂盒，再加贈超可愛的小熊捲線器！此外，為讓消費者享受最精緻便利的智慧行動生活，中華電持續深耕各項加值服務，推出 「免費下載Hami Music App再抽炎亞綸音樂會門票」的好康優惠，想要免費暢聽炎亞綸精選好歌？立刻下載Hami Music App，全曲聆聽不中斷，更有炎亞綸?睡不著音樂會?門票準備送給你！更多好康歡迎親洽中華電信攤位。

 

 

※註1：部份機型因原廠供貨有限，例如Samsung GALAXY Note II粉色每日限量供應20台。購機方案實際之適用機型及方案購機價、詳細購機預先繳納金額等包裝內容，請以展場活動現場公告為準。