小编use*_*912的帖子

我对Cuda相对较新,我正在尝试编写一个内核,用于计算查询向量和大型向量数据库之间的绝对差值之和.两者的元素必须是8位无符号整数.我已经将我的内核基于nvidias示例并行缩减内核,我也读过这个帖子.

我只得到大约5GB/s,这比快速CPU好多了,甚至没有接近我的DDR5 GT640的理论带宽大约80GB/s.

我的数据集包含1024字节的查询向量,100,000 x 1024字节的数据库

我有100,000个128个线程的块,如果每个块访问相同的1024字节query_vector,那会导致更糟糕的性能吗？由于每个块都访问相同的内存位置.

blockSize和共享内存都设置为128和128*sizeof(int),128是#define'd为THREADS_PER_BLOCK

template<UINT blockSize> __global__ void reduction_sum_abs( BYTE* query_vector, BYTE* db_vector, uint32_t* result )
{
    extern __shared__ UINT sum[]; 
    UINT db_linear_index = (blockIdx.y*gridDim.x) + blockIdx.x ; 
    UINT i = threadIdx.x; 

    sum[threadIdx.x] = 0; 

    int* p_q_int = reinterpret_cast<int*>(query_vector); 
    int* p_db_int = reinterpret_cast<int*>(db_vector); 

    while( i < VECTOR_SIZE/4 ) {

        /* memory transaction */
        int q_int = p_q_int[i]; 
        int db_int = p_db_int[db_linear_index*VECTOR_SIZE/4 + i]; 

        uchar4 a0 = *reinterpret_cast<uchar4*>(&q_int); 
        uchar4 b0 = *reinterpret_cast<uchar4*>(&db_int); 

        /* sum of …