CUDA:矩阵转置的GPU实现(Share Memory)

这篇具有很好参考价值的文章主要介绍了CUDA:矩阵转置的GPU实现(Share Memory)。希望对大家有所帮助。如果存在错误或未考虑完全的地方,请大家不吝赐教,您也可以点击"举报违法"按钮提交疑问。

本文参加2022CUDA on Platform线上训练营学习笔记

欢迎各位大犇提意见

一、矩阵转置(Matrix Transpose)基础

cuda矩阵转置,CUDA_On_Arm_2022夏令营,矩阵,算法,线性代数,人工智能
上图中将m * n的矩阵A通过矩阵转置变成了n * m的 AT,简单来讲矩阵转置即为将原始矩阵的第一行转置为目标矩阵的第一列,以此类推,相信基础扎实的你简单地看看CPU端的代码就能理解

二、矩阵转置的CPU端实现

__host__ void cpu_transpose(int *matrix,int *tr_matrix,int m,int n) {
	for (int i = 0; i < n; i++) {
		for (int j = 0; j < m; j++) {
			tr_matrix[i * m + j] = matrix[j * n + i];
		}
	}
	return;
}

定义一个名为cpu_transpose的函数,将矩阵matrix转置为矩阵tr_matrix,通过观察代码不难发现tr_matrix[i][j]=matrix[j][i],这里需要注意到的是坐标的转换,转置后的矩阵行数和列数发生变换,留意m和n不要乘错了。

原始矩阵:
cuda矩阵转置,CUDA_On_Arm_2022夏令营,矩阵,算法,线性代数,人工智能
CPU端执行结果:
cuda矩阵转置,CUDA_On_Arm_2022夏令营,矩阵,算法,线性代数,人工智能

三、矩阵转置的GPU端实现(share Memory)

1、核函数的编写

GPU端的实现与CPU端类似,首先根据各个线程的index(索引)计算出当前线程在原始矩阵中的位置rowcol,在原始矩阵中的row行,col

	int row = blockDim.y * blockIdx.y + threadIdx.y;
	int col = blockDim.x * blockIdx.x + threadIdx.x;

下边我们申请同一个block中的线程可以访问的shared Memory

__shared__ int smem_matrix[BLOCK_SIZE][BLOCK_SIZE];

在GPU中申请了一块名为smem_matrix大小为sizeof(int)*BLOCK_SIZE^2的共享内存,在执行赋值操作之前将当前block中的线程需要访问到的数据从Global_Memory中复制到share_Memory

smem_matrix[threadIdx.y][threadIdx.x] = row < m&& col < n ? matrix[row*n+col] : 0;

赋值时需要注意的是:由于我们为内核函数设置执行配置的时候通常会向上取整,会申请多于实际需求的线程数,所以在我们赋值之前需要判断当前线程的坐标是否是需求坐标,以此来防止访问matrixrow*n+col成为野指针,对我们的数据造成重大的危害
有了同一个block中的线程申请一个share Memory的概念后,需要做的是同步同一个BLock中的线程

__syncthreads();

通过上边一系列的操作,我们就可以开始真正的转置操作了,需要注意的是,我们已经把线程所需的数据赋值到share Memory当中,所以我们在赋值时只需调用smem_matrix,同样,赋值操作之前,我们需要判断当前的坐标是否实际有效

	if(blockIdx.x * blockDim.x + threadIdx.y < n && threadIdx.x + blockIdx.y * blockDim.x < m)
	tr_matrix[threadIdx.x+blockIdx.y*blockDim.x+m*(blockIdx.x*blockDim.x+threadIdx.y)] = smem_matrix[threadIdx.x][threadIdx.y];

上述分析使我们获得了完整的GPU代码

__global__ void cuda_transpose(int *matrix,int *tr_matrix,int m,int n) {
	int row = blockDim.y * blockIdx.y + threadIdx.y;
	int col = blockDim.x * blockIdx.x + threadIdx.x;
	__shared__ int smem_matrix[BLOCK_SIZE][BLOCK_SIZE];
	smem_matrix[threadIdx.y][threadIdx.x] = row < m&& col < n ? matrix[row*n+col] : 0;
	__syncthreads();
	if(blockIdx.x * blockDim.x + threadIdx.y < n && threadIdx.x + blockIdx.y * blockDim.x < m)
	tr_matrix[threadIdx.x+blockIdx.y*blockDim.x+m*(blockIdx.x*blockDim.x+threadIdx.y)] = smem_matrix[threadIdx.x][threadIdx.y];
	return;
}

2、核函数的启动

在设备端申请两个指针并为其分配内存

	int* d_matrix, *dtr_matrix;
	cudaMalloc((void**)&d_matrix, sizeof(int) * m * n);
	cudaMalloc((void**)&dtr_matrix, sizeof(int) * m * n);

手动将matrix中的数据通过Pcie复制到设备端的Global Memory当中

cudaMemcpy(d_matrix, matrix, sizeof(int) * m * n, cudaMemcpyHostToDevice);

核函数执行设置的设定,一个warp通常为32个线程所以我们一个Block中的线程数最好设置为32的整数倍,从此提高使用率,有效防止inactive code的出现

dim3 block = { BLOCK_SIZE,BLOCK_SIZE,1 }; //BLOCK_SIZE = 16

gridDim的设置最需关注的就是申请的线程能够有效的覆盖真个矩阵,宁可多申请,通过核函数中的if屏蔽,也不少申请,导致计算的缺失,所以我们在计算中采用向上取整的方法
需要注意的使 dim3 类型中的三个成员都是要求unsigned int 类型的所以我们在前面添加(unsigned int)来强制将我们的计算结果转换为无符号

dim3 gird = { (unsigned int)(n - 1 + BLOCK_SIZE) / BLOCK_SIZE, (unsigned int)(m - 1 + BLOCK_SIZE) / BLOCK_SIZE,1 };

核函数启动!

cuda_transpose << < gird , block  >> > (d_matrix, dtr_matrix, m, n);

3、核函数性能计数

在CUDA中有一种特殊的类型cudaEvent_t,可以帮助我们记录核函数的执行信息

	cudaEvent_t kernel_start;
	cudaEvent_t kernel_end;

	cudaEventCreate(&kernel_start);
	cudaEventCreate(&kernel_end);

kernel_start用于记录核函数开始执行时的信息,kernel_end用来记录核函数运行结束时的信息,这里使用到了两个函数cudaEventQuery(kernel_start);,cudaEventSynchronize(kernel_end);,前者是非阻塞的,只要执行到就直接记录,后者是阻塞式的,需要前面的执行完毕才能运行,具体的性能计数函数如下
cuda矩阵转置,CUDA_On_Arm_2022夏令营,矩阵,算法,线性代数,人工智能
通过简单的逻辑组合,就可以得到核函数的实际运行时间,具体代码如下

	cudaEventCreate(&kernel_start);
	cudaEventCreate(&kernel_end);
	cudaEventRecord(kernel_start);
	cudaEventQuery(kernel_start);
	cuda_transpose << < gird , block  >> > (d_matrix, dtr_matrix, m, n);
	cudaEventRecord(kernel_end);
	cudaEventSynchronize(kernel_end);
	float ms;
	cudaEventElapsedTime(&ms, kernel_start, kernel_end);

四、代码参考

#include <cuda_runtime.h>
#include <device_launch_parameters.h>
#include <stdio.h>
#include <stdlib.h>
#include <string>
#include <iostream>
#define BLOCK_SIZE 32
using namespace std;


__global__ void cuda_transpose(int *matrix,int *tr_matrix,int m,int n) {
	int row = blockDim.y * blockIdx.y + threadIdx.y;
	int col = blockDim.x * blockIdx.x + threadIdx.x;
	__shared__ int smem_matrix[BLOCK_SIZE][BLOCK_SIZE];
	smem_matrix[threadIdx.y][threadIdx.x] = row < m&& col < n ? matrix[row*n+col] : 0;
	__syncthreads();
	if(blockIdx.x * blockDim.x + threadIdx.y < n && threadIdx.x + blockIdx.y * blockDim.x < m)
	tr_matrix[threadIdx.x+blockIdx.y*blockDim.x+m*(blockIdx.x*blockDim.x+threadIdx.y)] = smem_matrix[threadIdx.x][threadIdx.y];
	return;
}

__host__ void cpu_transpose(int *matrix,int *tr_matrix,int m,int n) {
	for (int i = 0; i < n; i++) {
		for (int j = 0; j < m; j++) {
			tr_matrix[i * m + j] = matrix[j * n + i];
		}
	}
	return;
}

__host__ void init_matrix(int* matrix,int m,int n) {
	for (int i = 0; i < m; i++) {
		for (int j = 0; j < n; j++) {
			matrix[i*n+j] = rand();
		}
	}
}

void print(int*, string,int,int);
bool check(int*, int*, int, int);

int main() {
	int m = 1111;
	int n = 113;
	int *matrix;
	cudaMallocHost((void**)&matrix, sizeof(int) * m * n);
	init_matrix(matrix,m,n);
	//print(matrix, "init matrix", m, n);


	int* htr_matrix;
	cudaMallocHost((void**)&htr_matrix, sizeof(int) * m * n);
	cpu_transpose(matrix, htr_matrix, m, n);
	//print(htr_matrix, "CPU", n, m);
	//将CPU端执行的结果存放在htr_matrix中 

	int* d_matrix, *dtr_matrix;
	cudaMalloc((void**)&d_matrix, sizeof(int) * m * n);
	cudaMalloc((void**)&dtr_matrix, sizeof(int) * m * n);
	cudaMemcpy(d_matrix, matrix, sizeof(int) * m * n, cudaMemcpyHostToDevice);
	dim3 gird = { (unsigned int)(n - 1 + BLOCK_SIZE) / BLOCK_SIZE, (unsigned int)(m - 1 + BLOCK_SIZE) / BLOCK_SIZE,1 };
	dim3 block = { BLOCK_SIZE,BLOCK_SIZE,1 };


	cudaEvent_t kernel_start;
	cudaEvent_t kernel_end;

	cudaEventCreate(&kernel_start);
	cudaEventCreate(&kernel_end);
	cudaEventRecord(kernel_start);
	cudaEventQuery(kernel_start);
	cuda_transpose << < gird , block  >> > (d_matrix, dtr_matrix, m, n);
	cudaEventRecord(kernel_end);
	cudaEventSynchronize(kernel_end);
	float ms;
	cudaEventElapsedTime(&ms, kernel_start, kernel_end);


	int* hdtr_matrix;
	cudaMallocHost((void**)&hdtr_matrix, sizeof(int) * m * n);
	cudaMemcpy(hdtr_matrix, dtr_matrix, sizeof(int) * m * n, cudaMemcpyDeviceToDevice);
	//print(hdtr_matrix, "GPU", n, m);
	
	if (check(hdtr_matrix, htr_matrix, n, m)) {
		cout << "pass\n";
	}
	else {
		cout << "error\n";
	}
	
	printf("GPU time is : %f \n", ms);

	cudaFree(hdtr_matrix);
	cudaFree(dtr_matrix);
	cudaFree(matrix);
	cudaFree(htr_matrix);
	cudaFree(d_matrix);
	return 0;
}


void print(int* a, string name,int m,int n) {
	cout << "NAME : " << name << endl;
	for (int i = 0; i < m; i++) {
		for (int j = 0; j < n; j++) {
			printf("%6d ", a[i * n + j]);
		}
		printf("\n");
	}
}

bool check(int* a, int* b, int m, int n) {
	bool check_flag = true;
	for (int i = 0; i < m; i++) {
		for (int j = 0; j < n; j++) {
			if (a[i * n + j] != b[i * n + j]) {
				return false;
			}
		}
	}
	return check_flag;
}


执行结果如图
cuda矩阵转置,CUDA_On_Arm_2022夏令营,矩阵,算法,线性代数,人工智能

五、实践心得

本次实践通过GPU端中的share Memory对核函数运行时的读写问题做了优化,当线程与线程之间为连续读写时,global Memory的效率是比较高的,不使用share Memory时,使用GPU进行矩阵转置会出现两难问题(1.读row-major 写col-major,2写col-major 读row-major),而在share Memoryrow-majorcol-major的效率几乎相同,很好地解决了global memory上的问题,在编写过程中,需要注意的是,要顺着global memory写首先保证global memory读写时是row-major,以达到最高的优化效率。
遇到的最大问题是,边界的判断问题,GPU转置过程中,由于要保证global memoryrow-major,所以坐标不像是CPU端中的简单调换,具体表现为(在对share 数字赋值时该线程无意义,而在写global操作中该线程有意义),所以在__syncthreads();后需要判断当前线程是否有意义

鄙人第一次写实操博客,有建议必洗耳恭听
再次感谢伟大的NV 开发者社区文章来源地址https://www.toymoban.com/news/detail-532243.html

到了这里,关于CUDA:矩阵转置的GPU实现(Share Memory)的文章就介绍完了。如果您还想了解更多内容,请在右上角搜索TOY模板网以前的文章或继续浏览下面的相关文章,希望大家以后多多支持TOY模板网!

本文来自互联网用户投稿,该文观点仅代表作者本人,不代表本站立场。本站仅提供信息存储空间服务,不拥有所有权,不承担相关法律责任。如若转载,请注明出处: 如若内容造成侵权/违法违规/事实不符,请点击违法举报进行投诉反馈,一经查实,立即删除!

领支付宝红包 赞助服务器费用

相关文章

  • CUDA编程模型系列六(利用shared memory和统一内存优化矩阵乘)

    CUDA编程模型系列六(利用shared memory和统一内存优化矩阵乘) 本系列教程将介绍具体的CUDA编程代码的细节 CUDA编程模型系列六(利用shared memory和统一内存优化矩阵乘)

    2024年02月11日
    浏览(50)
  • CUDA编程:矩阵乘运算从CPU到GPU

    本文内容涉及到CUDA矩阵1D运算、2D运算、共享内存、CUBLAS的使用。 文中的全部code: https://github.com/CalvinXKY/BasicCUDA/tree/master/matrix_multiply V100上的测试对比: 运行内容“./matMul wA=1024 hA=256 wB=128 hB=1024” 矩阵 C = A x B的数学运算,是线性代数里面最基本的内容, 计算的基本公式如下

    2024年04月08日
    浏览(55)
  • 解决:RuntimeError: CUDA out of memory. Tried to allocate 160.00 MiB (GPU 0; 10.76 GiB total capacity..

    完整报错:   问题分析: 内存分配不足: 需要160MB,,但GPU只剩下135.31MB。 解决办法: 1.减小batch_size。注意batchsize的调整要配合学习率的调整,一般是正比关系,BS增大两倍,LR增大两倍或者根号二倍。减小也是相应更改。 2.运行torch.cuda.empty_cache()函数。加在训练开始前即可

    2024年02月16日
    浏览(50)
  • torch.cuda.OutOfMemoryError: CUDA out of memory.

    训练清华ChatGLM-6B时报错, 原因是显存不够 torch.cuda.OutOfMemoryError: CUDA out of memory. Tried to allocate 96.00 MiB (GPU 0; 23.70 GiB total capacity; 4.37 GiB already allocated; 64.81 MiB free; 4.37 GiB reserved in total by PyTorch) If reserved memory is allocated memory try setting max_split_size_mb to avoid fragmentation.  See documentatio

    2024年02月06日
    浏览(48)
  • Pycharm报错torch.cuda.OutOfMemoryError: CUDA out of memory.

    报错 做深度学习相关的实验,可以看到我的显卡内存很小(哭了,不过我有时候是在别的电脑上做的,那个电脑比这个好用),网上搜到的说的 max_split_size_mb:128 这个方法我贴到我代码上之后没有效果。 因为我在这个电脑上做的是主实验后面的一些对比实验,也就是代码中很

    2024年02月05日
    浏览(46)
  • 【CUDA OUT OF MEMORY】【Pytorch】计算图与CUDA OOM

    在实践过程中多次碰到了CUDA OOM的问题,有时候这个问题是很好解决的,有时候DEBUG一整天还是头皮发麻。 最近实践对由于计算图积累导致CUDA OOM有一点新的看法,写下来记录一下。 包括对计算图的一些看法和一个由于计算图引发错误的简化实例记录。 本人能力有限,认识片

    2024年02月09日
    浏览(39)
  • CUDA报错:Out of Memory

    如果报错里提示Pytorch reserved的内存远大于Already allocated的内存,那么就是因为分配显存时单位过大,导致出现大量内存碎片无法继续分配(与操作系统内存管理同理)。 我们可以限制一次分配的最大单位来解决这个问题。 随后代码便可正常运行了。

    2024年02月15日
    浏览(57)
  • 部署stable diffusion 错误torch.cuda.OutOfMemoryError: CUDA out of memory.

    以来安装完毕,开始执行web_ui.bat 错误截图:  猜测原因:GPU用错了 webUI.py加一行代码 在此启动web_ui.bat,成功打开网页页面

    2024年02月11日
    浏览(51)
  • RuntimeError: CUDA out of memory See documentation for Memory Management and PYTORCH_CUDA_ALLOC_CONF

    报错: If reserved memory is allocated memory try setting max_split_size_mb to avoid fragmentation. See documentation for Memory Management and PYTORCH_CUDA_ALLOC_CONF 当reserved memory is allocated memory,进行如下设置,可解决此bug: 代码如下:

    2024年02月11日
    浏览(52)
  • 【CUDA】GPU 算力与 CUDA 版本对应关系

    官方算力表:https://developer.nvidia.com/cuda-gpus#compute 2.1. 信息来源 1 https://docs.nvidia.com/datacenter/tesla/drivers/index.html#cuda-arch-matrix 2.2. 信息来源 2 https://en.wikipedia.org/wiki/CUDA#GPUs_supported

    2024年01月19日
    浏览(39)

觉得文章有用就打赏一下文章作者

支付宝扫一扫打赏

博客赞助

微信扫一扫打赏

请作者喝杯咖啡吧~博客赞助

支付宝扫一扫领取红包,优惠每天领

二维码1

领取红包

二维码2

领红包