一种全局内存访问的方法和设备的制造方法_3

文档序号:9471303阅读:来源:国知局
,与上一次half-warp线程束所访问每个元素对应的索引值中的最大索引值 按照第一公式得到的第二列号进行比较,若第一列号与第二列号相等,且将最小索引值和 最大索引值满足第三公式:
[0113]
[0114] maxlndex表示最大索引值,minlndex表示最小索引值,则确定访问模式为依次访 问,也就是说,此次half-warp线程束所访问的最小索引值的元素,和上一次half-warp线 程束所访问的最大索引值的元素属于同一列,且此次half-warp线程束所访问的最小索引 值的元素,和上一次half-warp线程束所访问的最大索引值的元素位于相邻的两行,那么 就确定访问的是连续的子数据块。
[0115] 若第一列号与第二列号相差为1,且最大索引值按照第二公式得到的行号为〇,〇 代表第一行,最小索引值按照第二公式得到的行号为M-l,M-I代表最后一行,则确定访问 模式为依次访问,为连续的子数据块。
[0116] 204、在访问数据块的同时对数据块进行转置,得到转置数据块并对转置数据块进 行存储。
[0117] 具体而言,在确定了此次访问为按列依次访问后,则在此次访问的同时,对该数 据块Data进行转置,并将该数据块的标志位flag更新为第二标识true,以表示该数据块 Data存在转置数据块Data'。
[0118] 其中,对数据块Data进行转置,是通过将当前half-warp线程束访问的元素通过 局部内存进行转置,并存至新的数据空间中实现的。示例性的,先获取该数据块Data的空 间大小,在全局内存中分配同等大小的数据空间,用来存储转置后的新数据块Data',同时 分配局部内存block用于存储待转置的元素,而后在当前half-warp线程束访问Data元素 的同时将访问的元素进行转置,这里是通过每访问一个元素,对该元素进行转置实现的,并 将转置后的元素存储在局部内存中,待此次访问并转置完成后,将转置后的元素形成转置 数据块写回全局内存分配的同等大小的数据空间。也即此次访问还是访问的原数据块,为 非合合并访问,形成转置数据块,是为了方便再次将要按列访问原数据块时,直接访问其转 置数据块即可,也即下一次访问就会转化为合并访问。指的是CPU中内存模型的其中一种 当事件过程被触发时,局部内存便会分配内存空间给待转置数据块。
[0119] 其中,局部内存的大小可以为:
[0120] Block_dim*(Block_dim+l)*sizeof(type of Data)
[0121] Block_dim表示当前half-warp线程束的线程个数,sizeof(typeofData)表示 数据块中的一个元素的存储空间大小。这里的Block_dim+l之所以要加I是为了防止局 部内存出现存储冲突(bankconflict)的情况出现。具体而言,bank是指局部内存被划分 为大小相等,能被同时访问的存储器模块,不同的存储器模块可以互不干扰同时工作,但当 half-warp请求访问的多个地址位于同一bank时,由于存储器模块在一个时刻无法响应多 个请求,因此这些请求就必须被串行的完成,会出现bankconflict情况。Block-dim+1之 后可以保证half-warp请求访问的多个地址位于不同bank。
[0122] 上述转置过程中,根据GPU的内置编程模型,先将Data中的数据存放到block中 可以通过下列语言实现:
[0123] xlndex = blockldx. x*Block_dim+threadIdx. x;
[0124] ylndex = blockldx. y*Block_dim+threadIdx. y;
[0125]Index=yIndex*N+xIndex;
[0126]block[threadldx.y][threadldx.x] =Data[index];
[0127] 再将转置后的矩阵写回在全局内存中分配好的Data'中可以通过以下语言实现:
[0128] xlndex = blockldx. y*Block_dim+threadIdx. x;
[0129] ylndex = blockldx. x*Block_dim+threadIdx. y;
[0130]Data' [yIndex*M+xIndex] = block[threadldx. x][threadldx. y]〇
[0131] 205、判断访问是否结束,若未结束,则进入步骤202;若结束,则进入步骤211。
[0132] 在确定了数据块没有发生转置时,在访问数据块时,如果发生访问模式是按列访 问的情况,要在访问数据元素的同时进行转置,在转置时同时判断转置是否结束。
[0133] 这里可以根据当前half-warp线程束访问的子数据块中的每个元素的索引值中 的最大值,判断转置是否结束,若满足maxlndex=M*N_1,则确定此次访问结束,maxlndex 表示当前half-warp线程束访问的每个元素的索引值中的最大值。
[0134] 206、访问未进行转置处理之前的数据块。
[0135] 这里的访问转置处理之前的数据块,可以是由于前述步骤203判定了此次访问不 是按列访问,或者是由于前述步骤204判定了此次访问不是依次访问,都要访问Data数据 块中的元素,并将Data的标志位flag设为第一标识false,标识该数据块未进行转置。
[0136] 207、判断访问模式是否为按列依次访问,若为按列依次访问,则进入步骤208;若 不为按列依次访问,则进入步骤209。
[0137] 当确定了数据块Data的标志位为true后时,说明该数据块Data存在转置数据块 Data',这时,再判断当前访问是否为按列依次访问,这里的按列依次访问的实现方式与步 骤203和步骤204类似,不再赘述。
[0138] 208、访问转置数据块,而后进入步骤210。
[0139]如果当前访问是按列依次访问,则访问Data'中的数据。具体可以是:根据当 前half-warp线程束所获得的Data数据块中的元素的索引值index获得对应的Data' 数据块中的相应元素的索引值index',并访问Data[index'] :index' = (int) (index/ N) + (index%N)*M〇
[0140] 209、访问转置之前的数据块,而后进入步骤210。
[0141] 如果当前half-warp线程束不是按列依次访问,就访问转置之前的Data中的元 素,这里包括不是按列访问,或者是按列但不是依次访问的情况。
[0142] 210、判断此次访问是否结束,若未结束,则进入步骤207;若结束,则进入步骤 211。
[0143] 这里当前half-warp线程束访问当前的元素完成后,都要判断访问是否结束,判 断的依据是根据记录的当前half-warp线程束访问Data元素的最大索引值maxdex,是否 满足maxlndex=M*N-1,如果满足,则访问结束,如果不满足,则继续访问,进入步骤207。
[0144] 211、结束。
[0145] 这样一来,对于Data数据块以行优先存储的情况,当全局内存只读数据块被改变 存储布局时,无论此后以何种模式(按列、按行、乱序)访问该Data数据块,只要判断其为 按列访问,就直接访问其对应的转置后的数据块,避免了非合并访问的情况,提升了存储器 的访问带宽。
[0146] 需要说明的是,本发明是针对按列依次访问GPGPU内存模型的全局存储器而提出 的实施方案,对于全局存储器只读单元可能还有其他的访问方式,如斜对角访问等等,都可 以应用本发明的实施思维来解决其它的访问方式对应的问题。
[0147] 本发明实施例提供的全局内存访问的方法,在访问只读全局内存数据块时,根据 数据块的标志位判断数据块是否已进行转置;若未进行转置,则判断访问模式是否为按列 依次访问,若访问模式为按列依次访问,则在访问数据块的同时对数据块进行转置,得到转 置数据块并对转置数据块进行存储;若已进行转置,则判断访问模式是否为按列依次访问, 若访问模式为按列依次访问,则访问转置数据块,使得访问转置数据块时能够进行合并访 问,若访问模式不为按列依次访问,则访问转置之前的数据块,解决了现有技术中,访问全 局内存过程中,可能会出现按列访问时非合并访问的情况,而导致的全局内存访问带宽降 低的问题。
[0148] 实施例三
[0149] 本发明实施例提供一种设备01,如图3所示,包括:
[0150] 第一判断单元011,用于在访问只读全局内存数据块时,根据数据块的标志位判断 数据块是否已进行转置。
[0151] 第二判断单元012,用于若未进行转置,则判断访问模式是否为按列依次访问。
[0152] 转置单元013,用于若访问模式为按列依次访问,则在访问数据块的同时对数据块 进行转置,得到转置数据块并对转置数据块进行存储;
[0153] 第二判断单元012,还用于若已进行转置,则判断访问模式是否为按列依次访问。
[0154] 访问单元014,用于若访问模式为按列依次访问,则访问转置数据块,使得访问转 置数据块时能够进行合并访问,若访问模式不为按列依次访问,则访问转置之前的数据块。
[0155] 可选的,所提供的设备,还包括:插入单元015,用于在判断访问模式是否为按列 访问之前,在GPU编译平台的访问语句前插入桩代码,桩代码用于指示判断访问模式是否 为按列依次访问。
[0156] 可选的,第一判断单元011可以具体用于:
[0157] 若标志位为第一标识,则确定数据块未进行转置;
[0158] 若标志位为第二标识,则确定数据块已进行转置。
[0159] 可选的,第二判断单元012可以具体用于:
[0160] 判断访问模式是否为按列访问;
[0161] 若判断访问模式为按列访问,则再判断访问模式是否为依次访问。
[0162] 可选的,第二判断单元012可以具体用于:
[0163] 获取当前half-warp线程束访问数据块时所访问的每个元素的索引值,根据索引 值并按照第一公式获取每个元素对应的列号;
[0164] 若每个元素对应的列号相等,且相邻索引值之间相差为N,N表示数据块的列数, 则确定访问模式为按列访问;
[0165] 若每个兀素对应的列号中有两个列号相差为1,同时列号相等的相邻索引值相差 为N,其中的列号大者对应的每个元素按照第二公式得出的行值中最小值为0,列号小者对 应的每个元素按照第二公式得出的行值中最大值为M-I,M表示数据块的行数,则确定访问 模式为按列访问;
[0166] 其中,第一公式包括
,_ columnlndex表
当前第3页1 2 3 4 5 
网友询问留言 已有0条留言
  • 还没有人留言评论。精彩留言会获得点赞!
1