1

我试图在这个内核函数中利用共享内存，但性能不如我预期的那么好。这个函数在我的应用程序中被调用了很多次（大约 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 个线程，结果是正确的，但内核不是并行的。
欢迎每一个建议:)
谢谢

4

2 回答 2

4

当你声明一个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]; 
}
于 2012-04-21T23:32:07.057 回答
4

我真的不明白你的内核想要做什么。

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

我试图指出一些错误。

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

    一种。你读了两次全局内存，写了两次 sm
    b. （废话循环被忽略，不使用索引）你读了两次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];
}
于 2012-04-22T02:52:20.330 回答