在GPU进行最小值归约运行时间比CPU长

Question

函数实现目的：求大小为100的一维数组最小值的下标，使用的归约运算

出现问题：GPU总比CPU运行时间久（GPU：MX150，CPU：i5 8th）

@cuda.jit
def arggetmin(Fitness, IN_index, OutResult, OutIndex, n):
    tid = cuda.threadIdx.x;
    idx = cuda.threadIdx.x + cuda.blockIdx.x * cuda.blockDim.x
    tmp = cuda.shared.array(shape=BLOCK_SIZE, dtype=float32)
    index = cuda.shared.array(shape=BLOCK_SIZE, dtype=int32)
    tmp[tid] = MAX
    if (tid > n): return

    #将数据传入共享内存
    if (idx < n):
        tmp[tid] = Fitness[idx]
        index[tid] = IN_index[idx]
    else:
        tmp[tid] = MAX
        index[tid] = IN_index[idx]

    cuda.syncthreads();

    #没个块内进行归约处理
    stride = int(cuda.blockDim.x / 2)
    while stride > 0:
        if (tid < stride):
            if (tmp[tid] > tmp[tid + stride]):
                tmp[tid] = tmp[tid + stride]
                index[tid] = index[tid + stride]
        cuda.syncthreads()
        stride = int(stride / 2)

    if (tid == 0):
        OutResult[cuda.blockIdx.x] = tmp[0]
        OutIndex[cuda.blockIdx.x] = index[0]

def main():
    n = 100        #数组大小
   
   #生成数组和下标
    a = getRamdomlist(n)
    index = getIndex(n)
    
    #传入到GPU中
    a_device = cuda.to_device(a)
    index_device = cuda.to_device(index)

    threads_per_block = BLOCK_SIZE
    block_per_grid = math.ceil(n / threads_per_block)
    gpu_result = cuda.device_array(shape=block_per_grid, dtype=float) 
    gpu_index = cuda.device_array(shape=block_per_grid, dtype=int)

    #GPU部分
    time1 = perf_counter()
    arggetmin[block_per_grid, threads_per_block](a_device, index_device, gpu_result, gpu_index, n)
    cuda.synchronize()
    time2 = perf_counter()
    print("matmul GPU time :", (time2 - time1))
  
    #CPU部分
    start = perf_counter()
    innn = np.argmin(a)
    end = perf_counter()
    print("matmul CPU time :", (end - start))


if __name__ == "__main__":
    main()

GPU运算时间0.5s左右，CPU运算时间0.0005s左右

使用过“循环展开”的方法，但结果更加糟糕。（运行Nvidia官方的“矩阵乘法”的代码时，使用共享内存的代码比不使用共享内存的代码要慢）

为什么会出现这种情况，一直不太明白

Answer 1

以下答案引用自GPT-3大模型,请合理使用：

示例，谢谢

这可能是因为使用了不正确的共享内存的实现，或者由于没有考虑到适当的优化。如果要使用共享内存，则需要确保使用一致的数据结构和传输方式，以便每个线程可以快速访问内存中的数据。这样可以减少CPU和GPU之间的交互，大大加快计算性能。

此外，还有一些实用的优化技巧，比如常量表达式计算，以及使用特定于 GPU 的指令集和函数，可以使代码更加有效。以下是一个示例，实现本问题的优化版本：

@cuda.jit
def arggetmin_opt(Fitness, IN_index, OutResult, OutIndex, n):
    tid = cuda.threadIdx.x;
    idx = cuda.threadIdx.x + cuda.blockIdx.x * cuda.blockDim.x
    tmp = cuda.shared.array(shape=BLOCK_SIZE, dtype=float32)
    index = cuda.shared.array(shape=BLOCK_SIZE, dtype=int32)
    tmp[tid] = MAX
    if (tid > n): return

    #将数据传入共享内存
    if (idx < n):
        tmp[tid] = Fitness[idx]
        index[tid] = IN_index[idx]
    else:
        tmp[tid] = MAX
        index[tid] = IN_index[idx]

    cuda.syncthreads();

    #没个块内进行归约处理
    stride = cuda.blockDim.x 
    while stride > 0:
        if (tid < stride):
            if (tmp[tid] > tmp[tid + (stride >>1)]):
                tmp[tid] = tmp[tid + (stride >>1)]
                index[tid] = index[tid + (stride >>1)]
        cuda.syncthreads()
        stride = (stride >> 1) 

    if (tid == 0):
        OutResult[cuda.blockIdx.x] = tmp[0]
        OutIndex[cuda.blockIdx.x] = index[0]

如果我的回答解决了您的问题，请采纳我的回答