小编Ade*_*ick的帖子

根据"CUDA C编程指南",只有在多处理器常量高速缓存被命中时,常量内存访问才会受益(见第5.3.2.4节)¹.否则,对于半翘曲,可能存在比合并的全局存储器读取更多的存储器请求.那么为什么常量内存大小限制为64 KB？

还有一个问题是为了不要问两次.据我所知,在Fermi架构中,纹理缓存与L2缓存相结合.纹理使用是否仍然有意义或全局内存读取是否以相同的方式缓存？

¹ 恒定存储器(第5.3.2.4节)

常量存储空间驻留在设备存储器中,并缓存在F.3.1和F.4.1节中提到的常量高速缓存中.

对于计算能力为1.x的设备,对于warp的常量内存请求首先被分成两个请求,每个半warp一个,独立发出.

然后,请求被分成多个单独的请求,因为初始请求中存在不同的内存地址,吞吐量减少的因子等于单​​独请求的数量.

然后,在高速缓存命中的情况下,或者在设备存储器的吞吐量下,以恒定高速缓存的吞吐量来服务所得到的请求.

我致力于优化算法,因此性能非常重要.与VS 2008相比,在VS 2010中编译时,该算法的速度提高了约8倍.谷歌搜索显示这不是我的错(参见例如/sf/answers/389212911/).问题是最终项目必须在VS 2008下构建.

我倾向于在VS 2010中将我的算法构建为DLL,然后将其链接到主项目.是否可以在VS 2008下使用VC++ 2010运行时库和我的DLL？如果是这样,那么最痛苦的方法是什么？还有其他想法吗？谢谢.

我实现了一个简单的内核,它是某种卷积.我在NVIDIA GT 240上进行了测量.在CUDA上写入时需要70 ms,在OpenCL上写入时需要100 ms.好吧,我想,NVIDIA编译器更适合CUDA(或者我做错了).我需要在AMD GPU上运行它,所以我迁移到了AMD APP SDK.完全相同的内核代码.

我做了两次测试,他们的结果对我来说令人沮丧:HD 6670为200 ms,HD 5850为70 ms(与GT 240 + CUDA同时).而且我对这种奇怪行为的原因很感兴趣.

所有项目都是使用NVIDIA和AMD的示例项目中的设置在VS2010上构建的.

请不要将我的帖子视为NVIDIA广告.我很明白HD 5850比GT 240更强大.我唯一想知道的是为什么会出现这种差异以及如何解决问题.

更新.下面是内核代码,它在基础代码中查找6个大小相同的模板图像.基本图像的每个像素都被视为其中一个模板的可能来源,并由单独的线程处理.内核比较基本图像和模板1的每个像素的R,G,B值,并且如果至少一个差异超过diff参数,则对应的像素被计数为不匹配.如果不匹配像素的数量小于maxNonmatchQt相应的模板被击中.

__constant int tOffset = 8196; // one template size in memory (in bytes)
__kernel void matchImage6( __global unsigned char* image, // pointer to the base image
            int imgWidth, // base image width
            int imgHeight, // base image height
            int imgPitch, // base image pitch (in bytes)
            int imgBpp, // base image bytes (!) per pixel …

我将在CUDA上并行化一个局部搜索算法来解决一些优化问题.问题非常严重,因此实际可解决的问题的规模非常小.我担心的是计划在一个内核中运行的线程数量不足以在GPU上获得任何加速(即使假设所有线程都已合并,没有银行冲突,非分支等).假设为100个线程启动了一个内核.期望使用GPU获得任何利润是否合理？如果线程数是1000怎么办？分析案例还需要哪些其他信息？

这是我在GT 440上的内核的Compute Visual Profiler的输出:

• 内核详细信息:网格大小:[100 1 1],块大小:[256 1 1]
• 寄存器比率:0.84375(27648/32768)[每个线程35个寄存器]
• 共享内存比率:0.336914(16560/49152)[每块5520字节]
• 每SM的活动块数:3(每个SM的最大活动块数:8)
• 每个SM的活动线程数:768(每个SM的最大活动线程数:1536)
• 潜在占用率:0.5(24/48)
• 占用限制因素:登记

请注意标记为粗体的子弹.内核执行时间是121195 us.

我通过将一些局部变量移动到共享内存来减少每个线程的一些寄存器.Compute Visual Profiler输出变为:

• 内核详细信息:网格大小:[100 1 1],块大小:[256 1 1]
• 寄存器比率:1(32768/32768)[每个线程30个寄存器]
• 共享内存比率:0.451823(22208/49152)[每块5552字节]
• 每SM的活动块数:4(每个SM的最大活动块数:8)
• 每个SM的活动线程数:1024(每个SM的最大活动线程数:1536)
• 潜在入住率:0.666667(32/48)
• 占用限制因素:登记

因此,现在4块在单个SM上与3先前版本中的块同时执行.但是,执行时间115756 us几乎相同!为什么？是不是完全独立的块在不同的CUDA核心上执行？

我在Haskell中创建了一个数据类型:

type Name = String
data ModelNode = NodeAttribute Name | 
                 NodeRelation (Name,Name) |
                 NodeConstraint Name |
                 NodeKPI Name

我需要这种类型的实例Ord.我想出的实现是:

instance Ord ModelNodeKind where
    compare (NodeAttribute n)  (NodeAttribute n') = compare n n'
    compare (NodeAttribute _)  _                  = LT
    compare _                  (NodeAttribute _)  = GT
    compare (NodeRelation n)   (NodeRelation n')  = compare n n'
    compare (NodeRelation _)   _                  = LT
    compare _                  (NodeRelation _)   = GT
    compare (NodeConstraint n) (NodeConstraint n')= compare n n'
    compare (NodeConstraint _) _ …

我有数组A[0...N]的double和阵列B[0...N]的int.每个B[i]变化都有[0...P].我只需要计算数组C[0...P]:

C[j] = SUM( A[i] : B[i] = j)

据我所知,我不能使用N具有atomicAdd()函数的线程,因为它不支持double.使用P线程的直接实现非常不同.有没有更好的办法？

我们假设我在GHCi中执行以下代码:

[unsafePerformIO (randomRIO (0,1)) | _ <- [1..10]].

[1,1,0,0,1,0,0,0,0,1]每次执行此操作时,结果可能看起来都不同.现在假设我定义了以下函数:

f :: Int
f = unsafePerformIO (randomRIO (0,1))

并试着打电话:

[f | _ <- [1..10]].

结果总是如此[0,0,0,0,0,0,0,0,0,0].如果我只是打一个电话f,它总是返回0.暂且说我不应该把IO扔掉,我做错了什么？