cuda合並訪問

⑴ 什麼是CUDA

CUDA(Compute Unified Device Architecture)，顯卡廠商NVidia推出的運算平台。 CUDA™是一種由NVIDIA推出的通用並行計算架構，該架構使GPU能夠解決復雜的計算問題。 它包含了CUDA指令集架構(ISA)以及GPU內部的並行計算引擎。
隨著顯卡的發展，GPU越來越強大，而且GPU為顯示圖像做了優化。在計算上已經超越了通用的CPU。如此強大的晶元如果只是作為顯卡就太浪費了，因此NVidia推出CUDA，讓顯卡可以用於圖像計算以外的目的。
簡單來講，比如通過CUDA架構，視頻播放軟體可以充分挖掘NVIDIA系列顯卡的GPU並行計算能力，輕松進行高清影片的播放，與軟體高清解碼相比，CPU佔用可以下降一半以上。當然，CUDA的應用領域絕不僅僅是視頻、圖形、游戲，包括各種3D和建模，醫療、能源、科學研究等，到處都可見到這種技術架構的應用。
支持CUDA的硬體環境需要有NVidia GF8系列及以上型號的顯卡，並且安裝185版本以上的顯卡驅動程序。以QQ影音播放器來講，要想開啟CUDA硬體解碼加速，可以打開QQ影音的「播放器設置」，進入「高清加速」面板，在「硬體優化」中選擇「自定義優化模式」，然後在「濾鏡配置」中的「視頻解碼器」中自定義選擇相應的「QQ CUDA Video Decoder(CUDADecFilter.ax)」即可。而關閉CUDA加速，只需取消選擇「QQ CUDA Video Decoder(CUDADecFilter.ax)」，或者切換到「智能高清模式」或「穩定兼容模式」通過這種高清解碼定義的開啟，並不是說你的畫質能夠提升多少，而是提升高清視頻播放時的流暢以及降低CPU的佔用。這個時候，節約下來的CPU空間，可以允許你再去做別的工作，這樣就會大大提升你的工作效率，而不至於除了看視頻，其他的什麼都不能做了。

⑵ CUDA——內存層級

通過cache L2(CC>=3.0)訪問，cache line 大小128 bytes ，每個線程操作盡量少的cache line，速度更快

共享內存分成相同大小的內存塊，實現高速並行訪問，但是當多個線程的請求地址映射到同一個內存塊block時，訪問是串列的

步幅stride為n時 最大公約數為1，即gcd(n,32)==1 ,訪問共享內存可以避免塊沖突

位於堆棧中，不在寄存器中的所有內容
作用域為特定線程
存儲在global內存空間中，速度比寄存器慢很多

內核使用的寄存器比可用的寄存器多,存儲到local memory中

類似constant memory，是只讀內存，以某種形式訪問的時候可以提升性能。原本是用在OpenGL和DirectX渲染管線中的。
有用的特點：

用例：

CC ≥ 3.5 大多數的 __restrict__ 變數自動載入到紋理緩存中了
通過 __ldg函數強行載入到緩存

⑶ cuda 內核函數中的多線程訪問變數是如何處理

你好，

樓主是說block級的同步嗎？CUDA沒有提供這樣的函數，因為一是這種操作太耗時間，二是凡是這種操作都可以以其他方式解決。CUDA提供了一下兩個同步函數：

• __syncthreads(); 用來同步每個block內的線程，用於kernel當中。

• cudaDeviceSyncronize(); 用來同步設備上所有之前的操作。在調用cudaStream或者多個GPU的時候需要用到這個函數來保證每個stream或每個GPU運行到同一點。

⑷ HartSift: 一種基於GPU的高准確性和實時SIFT

尺度不變特徵變換 (SIFT) 是最流行和最強大的特徵提取演算法之一，因為它對尺度、旋轉和光照保持不變。它已被廣泛應用於視頻跟蹤、圖像拼接、同時定位和映射（SLAM）、運動結構（SFM）等領域。然而，高計算復雜度限制了其在實時系統中的進一步應用。這些系統必須在准確性和性能之間進行權衡以實現實時特徵提取。他們採用其他更快但精度較低的演算法，如 SURF 和 PCA-SIFT。為了解決這個問題，本文提出了一種使用 CUDA 的 GPU 加速 SIFT，命名為 HartSift，充分利用單機CPU和GPU的計算資源，實現高精度、實時的特徵提取。實驗表明，在 NIVDIA GTX TITAN Black GPU 上，HartSift 可以根據圖像的大小在 3.14-10.57ms (94.61-318.47fps) 內處理圖像。此外，HartSift 分別比 OpenCV-SIFT（CPU 版本）和 SiftGPU（GPU 版本）快 59.34-75.96 倍和 4.01-6.49 倍。同時，HartSift 的性能和 CudaSIFT（迄今為止最快的 GPU 版本）的性能幾乎相同，而 HartSift 的准確度遠高於 CudaSIFT。

SIFT演算法可以提取大量顯著特徵，這些特徵在縮放、旋轉、光照和3D視點保持不變，還提供了跨越雜訊和仿射失真的穩健匹配。但SIFT的高計算復雜度限制了其在大規模數據和實時系統中的進一步應用。而復雜度較低的演算法，如SURF、PCA-SIFT的准確性又不太高。因此，在主流計算平台上實現高精度、實時的SIFT是一個重要而有意義的研究課題。

而SIFT演算法具有很好的並行性，可以正確移植到GPU上。因此，在配備GPU的異構計算系統上實現高性能的SIFT具有重要的實用價值。

SIFT 演算法包含三個階段，包括高斯差分（DoG）金字塔的構建、精確的關鍵點定位和 128 維描述符生成。由於每個階段都有自己的並行特性，因此必須使用不同的並行粒度和優化來實現高性能。尤其是後兩個階段，負載不平衡不利於GPU優化，會導致性能下降。

本文的主要貢獻和創新可以概括如下：

有許多工作嘗試在GPU上使用SIFT演算法。

然而，為了實現高性能，他們省略了 SIFT 演算法的一些重要步驟，例如將輸入圖像加倍、保持尺度變化的連續性和擬合二次函數以定位準確的關鍵點信息。作者的實驗表明，這些遺漏會導致 SIFT 丟失很多關鍵點和准確性。

Lowe將輸入圖像尺寸加倍作為高斯金字塔 的最底層，每個尺度 通過高斯卷積產生：

高斯金字塔確定之後，利用相同Octave的層級相減，得到差分金字塔：

其中 ，在本文中， .

檢測尺度空間極值

將DoG金字塔每個像素與相鄰像素比較，同層8個，上下層9個，若像素是局部最大值或局部最小值，將其視為關鍵點候選。

去除無效關鍵點

去除較低對比度和不穩定邊緣響應的候選關鍵點，通過將3D二次函數擬合到附近數據執行子像素插值，以獲取精確的位置、比例和主曲率比。

方向分配

將候選關鍵點周圍的梯度累積到36 bins的直方圖中，根據每層的尺度計算搜索半徑。每個緊鄰像素由一個高斯加權窗口加權，梯度方向累計到36 bins的方向直方圖中。峰值為主要梯度方向，同時超過峰值80%的局部峰值bin也被視為關鍵點方向。

對關鍵點周圍像素計算梯度直方圖，搜索半徑比上一步驟大得多，同樣用一個高斯加權函數用於為每個鄰居的梯度值分配權重。

根據梯度方向將最終的梯度值累積到一個 360-bin 的圓形方向直方圖。最後，直方圖將被歸一化、平滑並轉換為 128D 描述符。

構建金字塔應該保持順序，以保證尺度空間變化連續性。Acharya和Bjorkman為加快這一過程，犧牲准確性打破構建順序。考慮到不能使准確性降低，構建順序在HartSift中保留。

分離卷積核

對於 大小的卷積核處理 大小的圖像需要進行 次運算，如果將2D卷積核拆解為兩個1D的卷積核，計算量減少至 . 通過使用共享內存和向量化方法，更容易實現合並全局內存訪問並減少一維卷積的冗餘訪問。

Uber 內核

Uber內核將多個不同任務放到一個物理內核中，在一個內核中並行處理任務，而不需要在內核之間切換。差分金字塔第 層由高斯金字塔第 和第 層決定。將高斯差分金字塔和高斯卷積核封裝在單個核中，可以充分挖掘並行性。

線程不需要重復讀取高斯金字塔第 層的值，這是由於第 層的值計算完後，結果會放在寄存器內而不是全局內存中。藉助Uber內核的優勢，我們可以節省 的空間和 的內核運行時間

異構並行

HartSift 採用異構並行方法來加速這一階段。CPU 和 GPU 將並行協作，構建 DoG 金字塔。

由於GPU處理小圖像沒有優勢，作者將 以下的圖像放到CPU處理，大圖像放到GPU處理。用戶也可以自行設置分離點，確保CPU和GPU負載平衡。

存在兩個問題：

負載均衡

Warp是GPU最小並行執行單元，即以鎖步方式執行的 32 個線程的集合。若負載不均衡，則warp執行時間取決於最後一個線程完成的時間，warp負載不均衡會導致GPU效率降低。

由於候選關鍵點分布的隨機性，幾乎所有經線都包含不同數量的空閑線程。如果這些warp繼續處理以下部分，就會出現兩個級別的負載不平衡.

在去除無效的候選關鍵點部分時，線程將進行亞像素插值以獲得准確的候選關鍵點信息，從而去除具有低對比度或不穩定邊緣響應的關鍵點候選。換句話說，一些線程會比其他線程更早返回一次。負載不平衡會變得更加嚴重。

為了突破性能瓶頸，HartSift 引入了重新平衡工作負載和多粒度並行優化。

重新平衡工作負載

當檢測到負載不平衡時，HartSift 將通過啟動具有適當粒度的新內核並分派每個具有 32 個活動線程的新經線來重新平衡工作負載。

此外，啟動三個內核分別處理這三個部分，不僅可以重新平衡工作量，還可以根據不同部分的並行特性提供多粒度的並行。

多粒度並行

重新平衡工作負載優化保證每個內核中的線程和經線被完全載入，多粒度並行優化保證工作負載將平均分配到線程和經線。此外，不同內核的並行粒度取決於工作負載的特性。

HartSift通過將一個線程映射到一個或多個像素，採用與關鍵點候選檢測部分和無效關鍵點去除部分並行的線程粒度。然而，線程粒度並行會導致方向分配部分的負載不平衡，因為不同關鍵點的相鄰區域半徑不同。線程粒度並行會為單個線程分配過多的工作，這在某些情況下限制了硬體資源的利用率。所以在這部分應用混合粒度並行：扭曲粒度構建直方圖，線程粒度找出並將主導方向分配給相應的關鍵點。

基於扭曲的直方圖演算法

作者針對每個關鍵點提出了一種基於扭曲粒度和原子操作的高性能直方圖演算法，以充分利用局部性。

該階段關鍵點的鄰域半徑遠大於前一階段。需要為每個關鍵點累積數千個鄰居到一個 360-bin 直方圖。如果採用前一階段的基於原子扭曲的直方圖演算法，會對這一階段的性能產生不同的影響。

HartSift引入了一種atomic-free的直方圖演算法，進一步提升了這一階段的性能。

該演算法包含三個步驟：

為了消除線程間的負載不平衡，實現全局合並訪問，HartSift 使用一個warp 來處理一個keypoint 的所有鄰居。當線程計算出它們的方向 bin 時，它們需要根據bin變數的值將梯度值累加到局部直方圖。考慮到有如此多的鄰居並且一個經線的一些線程可能具有相同的 bin，演算法1引入了一種無原子的多鍵約簡方法來累積每個經線的部分和。這種方法可以利用warp級shuffle和vote指令完全消除原子操作和本地同步。它根據bin對經紗的線程進行分組並指定每組具有最低車道的線程作為隊長線程。隊長線程將保存他們自己的 bin 的部分總和，並將它們並行地累積到駐留在共享內存中的本地直方圖，而不會發生 bank 沖突和同步。在遍歷所有鄰居後，HartSift 將最終的局部直方圖復制到駐留在全局內存中的全局直方圖。

本文提出了一種GPU上的並行SIFT，命名為Hart-Sift，它可以在單機內同時使用CPU和GPU來實現高精度和實時的特徵提取。HartSift根據每個階段的不同特點，通過適當採用不同的優化策略來提升性能，例如負載均衡、基於warp的直方圖演算法和不同尺度樣本的atomic-free直方圖演算法等。在NVIDIA GTX TITAN Black GPU上，HartSift可以在3.14 ~ 10.57ms（94.61 ~ 318.47fps）內提取高精度特徵，輕松滿足高精度和實時性的苛刻要求。另外，與OpenCV-SIFT和SiftGPU相比，HartSift獲得了59.34 ~ 75.96倍和4.01 ~ 6.49倍加速分別。同時，HartSift 和 CudaSIFT 的性能幾乎相同，但 HartSift 遠比 CudaSIFT 准確。