我有兩個計算類似內容的CUDA內核。一種是使用全局內存(myfun
是一種從全局內存中讀取很多內容並進行計算的設備函數)。第二個內核將該塊數據從全局內存傳輸到共享內存,以便數據可以在塊的不同線程之間共享。使用全局內存的內核比共享內存的內核快得多。可能的原因是什麼?CUDA中的全局vs共享內存
loadArray只是將d_x
的一小部分複製到m
。
__global__ void mykernel(float *d_x, float *d_y, int *d_z, float *d_u, int N, int K, int D)
{
int tid = blockIdx.x*blockDim.x + threadIdx.x;
int index = 0;
float max_s = 1e+37F;
if (tid < N)
{
for (int i = 0; i < K; i++)
{
float s = myfun(&d_x[i*D], d_y, tid);
if (s > max_s)
{
max_s = s;
index = i;
}
}
d_z[tid] = index;
d_u[tid] = max_s;
}
}
使用共享內存:
__global__ void mykernel(float *d_x, float *d_y, int *d_z, float *d_u, int N, int K)
{
int tid = blockIdx.x*blockDim.x + threadIdx.x;
int index = 0;
float max_s = 1e+37F;
extern __shared__ float m[];
if(threadIdx.x == 0)
loadArray(m, d_x);
__syncthreads();
if (tid < N)
{
for (int i = 0; i < K; i++)
{
float s = myfun(m, d_y, tid);
if (s > max_s)
{
max_s = s;
index = i;
}
}
d_z[tid] = index;
d_u[tid] = max_s;
}
}