cuda中的共享内存

我有一个大小为NxN的复杂数据的矩阵u，我想每行乘以一个大小为1xN的向量k。 u中的数据按行存储。cuda中的共享内存

我有两个这样的实现。一个利用共享内存，将矩阵分成瓦片，另一个没有。

我发现共享内存实现multiply1不会更快，并且系统性地一样快，甚至比multiply2慢。

共享存储器实现如下，

__global__ void multiply1(cufftComplex *u, cufftComplex *k) { 
    __shared__ cufftComplex k_s[BLOCK_WIDTH]; 
    int idx = blockDim.x*blockIdx.x + threadIdx.x; 
    int idy = blockDim.y*blockIdx.y + threadIdx.y; 
    int index; 

    if (threadIdx.y == 0 && idx < N) { 
     k_s[threadIdx.x] = k[idx]; 
    } 
    __syncthreads(); 

    if (idx < N && idy < N) { 
     index = N*idy + idx; 
     u[index] = cuCmulf(k_s[threadIdx.x],u[index]); 
    } 

}

鉴于全球存储器实现如下，

__global__ void multiply2(cufftComplex *u, cufftComplex *k) { 
     int idx = blockDim.x * blockIdx.x + threadIdx.x; 

     if (idx < N*N) { 
      u[idx] =cuCmulf(k[idx % N],u[idx]); 
     } 
    }

和主函数调用，对于大小的矩阵64×64

dim3 block(16,16); 
dim3 grid(4,4); 
multiply1<<<grid, block>>>(d_u, d_k); 
multiply2<<<16, 256>>>(d_u, d_k);

如何使用探查器找出为什么multiply1不是gett速度至少略有增加？哪些指标可以阐明究竟发生了什么？

分析器告诉我multiply1，我得到152 GB/s的全局内存加载吞吐量，而multiply2我得到81 GB /秒。这是合乎逻辑的，因为我从全局内存中加载较少。这不应该转化为更快的执行吗？

来源

2014-04-03 user3495341

如果你多次使用它会更快，但在这里你只用了一次。您在变换你的问题：

copy from global memory to shared memory 
read from shared memory

代替：

read from global memory

所以是的，它肯定是比以前的算法只使用全局内存要慢。如果你想利用共享内存，你的算法要多次读取它，否则你不会花费全局内存。

来源

2014-04-03 20:50:44 user3018144

我从全局存储器复制每块一次k的适当的元件，而不是每一次元素，因为if语句的：如果（threadIdx.y == 0 && IDX user3495341

每块一次，但同一块中的线程正在等待复制完成。 –

我该如何检查？会有一种方法可以在探查器中执行此操作吗？ – user3495341

cuda中的共享内存

回答

相关问题