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 / 请教有关FP32和FP64相关问题

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

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

3

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

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

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

4

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]的指导

5

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

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

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

6

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

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

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

7

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;
}
}