网站手机优化,wordpress让投稿,自治区住房和城乡建设部网站,韩国的 电子商务网站CUDA内存组织
CUDA设备内存的分类与特征
内存类型物理位置访问权限可见范围生命周期1全局内存芯片外可读写所有线程和主机端由主机分配与释放2常量内存芯片外只读所有线程和主机端由主机分配与释放3纹理和表面内存芯片外一般只读所有线程和主机端由主机分配与释放4寄存器内存…CUDA内存组织
CUDA设备内存的分类与特征
内存类型物理位置访问权限可见范围生命周期1全局内存芯片外可读写所有线程和主机端由主机分配与释放2常量内存芯片外只读所有线程和主机端由主机分配与释放3纹理和表面内存芯片外一般只读所有线程和主机端由主机分配与释放4寄存器内存芯片内可读写单个线程所在线程5局部内存芯片外可读性单个线程所在线程6共享内存芯片内可读性单个线程块所在线程块 全局内存核函数中所有线程都能访问其中的数据。 用cudaMalloc()为全局内存变量分配设备内存 用cudaMemcpy()将主机数据复制到全局内存 常量内存一共64KB只读可见范围与生命周期与全局内存一样访问速度比全局内存快在核函数未满用 _constant_ 定义变量并使用cudaMemcpyToSymbol()将数据从主机端复制到设备的常量内存。 纹理内存与表面内存类似于常量内存可见范围与生命周期相同; 寄存器在核函数中定义的不加任何限定符的变量一般来说放在寄存器中核函数定义不加任何限定符的数组可能放于寄存器也可能放于局部内存中 局部内存寄存器放不下的变量索引值不能在编译时确定的数组 共享内存与寄存器类似存在于芯片上仅次于寄存器的读写速度
CUDA中的内存组织示意图 GPU设备规格查询
#include stdio.h
#include cuda_runtime.h
#include device_launch_parameters.hint main()
{int device_id 0;cudaDeviceProp prop;cudaGetDeviceProperties(prop, device_id);printf(Device id: %d\n, device_id);printf(Device name: %s\n, prop.name);printf(Compute capability: %d.%d\n, prop.major, prop.minor);printf(Amount of global memory: %g GB\n, prop.totalGlobalMem / 1024.0);printf(Amount of constant memory: %g KB\n, prop.totalConstMem / 1024.0);printf(Maximum grid size: %d %d %d\n,prop.maxGridSize[0], prop.maxGridSize[1], prop.maxGridSize[2]);printf(Maximum block size: %d %d %d\n, prop.maxThreadsDim[0], prop.maxThreadsDim[1], prop.maxThreadsDim[2]);printf(Number of SMs: %d\n, prop.multiProcessorCount);printf(----------------------------- \n);printf(Maximum amount of shared memory per block: %g KB\n, prop.sharedMemPerBlock / 1024.0);printf(Maximum amount of shared memory per SM: %g KB\n,prop.sharedMemPerMultiprocessor / 1024.0);printf(Maximum number of registers per block: %d K\n, prop.regsPerBlock / 1024.0);printf(Maximum number of registers per SM: %d K\n, prop.regsPerMultiprocessor / 1024.0);printf(Maximum number of threads per block: %d \n, prop.maxThreadsPerBlock);printf(Maximum number of threads per SM: %d \n, prop.maxThreadsPerMultiProcessor);return 0;
} 全局内存的合并与非合并访问
合并访问一个线程束对全局内存的一次访问读/写导致最少数量的数据传输否则为非合并访问。
利用共享内存和统一内存优化矩阵乘 #include stdio.h
#include cuda_runtime.h
#include device_launch_parameters.h
#includemath.h
#include malloc.h
#include opencv2/opencv.hpp
#include stdlib.h//利用share memory 和统一内存优化矩阵乘#define M 1000
#define N 500
#define K 1000__managed__ int a[M*N];
__managed__ int b[N*K];
__managed__ int c_gpu[M*K];
__managed__ int c_cpu[M*N];#define BLOCK_SIZE 16__global__ void gpu_matrix(int* a, int* b, int* c, int m, int n, int k)
{__shared__ int sub_a[BLOCK_SIZE][BLOCK_SIZE];__shared__ int sub_b[BLOCK_SIZE][BLOCK_SIZE];int x blockIdx.x*blockDim.x threadIdx.x;int y blockIdx.y*blockDim.y threadIdx.y;int tmp 0;int idx;for (int step 0; step N/BLOCK_SIZE; step){int step_x step*BLOCK_SIZE threadIdx.x;int step_y y;idx step_y*n step_x;if (step_xn || step_ym){sub_a[threadIdx.y][threadIdx.x] 0;}else{sub_a[threadIdx.x][threadIdx.x] a[idx];}step_x x;step_y step*BLOCK_SIZE threadIdx.y;idx step * k step_x;if (step_x k || step_yn){sub_b[threadIdx.y][threadIdx.x] 0;}else{sub_b[threadIdx.y][threadIdx.x] b[idx];}__syncthreads();for (int i 0; i BLOCK_SIZE; i){tmp sub_a[threadIdx.y][i] * sub_b[i][threadIdx.x];}__syncthreads();}if (xk ym){c[y*k x] tmp;}}void cpu_matrix(int* a, int* b, int* c, int m, int n, int k)
{for (int y 0; y m; y){for (int x 0; x k; x){int tmp 0;for (int step 0; step n; step){tmp a[y*n step] * b[step*n x];}c[y*k x] tmp;}}}int main()
{for (int y 0; y M; y){for (int x 0; x N; x){a[y * N x] rand() % 1024;}}for (int y 0; y N; y){for (int x 0; x K; x){b[y*K x] rand() % 1024;}}unsigned int grid_x (K BLOCK_SIZE - 1) / BLOCK_SIZE;unsigned int grid_y (M BLOCK_SIZE - 1) / BLOCK_SIZE;dim3 dimGrid(grid_x, grid_y);dim3 dimBlock(BLOCK_SIZE, BLOCK_SIZE);gpu_matrixdimGrid, dimBlock(a, b, c_gpu, M, N, K);cpu_matrix(a, b, c_cpu, M, N, K);bool errors false;for (int y 0; y M; y){for (int x 0; x K; x){if (fabs(c_cpu[y*K x] - c_gpu[y*K x]) (1.0e-10)){errors true;}}}printf(Result: %s\n, errors ? Error : Pass);return 0;
}
