CUDA学习笔记7——CUDA内存组织

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.h"

int 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"
#include<math.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_x>n || step_y>m)
		{
			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_y>=n)
		{
			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 (x<k && y<m)
	{
		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_matrix<<<dimGrid, 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;
}