內存聲明：
設備和常量內存：
使用__device__或__constant__內存空間說明符在文件范圍內聲明的內存在使用設備運行時時具有相同的行為。 所有內核都可以讀取或寫入設備變量，無論內核是由主機還是設備運行時初始啟動。 等同地，所有內核將具有與在模塊范圍內聲明的__constant__相同的視圖。
紋理和表面內存：
CUDA支持動態創建的紋理和表面對象1，其中可以在主機上創建紋理參考，傳遞給內核，由內核使用，然后從主機中銷毀。 設備運行時不允許從設備代碼中創建或銷毀紋理或表面對象，但可以在設備上自由使用并自由傳送由主機創建的紋理和表面對象。 無論它們在哪里創建，動態創建的紋理對象總是有效的，并且可以從父級傳遞給子級內核。
設備運行時不支持從設備啟動的內核中的傳統模塊范圍（即費米架構）紋理和表面。 模塊范圍（遺留）紋理可以從主機創建并在設備代碼中用于任何內核，但只能由頂級內核（即從主機啟動的）使用。

1：動態創建的紋理和曲面對象是CUDA 5.0引入的CUDA內存模型的補充。 有關詳細信息，請參閱CUDA編程指南。
共享內存變量聲明：
在CUDA中，C / C ++共享內存可以聲明為靜態大小的文件范圍變量或函數范圍變量，也可以聲明為外部變量，其大小由內核的調用者在運行時通過啟動配置參數確定。 這兩種類型的聲明在設備運行時間下均有效。

__global__ void permute(int n, int *data) {extern __shared__ int smem[];if (n <= 1)return;smem[threadIdx.x] = data[threadIdx.x];__syncthreads();permute_data(smem, n);__syncthreads();// Write back to GMEM since we can't pass SMEM to children.data[threadIdx.x] = smem[threadIdx.x];__syncthreads();if (threadIdx.x == 0) {permute << < 1, 256, n / 2 * sizeof(int) >> >(n / 2, data);permute << < 1, 256, n / 2 * sizeof(int) >> >(n / 2, data + n / 2);} } void host_launch(int *data) {permute << < 1, 256, 256 * sizeof(int) >> >(256, data); }

符號地址：
由于所有全局范圍的設備變量都在內核的可見地址空間中，所以設備端符號（即標記為__ device_的那些符號）可以通過簡單的＆運算符從內核中引用。 這也適用于__constant__符號，但在這種情況下，指針將引用只讀數據。
考慮到設備端符號可以直接引用，引用符號的CUDA運行時API（例如cudaMemcpyToSymbol（）或cudaGetSymbolAddress（））是多余的，因此設備運行時不支持。 注意這意味著即使在子內核啟動之前，也不能在運行的內核中更改常量數據，因為對__constant__空間的引用是只讀的。

總結

以上是生活随笔為你收集整理的CUDA学习（六十三）的全部內容，希望文章能夠幫你解決所遇到的問題。

如果覺得生活随笔網站內容還不錯，歡迎將生活随笔推薦給好友。