利用共享存储单元优化应用利用共享存储单元进行矩阵转

简介: 利用共享存储单元优化应用利用共享存储单元进行矩阵转

Udacity的CUDA编程课程中介绍了CUDA实现矩阵转置的六种方式,本文介绍其中的一种方式

如果矩阵为N*N的方阵。该方式让每个线程处理一个矩阵元素,总共需要N*N个线程。首先,声明两个常量并配置blocks,threads:

const int N=1024;
const int K=32;
dim3 blocks(N/K,N/K); 
dim3 threads(K,K);

内核函数:

__global__ void 
transpose_parallel_per_element_tiled(float in[], float out[])
{
  // (i,j) locations of the tile corners for input & output matrices:
  int in_corner_i  = blockIdx.x * blockDim.x, in_corner_j  = blockIdx.y * blockDim.y;
  int out_corner_i = in_corner_j, out_corner_j = in_corner_i;
  int x = threadIdx.x, y = threadIdx.y;
  __shared__ float tile[K][K];
  // coalesced read from global mem, TRANSPOSED write into shared mem:
  tile[y][x] = in[(in_corner_i + x) + (in_corner_j + y)*N];
  __syncthreads();
  // read from shared mem, coalesced write to global mem:
  out[(out_corner_i + x) + (out_corner_j + y)*N] = tile[x][y];
}

内核涉及两个输入参数,in代表输入矩阵,out代表in的转置矩阵。为了更好地理解这段代码,我们将矩阵规模缩小并画图展示。假设N=4,K=2。blocks(2,2),threads(2,2)。

如图所示,矩阵被分成了4个区域,每个区域包含2*2=4个数据,分别由4个线程块进行处理。

矩阵左下那个方形区域由block(0,1)处理。该线程块corner(threads(0,0))所处理数据在in矩阵中的坐标为。S区域的其他数据坐标可由[(icj+y)*N+ici+x]得到,x,y是线程索引。blocks,threads的二维索引值x,y分别代表了矩阵的水平方向和竖直方向。

内核声明了一个和数据块大小相同的共享内存tile,用于存储S区域的数据。block读取S区域数据并按原样复制到tile当中,一个线程复制一个数据。随后block将tile中的数据转置,再把经过转置的tile中的数据写入到out矩阵当中区域。根据转置行列互换的性质可知,block(0,1) corner所处理数据在out矩阵中的坐标为(oci=icj,ocj=ici)即(2,0),区域的其他数据索引可由[(ocj+y)*N+oci+x]得到。转置后,in矩阵区域的数据会被block(0,1)转置到out矩阵。tile转置和写入数据到out是由最后一行代码完成的。

该例中,数据读取和写入都利用了全局内存合并访问,这也说明,对于二维矩阵块(x,y),CUDA按x优先顺序处理线程。


出于某种目的,我需要线程块的尺寸为非方形。假设blocks(4,2),threads(1,2)。内核函数__shared__ float tile[K][K];需要修改为__shared__ float tile[2][1]; 这样修改之后,问题在于无法在tile中对数据进行转置。所以我在将数据写入到out的时候,将位置索引[(oci + x) + (ocj + y)*N]转置为(oci + y) + (ocj + x)*N,实现非方矩阵块的矩阵转置。

这段解释不要看了。内核函数中的in_corner_i,in_corner_j分别代表线程块左上角线程(也就是threads(0,0))在矩阵中的x索引(水平方向)和y索引(竖直方向),作为线程块内部线程索引数据的参照点(线程块(0,0)线程所处理数据在矩阵中的位置)。对于block(0,1),in_corner_i=0,in_corner_j=2。转置后,S区域内的元素会被转置到右上方区域S’。所以block(0,1)的功能就是把in矩阵S区域的数据转置并复制到out矩阵S'区域。这就需要block(0,1)找到out矩阵右上方区域的参照点索引。根据转置行列互换的性质,得到out_corner_i(x索引)=in_corner_j=2,out_corner_j(y索引)=in_corner_i=0。x,y变量分别是线程在线程块中的二维索引。

共享内存tile的尺寸和block一样,block读取in数据并按原样复制到tile当中,一个线程复制一个数据。同步后,再把tile中的数据按转置顺序写入到out矩阵当中的对应位置。转置的过程相当于在共享内存中完成,所以线程读取和写入都用[(in_corner_i + x) + (in_corner_j + y)*N]来索引数据。

 

目录
打赏
0
0
0
0
0
分享
相关文章
程序员的自我修养—链接、装载与库--书签目录PDF
程序员的自我修养—链接、装载与库--书签目录PDF
1441 0
初探POC编写
初探POC编写
435 0
初探POC编写
深度揭秘超长序列生成任务训练技术
阿里自研的TorchAcc训练引擎提出了超长序列训练方案FlashSequence,针对超长文本理解、视频生成等场景。通过2D Context Parallel和Hybrid FSDP混合分布式策略,结合显存、计算和通信优化,实现了百万级别超长序列模型的高效训练。FlashSequence在算力、显存需求及分布式训练方面进行了多项创新,性能提升显著,最大可达48%。该方案大幅降低了企业创新成本,提升了业务应用的可能性。
基于qwen2.5的长文本解析、数据预测与趋势分析、代码生成能力赋能esg报告分析
Qwen2.5是一款强大的生成式预训练语言模型,擅长自然语言理解和生成,支持长文本解析、数据预测、代码生成等复杂任务。Qwen-Long作为其变体,专为长上下文场景优化,适用于大型文档处理、知识图谱构建等。Qwen2.5在ESG报告解析、多Agent协作、数学模型生成等方面表现出色,提供灵活且高效的解决方案。
750 49
|
10月前
SublimeText配置Markdown编辑及预览
本文详细介绍了如何配置Sublime Text及相关插件,使之成为Markdown编辑器并且能够在浏览器中实现预览功能。
|
8月前
|
Ubuntu禁止内核自动更新
通过上述步骤,您可以在Ubuntu系统中有效地禁用内核的自动更新。这些步骤包括锁定内核版本、禁用自动更新配置、移除不需要的内核包以及禁用相关的自动更新服务。这样可以确保系统在内核层面保持稳定,避免因内核自动更新导致的不必要问题。
1710 1
为何人们喜欢推理胜于训练大模型?
在AI和机器学习领域,越来越多的人转向重视推理而非大规模模型训练。推理的即时性和高效性使其在需要快速响应的场景中占优,如自然语言处理和图像识别。推理过程的可视化能帮助用户理解模型决策,便于调试和提升性能。此外,推理在边缘计算和移动设备上的应用降低了延迟和带宽成本,同时保护了用户隐私。相比于训练大模型的高资源消耗,推理更为节能且成本效益高,尤其在数据挖掘和新知识探索方面展现出创新潜力。推理在实际应用中与训练模型相结合,提供了性能与成本的有效平衡。随着技术进步,推理将在推动人工智能领域发展中发挥更大作用。
AI助理

你好,我是AI助理

可以解答问题、推荐解决方案等

登录插画

登录以查看您的控制台资源

管理云资源
状态一览
快捷访问