版权声明:林微原创,未经允许不得转载。 https://blog.csdn.net/Canhui_WANG/article/details/82913993
本文主要介绍基于CUDA的向量逆向,并结合三种CUDA异构内存进行相应优化。
1. 向量逆向
基于 Global memory 的核心算法
Global memory空间大,但是速度是所有Device内存中最慢的
__global__ void array_reverse(int *array_a_dev, int *array_a_rev_dev, int len)
{
int tid = threadIdx.x;
array_a_rev_dev[len - tid - 1] = array_a_dev[tid];
}
基于 Shared memory 的核心算法
当存在多个线程共同访问同一个数据的时候,可以将该数据放到shared memory中,以提高数据读取效率
__global__ void array_reverse_shared(int *array_a_dev, int *array_a_rev_dev, int len)
{
int tid = threadIdx.x;
__shared__ int array_shared[9];
array_shared[tid] = array_a_dev[tid];
__syncthreads();
array_a_rev_dev[len - tid - 1] = array_shared[tid];
}
基于Dynamic shared memory 核心算法
Dynamic shared memory 通过采用extern
的声明而不定义的方式,提高了代码的灵活性,比方说,dynamic shared memory可以在同一个程序的不同kernel中被定义,也可以动态地根据需求进行shared memory内存空间的分配。如下,
__global__ void array_reverse_dynamic_shared(int *array_a_dev, int *array_a_rev_dev, int len)
{
int tid = threadIdx.x;
extern __shared__ int array_shared[];
// __shared__ int array_shared[9];
array_shared[tid] = array_a_dev[tid];
__syncthreads();
array_a_rev_dev[len - tid - 1] = array_shared[tid];
}
Dynamic shared memory的kernel的调用有些不同,
array_reverse_dynamic_shared<<<dimGrid, dimBlock, len*sizeof(int)>>>(array_a_dev, array_a_rev_dev, len);
注:动态变化 (或者说,非预编译的变量) 的内存空间大小 (比如,这里的len*sizeof(int)) 需要在kernel函数调用时定义。
2. 编译调试
基于 Global memory 的代码
源代码:array_rev.cu
编译
nvcc array_rev.cu -o main
运行
./main
基于 Shared memory 的代码
源代码:array_rev_sm.cu
编译
nvcc array_rev_sm.cu -o main
运行
./main
基于 Dynamic shared memory 的代码
编译
nvcc array_rev_dynamic_sm.cu -o main
运行
./main