CUDA：共享内存的性能不佳且没有并行性

Question

我正在尝试利用此内核函数中的共享内存，但性能不如我预期的那么好。这个函数在我的应用程序中被调用了很多次（大约 1000 次或更多），所以我想利用共享内存来避免内存延迟。但显然有些问题是因为我使用共享内存后我的应用程序变得非常慢。
这是内核：

__global__ void AndBitwiseOperation(int* _memory_device, int b1_size, int* b1_memory, int* b2_memory){
int j = 0;

// index GPU - Transaction-wise
unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;
unsigned int tid = threadIdx.x;

// shared variable
extern __shared__ int shared_memory_data[];
extern __shared__ int shared_b1_data[];
extern __shared__ int shared_b2_data[];

// copy from global memory into shared memory and sync threads
shared_b1_data[tid] = b1_memory[tid];
shared_b2_data[tid] = b2_memory[tid];
__syncthreads();

// AND each int bitwise
for(j = 0; j < b1_size; j++)
    shared_memory_data[tid] = (shared_b1_data[tid] & shared_b2_data[tid]);

// write result for this block to global memory
_memory_device[i] = shared_memory_data[i];
}

共享变量被声明为 extern 因为我不知道 b1 和 b2 的大小，因为它们取决于我只能在运行时知道的客户数量（但两者的大小相同一直）。
这就是我调用内核的方式：

void Bitmap::And(const Bitmap &b1, const Bitmap &b2)
{

int* _memory_device;
int* b1_memory;
int* b2_memory;

int b1_size = b1.getIntSize();

// allocate memory on GPU
(cudaMalloc((void **)&b1_memory,  _memSizeInt * SIZE_UINT));
(cudaMalloc((void **)&b2_memory,  _memSizeInt * SIZE_UINT));
(cudaMalloc((void **)&_memory_device,  _memSizeInt * SIZE_UINT));

// copy values on GPU
(cudaMemcpy(b1_memory, b1._memory, _memSizeInt * SIZE_UINT, cudaMemcpyHostToDevice ));
(cudaMemcpy(b2_memory, b2._memory, _memSizeInt * SIZE_UINT, cudaMemcpyHostToDevice ));
(cudaMemcpy(_memory_device, _memory, _memSizeInt * SIZE_UINT, cudaMemcpyHostToDevice ));

dim3 dimBlock(1, 1);
dim3 dimGrid(1, 1);

AndBitwiseOperation<<<dimGrid, dimBlock>>>(_memory_device, b1_size, b1_memory, b2_memory);

// return values
(cudaMemcpy(_memory, _memory_device, _memSizeInt * SIZE_UINT, cudaMemcpyDeviceToHost ));

// Free Memory
(cudaFree(b1_memory));
(cudaFree(b2_memory));
(cudaFree(_memory_device));
}

b1 和 b2 是位图，每个元素有 4 位。元素的数量取决于客户的数量。另外，我对内核的参数有疑问，因为如果我添加一些块或线程，AndBitwiseOperation() 不会给我正确的结果。每个块只有 1 个块和 1 个线程，结果是正确的，但内核不是并行的。
欢迎每一个建议:)
谢谢

Answer 1

当你声明一个extern __shared__ 数组时，你还必须在内核调用中指定它的大小。

内核配置为：

Dg、Db、Ns、S >>>

Ns是extern __shared__数组的大小，默认为0。

我认为您不能在内核中定义多个extern __shared__ 数组。编程指南中的一个示例定义了一个 extern __shared__ 数组并手动设置其中包含偏移量的数组：

extern __shared__ float array[]; 
__device__ void func()      // __device__ or __global__ function 
{ 
    short* array0 = (short*)array;  
    float* array1 = (float*)&array0[128]; 
    int*   array2 =   (int*)&array1[64]; 
} 

Answer 2

我并没有真正理解你的内核想要做什么。

您应该阅读有关 CUDA 和 GPU 编程的更多信息。

我试图指出一些错误。

1. 共享内存 (sm) 应减少全局内存读取。 分析每个线程的全局内存 (gm) 读写操作。

   一个。你读了两次全局内存，写了两次 sm
   湾。 （废话循环被忽略，不使用索引）你读了两次sn，写了一次sm
   C。你读一次sm，写一次gm

   所以总的来说，你没有任何收获。可以直接使用全局内存。

2. 您使用所有线程在块索引“i”处写出一个值。 您应该只使用一个线程来写出这些数据。
   将序列化的多个线程输出相同的数据是没有意义的。

3. 您使用了循环，根本不使用循环计数器。

4. 你在“tid”写，在“i”随机读。

5. 这个分配是开销。

   unsigned int tid = threadIdx.x;
   

6. 超过一个块的结果不能正确，因为一个块 tid = i!
   所有错误的索引都会导致使用不止一个块的错误计算

7. “i”处的共享内存从未被写入！

   _memory_device[i] = shared_memory_data[i];
   

我的假设你的内核应该做什么

/*
 * Call kernel with x-block usage and up to 3D Grid
 */
__global__ void bitwiseAnd(int* outData_g, 
    const long long int inSize_s, 
    const int* inData1_g, 
    const int* inData2_g)
{
    //get unique block index
    const unsigned long long int blockId = blockIdx.x //1D
        + blockIdx.y * gridDim.x //2D
        + gridDim.x * gridDim.y * blockIdx.z; //3D

    //get unique thread index
    const unsigned long long int threadId = blockId * blockDim.x + threadIdx.x; 

    //check global unique thread range
    if(threadId >= inSize_s)
        return;

    //output bitwise and
    outData_g[thread] = inData1_g[thread] & inData2_g[thread];
}