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) 读写操作。

   a) 你读了两次GM，写了两次SM。

   b）（忽略无意义循环，不使用索引）您读取 SM 两次并写入 SM 一次。

   c) 读一次SM，写一次GM。

   所以总的来说你一无所获。直接用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];
}