2012-02-28 41 views
1

我有如下操作的工作流:CUDA 2D阵列 - 在分配之前考虑音高?

  1. 负荷初始值
  2. 过程值的中间结果甲
  3. 方法A中间体结果乙
  4. 方法B中间体结果Ç
  5. 过程C和B至中间结果D和E
  6. 总和部分D至最终结果F

我所有中间结果的自然结构都是2D数组,我用cudaMallocPitch()分配它。

不幸的是,我的算法要求我一次将D,E,C和B保存在内存中,并且单独地,内存比B大4倍。由于我的处理中有另一个限制(迭代在存储器中的图形结构上),A或B的尺寸受D和E的最大尺寸限制,而D和E的最大尺寸又由初始值+ B的内存消耗+ C的内存消耗的内存使用量决定。这种依赖性是因为我正在从主机中“分页”到/从设备内存的中间结果部分(以容纳非常大的问题集),并且我无法开始第4步,直到完成了整个步骤1-3问题集。

一旦我有B中所有的问题集,我可以删除A.

我目前确定d + E的最大尺寸用下面的函数:

int gpuCalculateSimulPatterns(int lines, int patterns) { 
    // get free memory 
    size_t free_mem, total_mem; 
    int allowed_patterns; 
    cudaMemGetInfo(&free_mem, &total_mem); 
    allowed_patterns = (free_mem - (lines*sizeof(int)))/(lines*(sizeof(int)*2.5) + lines*sizeof(char)*1.5); 
    return min(patterns, allowed_patterns -(allowed_patterns % 32)); 
} 

它“作品” ,但仅仅是因为我高估了D或E(它们的尺寸和内存使用率相同)的大小25%,并且使B的预期大小加倍。即使如此,我仍遇到边缘情况,因为它的内存分配失败内存不足。我想更有效地利用卡上的内存保持对齐,因为我的内核对全局内存进行多次读写。

不,使用共享内存不是一个选项,因为我在多个块之间使用多个内核,并且块内的线程根本不相互作用。

我发现cudaMallocPitch()只返回已成功分配的内存使用的音调。有没有办法给司机一个二维内存分配请求,只是要求它分配的音调?

我想搭建一个试验/错误优化程序,但A,B,D和E之间维度的链接依赖性(CI计算先验,因为它没有分配点距线性)使得这个解决方案很糟糕,它需要重新计算每个问题集。

有没有人有更好的方法,可以让我确定适合任意数量的设备内存的适当大小的中间数据集?

编辑:

为中间体A的存储器被重复使用,我的边界的计算使得假设C + d + E + B >>初始+ A + B(它凭借这样的事实为真甲& B是相同尺寸的1字节的字符,而C,d,E是整数)和这样我只需要保证有一个为B + C + d + E.

我只使用足够的空间计算能力2.x卡可以用(Quadro 2000,Tesla C2075,GTX460)进行测试。

+0

是间距分配浪费了多少内存?你有多少例子?是否足够担心?如果是这样,你可能会考虑做线性内存分配,并将你的索引线性化。 (另外,从你的依赖链看起来中间A的内存可以被重用)。 – harrism 2012-02-29 01:20:11

+0

我使用音调线性内存的主要原因是利用了合并,并且因为一个进程中位置(X,Y)的结果在另一个进程中再次用于(X,Y)。 音高分配浪费了可变数量的内存。我认为它保持到最接近的512字节,至少对于我正在测试的其中一张卡片。 – 2012-02-29 15:06:43

+0

你不需要间距线性来合并。您只需要(在Fermi上)查找经线中每个负载的地址,并将其放入单个对齐的128字节高速缓存行中。你永远不会跨进程获得缓存重用,所以我不确定为什么这是相关的。也许我误解了。如果经线内的局部性模式是2D(并且具有良好的空间局部性),那么使用纹理访问会更好,因为硬件纹理缓存和纹理内存布局针对2D局部性进行了优化,而间距线性仍然是,线性的。 – harrism 2012-02-29 23:16:19

回答

1

节距以字节为单位的计算是这样的:

pitch_size = DIV_UP(width_in_bytes, prop.textureAlignment); 

凡DIV_UP舍入的第一个参数到第二个参数的下一个较高的倍数。

您可以通过致电cudaGetDeviceProperties()查询prop.textureAlignment

+0

虽然我在cuda文档中找不到DIV_UP,但它与宏的效果相同: '#define DIV_UP(x,y)(((x%y)> 0)?((x/y)+((x%y)> 0))* y:x)' – 2012-03-07 18:28:25

+0

嗯,我不会使用模数,这是一个相当昂贵的方式来做到这一点。我将使用'#define DIV_UP(x,y)((y)*(((x)+(y)-1)/(y)))' – harrism 2012-03-08 00:23:43

+0

我没有在GPU代码中使用宏,代码谨慎使用。谢谢你的提醒。 – 2012-03-09 01:21:39

0

不应该以字节为单位的间距为:

pitch_size = DIV_UP(width_in_bytes, prop.textureAlignment)*prop.textureAlignment; 

代替:

pitch_size = DIV_UP(width_in_bytes, prop.textureAlignment);