1

CUDA / 有关对缓存优化的问题

« 于: 三月 01, 2023, 05:54:35 pm »

问题1：
请问L1/L2 Cache上只有local memory和global memory吗？如果SM读写数据是否和CPU一样先访问L1，再考虑L2，最后在内存？
问题2：
如下图所示，L1 Cache的hit Rate仅51%，请问这方面应该如何优化呢？

2

CUDA / Re: 请教有关FP32和FP64相关问题

« 于: 二月 13, 2023, 10:44:45 pm »

很感激您的回复，已找到问题所在，此问题有显著改善，多谢

3

CUDA / 请教有关FP32和FP64相关问题

« 于: 二月 08, 2023, 04:44:20 pm »

为什么我在代码中 所有的计算部分已经全都是用的float类型（可能一部分是FP64单元去计算了），如Nsight Compute分析所示，为什么还提示我让我用FP32，如何让FP64不去参加运算？在这种情况下请问怎么进行优化呢？

4

CUDA / 求问：有关架构方面的问题

« 于: 十二月 11, 2022, 07:33:07 pm »

1：书上所说：Fermi架构可以在每个SM上同时处理48个线程束，请问这个48个线程束是如何计算得到的？与什么参数有关呢？
2：在Fermi架构中，每个SM中含有两个线程束调度器和两个指令调度调度单元，这是不是意味着每个时间周期内有两个warp即64个线程从active thread变成计算线程开始运算操作？可这样不是和问题一中的同时处理48个线程束矛盾了？

5

CUDA / Re: 求问几个有关优化方面的问题

« 于: 十月 13, 2022, 02:11:13 pm »

看上去你的SM和Memory没有充分忙碌，你的图综合说明了：
occupancy较低，理论occupancy才能到大约12个warps左右（看图估计的）/ SM, 也就是如果是8.6的卡才20%~30%，achieved occupancy更是大约10个warps/SM（~20%）。如果你认为你已经用launch bounds限制了，则可能你的launch bounds设定有误（或者你因为没有说明的苦衷而不能更激进的设定）。这里可以考虑尝试提升一下，看看是否有效果。

此外，可以尝试消除blocks内部的长尾效应（部分warps提前结束，部分warps拖后结束，导致提前结束的warps依然占用部分资源，从而achieved occupancy比theoretical的occupancy低一些。这里也可以考虑修一下。不过这个不是主要因素。

你的图看来主要矛盾点在访存，而不是SM使用率。这样综上尝试提升了occupancy后，看看能否让SM和MEM的忙碌程度提升一下（主要是访存），这样看看能否提速。

N卡只支持从每个线程的角度的1B, 2B, 4B, 8B, 16B粒度的访存的。你的"一次性"24B每个线程的写入，因为粒度和对齐的原因，可能会被拆分成3次8B的写入的，因此不是一次。从这个拆分的角度看，每个warp的这个写入，不是合并的（warp的32个8B写入，中间有很多16B的空洞）。不过因为有cache的存在，在8.6上这种写入其实还好，不太要紧。不过如果能提升，可以尝试修改下。例如改成32B每个线程的写入量，这样1个warp会拆分成两条无空洞的合并的16B写入，不一定有正面效果（因为写入量变大了），不过你可以试试看，特别是在有ECC的卡上，能减少空洞导致的读取--修改--写入的访存放大。（普通卡一般情况下可能无此放大问题）。

你说的因为分支，每个warp实际有效的计算量是13.50/32 = 42%，这的确是个问题。如果算法的确会导致warp内分支的话，根据每个分支体的代码量不同，可能的确没办法。也可能你1个warp内部，通过任务重排，例如将40%的分支1，60%的分支2，重新排列成前40%的线程计算分支1，后60%的线程计算分支2，可能会提升效率。如果你假设block有8个warps，256个线程：原本每个warp都是40%效率分支1 ---> 60%效率分支2。变成了前40%的线程（也就是warp0，1，2无分支的全效率执行分支1，和warp4，5，6，7全效率的执行分支2，和warp 3依然部分效率的两个分支都执行），这样的任务重拍可以有效的减少divergent branch, 让数字非常好看。但是不一定总是有正面效果。（因为太小的分支不够block里头任务重排的的成本）。

此外，目前等待MEM的延迟依然是主要矛盾（Long SB，倒数第二个图），看看能重写代码的结果或者算法的逻辑，减少这一点。

十分感谢，看了您的建议受益匪浅

6

CUDA / 求问几个有关优化方面的问题

« 于: 十月 12, 2022, 05:04:09 pm »

问题1（图一所示）：
 如Nsight Compute提示所示，如何有效的提高计算吞吐量和内存带宽？
问腿2 ：
在这个项目中，我一个block内有128个线程，每个线程往全局内存中每次写一个ulonglong3（24字节），求问，这算是满足对齐的要求了吗？（每个线程写的地址都是连续的，相当于写了24*128字节）
问题3：（图二所示）
如Nsight Compute提示所示，如何提高活跃的线程束数量呢？我已经使用__launch_bounds__来限制寄存器的数量来提高了，请问还有其他方法来优化此问题吗？
问题4：（图三和图四）
如Nsight Compute提示所示，如何解决线程束发散导致的问题呢？（算法以及无法减少if分支）

感谢各位[名词6]的指导

7

CUDA / 小白求问，为什么在vs中(release)的结果与生成的exe文件在cmd中运行的结果相差如此大？

« 于: 十月 10, 2022, 10:50:27 am »

两者的计算时间不应该相差不多吗，为什么差距会这么大呢？

8

CUDA / 小白求问，有关本地内存和共享内存与全局内存的问题

« 于: 九月 16, 2022, 10:36:53 am »

问题描述: kernel核函数内，我将数据保存在本地内存或者共享内存中，请问cuda中是否有类似与memcpy函数可以将本地内存和共享内存中的数据一次性搬运至全局内存中呢？还是只能for遍历，一个个数据搬运？

9

CUDA / 如何提高全局内存的写入速率？

« 于: 八月 16, 2022, 04:25:52 pm »

代码如下所示，d_m_rgbbuffer指向的是全局内存，该函数的作用是YUV转RGB 注释部分（*lmbmp++部分）就是RGB数据写入全局内存中，这一部分所花时间占整个时间的70%，（一个线程处理192个数据）请问应该如何进行优化？

程序代码: [选择]

__device__ void StoreBuffer(short * QtZzMCUBuffer, short dx, short dy, unsigned char* d_m_rgbbuffer)
{
short i, j;
unsigned char* lpbmp;
unsigned char R, G, B;
int y, u, v, rr, gg, bb;
//unsigned char RGB[63 * 3];//得到该8*8像素的RGB
for (i = 0; i <8; i++)
{
if ((dy + i) < d_m_height)
{
lpbmp = ((unsigned char*)d_m_rgbbuffer + (dy + i) * 3 * d_m_width + dx * 3);
for (j = 0; j <  8; j++)
{
if ((dx + j) < d_m_width)
{
y = QtZzMCUBuffer[i * 8 + j];
u = QtZzMCUBuffer[i * 8 + j + 64];
v = QtZzMCUBuffer[i * 8 + j + 128];
rr = ((y << 8) + 18 * u + 367 * v) >> 8;
gg = ((y << 8) - 159 * u - 220 * v) >> 8;
bb = ((y << 8) + 411 * u - 29 * v) >> 8;
R = (unsigned char)rr;
G = (unsigned char)gg;
B = (unsigned char)bb;
if (rr & 0xffffff00) if (rr > 255) R = 255; else if (rr < 0) R = 0;
if (gg & 0xffffff00) if (gg > 255) G = 255; else if (gg < 0) G = 0;
if (bb & 0xffffff00) if (bb > 255) B = 255; else if (bb < 0) B = 0;
//*lpbmp++ = B;
//*lpbmp++ = G;
//*lpbmp++ = R;
}
else  break;
}
}
else break;
}
}