共享内存数组大小及bank conflict

  • 2 replies
  • 299 views
*

queqd

  • *
  • 10
    • 查看个人资料
共享内存数组大小及bank conflict
« 于: 二月 28, 2023, 06:04:33 pm »
1:我在核函数中设置线程块的大小为128,我需要将float类型的变量定义为共享内存变量,发现当将共享内存数组定义写成数组长度__shared__ float S_A[128]时程序运行是错误的,当把定义写为__shared__ float S_A[sizeof(float)*320]时可以得到正确的结果。共享内存不是可以直接定义数组长度吗?
2:我在核函数中共享内存使用如下
程序代码: [选择]
kernel2(float* A,float* M, float* Y,float* X)
{
    const int n = blockDim.x * blockIdx.x + threadIdx.x; 
     
        __shared__ float S_A[sizeof(float) * 320];       
           
        for (int j = 0; j < 12; j++)
        {           
                S_A[threadIdx.x * 12 + j] = A[n * 12 + j];

                S_A[threadIdx.x * 12 + j]=M[Y[n]]+X[n * 12 + j];           
        }
__syncthreads();


}

因该怎么修改程序消除或减轻共享内存的bank冲突?

Re: 共享内存数组大小及bank conflict
« 回复 #1 于: 三月 01, 2023, 03:13:13 pm »
您好,您第一个问题是不是数组索引越界了,for循环里的S_A[threadIdx.x * 12 + j],其中threadIdx.x最大是127,当它取最大值(或没有取到最大值时),threadIdx.x * 12 + j都超过了您设置的S_A[128]里面的128

*

queqd

  • *
  • 10
    • 查看个人资料
Re: 共享内存数组大小及bank conflict
« 回复 #2 于: 三月 01, 2023, 05:40:09 pm »
您好,您第一个问题是不是数组索引越界了,for循环里的S_A[threadIdx.x * 12 + j],其中threadIdx.x最大是127,当它取最大值(或没有取到最大值时),threadIdx.x * 12 + j都超过了您设置的S_A[128]里面的128
您好,我按照您说的索引越界问题进行了修改将S_A定义设置为__shared__ float S_A[1536],程序可以正确运行,但是将定义设置为__shared__ float S_A[1280]程序也能正确运行,所以对共享内存数组大小问题还是有疑问。


另外,在我的核函数中对共享内存的使用如下
程序代码: [选择]
kernel2(float* A,float* M, float* Y,float* X)
{
    const int n = blockDim.x * blockIdx.x + threadIdx.x; 
       E[n]=0;

        __shared__ float S_E[sizeof(float) * 320];       
        __shared__ float S_A[sizeof(float) * 320];   

       S_E[threadIdx.x] = E[n];
__syncthreads();
        for (int j = 0; j < 12; j++)
        {           
                S_A[threadIdx.x * 12 + j] = A[n * 12 + j];

                S_A[threadIdx.x * 12 + j]=M[Y[n]]+X[n * 12 + j];     
atomicAdd(&S_E[threadIdx.x],S_A[threadIdx.x * 12 + j]/2); 
        }
__syncthreads();


}
请问如何才能最大程度减轻共享内存S_A和S_E共享内存的bank冲突问题?