我有两个 CUDA 内核可以计算类似的东西。一种是使用全局内存(myfun
是一个设备函数,它从全局内存中读取大量数据并进行计算)。第二个内核将该数据块从全局内存传输到共享内存,以便数据可以在块的不同线程之间共享。我使用全局内存的内核比使用共享内存的内核快得多。可能的原因有哪些?
loadArray
只复制一小部分d_x
to 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;
}
}