CUDA共享記憶體的使用示例

cuda共享記憶體使用示例如下：參考教材《gpu高效能程式設計cuda實戰》。p54-p65

1 #include 2 #include 3 #include 4 #include 5 #include 6 #include 78

using
namespace
std;910
#define imin(a,b) (a11
const
int n = 33 * 1024;12
const
int threadsperblock = 256;13
const
int blockspergrid = imin(32, (n + threadsperblock - 1) /threadsperblock);
1415 __global__  void dot(float *a, float *b, float *c)
1628
29//
每個執行緒塊中線程計算的加和儲存到緩衝區cache，一共有blockspergrid個緩衝區副本
30     cache[cacheindex] =temp;
31//
對執行緒塊中的執行緒進行同步
32__syncthreads();
3334
//歸約運算，將每個緩衝區中的值加和，存放到緩衝區第乙個元素位置
35int i = blockdim.x / 2;36
while (i != 0)37
42__syncthreads();
43         i /= 2;44
}45//使用第乙個執行緒取出每個緩衝區第乙個元素賦值到c陣列
46if (cacheindex == 0)47
50}5152
void
main()
5373
74//
將陣列上傳到gpu
75     cudamemcpy(dev_a, a, n * sizeof(float
), cudamemcpyhosttodevice);
76     cudamemcpy(dev_b, b, n * sizeof(float
), cudamemcpyhosttodevice);
7778     dot << > >(dev_a, dev_b, dev_partial_c);
7980     cudamemcpy(partial_c, dev_partial_c, blockspergrid * sizeof(float
), cudamemcpydevicetohost);
8182
//cpu 完成最終求和
83     c = 0;84
for (int i = 0; i < blockspergrid; i++)
8588
89#define sum_squares(x) (x*(x+1)*(2*x+1)/6)
90     printf("
does gpu value %.6g = %.6g?\n
", c, 2 * sum_squares((float)(n - 1
)));
9192
cudafree(dev_a);
93cudafree(dev_b);
94cudafree(dev_partial_c);
9596
free
(a);
97free
(b);
98free
(partial_c);
99 }

CUDA 共享記憶體

於gpu上啟動的每個執行緒塊上的共享記憶體，cuda c編譯器都會建立該變數的乙個副本。同一執行緒塊的每個執行緒都共享這塊記憶體，但是執行緒無法看到也不能修改其他執行緒塊中的共享記憶體。這樣做可以使得乙個執行緒塊中的多個執行緒能夠在計算上通訊和協作。共享記憶體緩衝區駐留在物理gup上，因此訪問共享記...

共享CUDA記憶體

共享cuda記憶體 程序間共享 此功能僅限於linux。將裝置陣列匯出到另乙個程序 使用cuda ipc api，可以與同一臺計算機上的另乙個程序共享裝置陣列。為此，請使用.get ipc handle 裝置陣列上的方法獲取乙個ipcarrayhandle物件，該物件可以轉移到另乙個程序。devic...

cuda 共享記憶體bank conflict詳解

在cuda平行計算中，共享記憶體在gpu速度優化上扮演著重要作用，但是如果共享記憶體使用不當，也會導致速度不快反降或者提速效果不佳，如發生bank conflict bank的中文翻譯為儲存體，gpu 共享記憶體是基於儲存體切換的架構 bank switched architecture 一般現在的...