cuda共享記憶體
阿新 • • 發佈:2020-12-24
共享記憶體
一、共享記憶體的概念
- 共享記憶體實際上是可受使用者控制的一級快取。申請共享記憶體後,其內容在每一個用到的
block
被複制一遍,使得在每個block
內,每一個thread
都可以訪問和操作這塊記憶體,而無法訪問其他block
內的共享記憶體。這種機制就使得一個block
之內的所有執行緒可以互相交流和合作。 - 在裝置程式碼中宣告共享記憶體要使用
__shared__
變數宣告說明符。 - 共享記憶體有兩種方法:靜態與動態
//動態分配共享記憶體
dynamicReverse<<<1, n, n*sizeof(int)>>>(d_d, n);
二、程式碼
程式碼如下(示例):
#include <stdio.h>
__global__ void staticReverse(int *d, int n)
{
__shared__ int s[64];
int t = threadIdx.x;
int tr = n-t-1;
s[t] = d[t];
__syncthreads();
d[t] = s[tr];
}
__global__ void dynamicReverse(int *d, int n)
{
extern __shared__ int s[];
int t = threadIdx.x;
int tr = n-t-1;
s[t] = d[t];
__syncthreads();
d[t] = s[tr];
}
int main(void)
{
const int n = 64;
int a[n], r[n], d[n];
for (int i = 0; i < n; i++) {
a[i] = i;
r[i] = n-i-1;
d[i] = 0;
}
int *d_d;
cudaMalloc(&d_d, n * sizeof(int));
// run version with static shared memory
cudaMemcpy(d_d, a, n*sizeof(int), cudaMemcpyHostToDevice);
staticReverse<<<1,n>>>(d_d, n);
cudaMemcpy(d, d_d, n*sizeof(int), cudaMemcpyDeviceToHost);
for (int i = 0; i < n; i++)
if (d[i] != r[i]) {
printf("Error: d[%d]!=r[%d] (%d, %d)\n", i, i, d[i], r[i]);
exit(-1);
}
printf("static success\n");
// run dynamic shared memory version
cudaMemcpy(d_d, a, n*sizeof(int), cudaMemcpyHostToDevice);
dynamicReverse<<<1,n,n*sizeof(int)>>>(d_d, n);
cudaMemcpy(d, d_d, n * sizeof(int), cudaMemcpyDeviceToHost);
for (int i = 0; i < n; i++) {
if (d[i] != r[i]) {
printf("Error: d[%d]!=r[%d] (%d, %d)\n", i, i, d[i], r[i]);
exit(-1);
}
}
printf("dynamic success\n");
}
執行結果
static success
dynamic success