CUDA將CPU作為主機（Host），GPU作為設備（Device）。一個系統中可以有一個主機和多個設備。CPU負責邏輯性強的事務處理和串行計算，GPU專注於執行高度線程化的並行處理任務。它們擁有相互獨立的存儲器（主機端的內存和顯卡端的顯存）。
運行在GPU上的函數稱為kernel（內核函數）。一個完整的CUDA程序是由一些列的kernel函數和主機端的串行處理步驟共同完成的。CPU串行代碼的工作包括在kernel啟動前進行的數據准備、設備初始化以及在kernel之間進行一些串行化計算。
kernel函數以函數類型限定符_global_定義，並且只能在主機端代碼中調用。調用時需要制定kernel的Grid中的block數目以及每個block中thread的數目。每個線程有自己的blockID和threadID，可用來區分其它的線程，這兩個內建變量是只讀的，由專用寄存器提供，只能有kernel函數調用。
CUDA將計算任務映射為大量的可並行執行的線程，並由硬件動態調度和執行這些線程。kernel是以線程網格Grid的形式組織，每個Grid又若干個線程塊block組成，每個block又由若干線程thread組成。本質上，kernel是以block為單位執行的，grid只是用來表示一些列並行block的集合。Block之間是不能彼此通信的。目前一個kernel只支持一個grid，在多指令多數據（MIMD）構架中會存在多個grid。
為方便編程，CUDA使用了dim3類型的內建變量threadIdx和threadIdx。
對於一維的block，線程的編號是threadIdx.x。
對於二維（Dx，Dy）的block，線程的編號是threadIdx.x+threadIdx.y * Dx。
對於三維（Dx，Dy，Dz）的block，線程的編號是threadIdx.x+threadIdx.y * Dx+hreadIdx.z* Dx * Dy。
GPU的計算核心是流多處理器（SM），每個SM包含8個標量流處理器（SP）以及其它的運算單元。Kernel是以block為單位執行的，同一個block的線程需要共享數據，因此它們共享同一個SM。一個block必須分配到一個SM，但是可以一個SM中同一個時刻有多個活動塊（active block）執行等待，即同一個SM可以有多個block上下文。實際運行中，block會被分為更小的線程束（wrap），線程束的大小由硬件的計算能力決定，Tesla的構架中一個wrap由32個線程組成。
CUDA采用了單指令多線程執行模型。這個模型是對單指令多數據的改進。CUDA中執行寬度可以在1——512個線程之間變化，但是在單指令多數據中執行寬度必須是一個wrap（32）。
每一個線程擁有自己的私有存儲器，每一個線程塊擁有一塊共享存儲器(Shared memory)；最后，grid中所有的線程都可以訪問同一塊全局存儲器(global memory)。除此之外，還有兩種可以被所有線程訪問的只讀存儲器：常數存儲器（constant memory）和紋理存儲器(Texture memory),它們分別為不同的應用進行了優化。全局存儲器、常數存儲器和紋理存儲器中的值在一個內核函數執行完成后將被繼續保持，可以被同一程序中其也內核函數調用。
| 存儲器 |
位置 |
擁有緩存 |
訪問權限 |
變量生存周期 |
| register |
GPU片內 |
N/A |
Device可讀/寫 |
與thread相同 |
| Local memory |
板載顯存 |
無 |
Device可讀/寫 |
與thread相同 |
| Shared memory |
GPU片內 |
N/A |
Device可讀/寫 |
與block相同 |
| Constant memory |
板載顯存 |
有 |
Device可讀，host要讀寫 |
可在程序中保持 |
| Texture memory |
板載顯存 |
有 |
Device可讀，host要讀寫 |
可在程序中保持 |
| Global memory |
板載顯存 |
無 |
Device可讀/寫, host可讀/寫 |
可在程序中保持 |
| Host memory |
Host內存 |
無 |
host可讀/寫 |
可在程序中保持 |
| Pinned memory |
Host內存 |
無 |
host可讀/寫 |
可在程序中保持 |
CUDA存儲器模型：
GPU片內：register,shared memory;
板載顯存：local memory,constant memory,texture memory,global memory;
host 內存： host memory, pinned memory.
register: 訪問延遲極低；
基本單元：register file （32bit/each）
計算能力1.0/1.1版本硬件：8192/SM；
計算能力1.2/1.3版本硬件： 16384/SM;
每個線程占有的register有限，編程時不要為其分配過多私有變量；
local memory：寄存器被使用完畢，數據將被存儲在局部存儲器中；
大型結構體或者數組；
無法確定大小的數組；
線程的輸入和中間變量；
定義線程私有數組的同時進行初始化的數組被分配在寄存器中；
shared memory：訪問速度與寄存器相似；
實現線程間通信的延遲最小；
保存公用的計數器或者block的公用結果；
硬件1.0~1.3中，16KByte/SM,被組織為16個bank；
聲明關鍵字 _shared_ int sdata_static[16];
global memory:存在於顯存中，也稱為線性內存（顯存可以被定義為線性存儲器或者CUDA數組）；
cudaMalloc（）函數分配，cudaFree（）函數釋放，cudaMemcpy（）進行主機端與設備端的數據傳輸；
初始化共享存儲器需要調用cudaMemset（）；
二維三維數組：cudaMallocPitch（）和cudaMalloc3D（）分配線性存儲空間，可以確保分配滿足對齊要求；
cudaMemcpy2D（），cudaMemcpy3D（）與設備端存儲器進行拷貝；
host內存：分為pageable memory 和 pinned memory
pageable memory： 通過操作系統API（malloc（），new（））分配的存儲器空間；
pinned memory：始終存在於物理內存中，不會被分配到低速的虛擬內存中，能夠通過DMA加速與設備端進行通信；cudaHostAlloc(), cudaFreeHost()來分配和釋放pinned memory；
使用pinned memory優點：主機端-設備端的數據傳輸帶寬高；某些設備上可以通過zero-copy功能映射到設備地址空間，從GPU直接訪問，省掉主存與顯存間進行數據拷貝的工作；pinned memory 不可以分配過多：導致操作系統用於分頁的物理內存變，導致系統整體性能下降；通常由哪個cpu線程分配，就只有這個線程才有訪問權限；
cuda2.3版本中，pinned memory功能擴充：
portable memory：讓控制不同GPU的主機端線程操作同一塊portable memory，實現cpu線程間通信；
使用cudaHostAlloc（）分配頁鎖定內存時，加上cudaHostAllocPortable標志；
write-combined Memory：提高從cpu向GPU單向傳輸數據的速度；不使用cpu的L1，L2 cache對一塊pinned memory中的數據進行緩沖，將cache資源留給其他程序使用；在pci-e總線傳輸期間不會被來自cpu的監視打斷；在調用cudaHostAlloc（）時加上cudaHostAllocWriteCombined標志；cpu從這種存儲器上讀取的速度很低；
mapped memory：兩個地址：主機端地址（內存地址），設備端地址（顯存地址）。
可以在kernnel程序中直接訪問mapped memory中的數據，不必在內存和顯存之間進行數據拷貝，即zero-copy功能；在主機端可以由cudaHostAlloc（）函數獲得，在設備端指針可以通過cudaHostGetDevicePointer（）獲得；通過cudaGetDeviceProperties（）函數返回的canMapHostMemory屬性知道設備是否支持mapped memory；在調用cudaHostAlloc（）時加上cudaHostMapped標志，將pinned memory映射到設備地址空間；必須使用同步來保證cpu和GPu對同一塊存儲器操作的順序一致性；顯存中的一部分可以既是portable memory又是mapped memory；在執行CUDA操作前，先調用cudaSetDeviceFlags（）（加cudaDeviceMapHost標志）進行頁鎖定內存映射。
constant memory：只讀地址空間；位於顯存，有緩存加速；64Kb；用於存儲需要頻繁訪問的只讀參數 ；只讀；使用_constant_ 關鍵字，定義在所有函數之外；兩種常數存儲器的使用方法：直接在定義時初始化常數存儲器；定義一個constant數組，然后使用函數進行賦值；
texture memory：只讀；不是一塊專門的存儲器，而是牽涉到顯存、兩級紋理緩存、紋理拾取單元的紋理流水線；數據常以一維、二維或者三維數組的形式存儲在顯存中；緩存加速；可以聲明大小比常數存儲器大得多；適合實現圖像樹立和查找表；對大量數據的隨機訪問或非對齊訪問有良好的加速效果；在kernel中訪問紋理存儲器的操作成為紋理拾取（texture fetching）；紋理拾取使用的坐標與數據在顯存中的位置可以不同，通過紋理參照系約定二者的映射方式；將顯存中的數據與紋理參照系關聯的操作，稱為將數據與紋理綁定（texture binding）；顯存中可以綁定到紋理的數據有：普通線性存儲器和cuda數組；存在緩存機制；可以設定濾波模式，尋址模式等；