gpt4 book ai didi

cuda寄存器压力

转载 作者:行者123 更新时间:2023-12-02 01:21:13 32 4
gpt4 key购买 nike

我有一个内核进行线性最小二乘拟合。事实证明,线程使用了太多寄存器,因此占用率很低。这是内核,

__global__
void strainAxialKernel(
float* d_dis,
float* d_str
){
int i = threadIdx.x;
float a = 0;
float c = 0;
float e = 0;
float f = 0;
int shift = (int)((float)(i*NEIGHBOURS)/(float)WINDOW_PER_LINE);
int j;
__shared__ float dis[WINDOW_PER_LINE];
__shared__ float str[WINDOW_PER_LINE];

// fetch data from global memory
dis[i] = d_dis[blockIdx.x*WINDOW_PER_LINE+i];
__syncthreads();

// least square fit
for (j=-shift; j<NEIGHBOURS-shift; j++)
{
a += j;
c += j*j;
e += dis[i+j];
f += (float(j))*dis[i+j];
}
str[i] = AMP*(a*e-NEIGHBOURS*f)/(a*a-NEIGHBOURS*c)/(float)BLOCK_SPACING;

// compensate attenuation
if (COMPEN_EXP>0 && COMPEN_BASE>0)
{
str[i]
= (float)(str[i]*pow((float)i/(float)COMPEN_BASE+1.0f,COMPEN_EXP));
}

// write back to global memory
if (!SIGN_PRESERVE && str[i]<0)
{
d_str[blockIdx.x*WINDOW_PER_LINE+i] = -str[i];
}
else
{
d_str[blockIdx.x*WINDOW_PER_LINE+i] = str[i];
}
}

我有 32x404 block ,每个 block 中有 96 个线程。在 GTS 250 上,SM 应能够处理 8 个 block 。然而,视觉分析器显示每个线程有 11 个寄存器,因此占用率为 0.625(每个 SM 5 个 block )。 BTW,每个 block 使用的共享内存是792 B,所以寄存器是问题所在。演出并不是世界末日。我只是好奇是否有办法可以解决这个问题。谢谢。

最佳答案

快速但有限的寄存器/共享内存与缓慢但大的全局内存之间始终需要进行权衡。没有办法“绕过”这种权衡。如果您通过使用全局内存来减少寄存器的使用,您应该会获得更高的占用率,但内存访问速度会更慢。

也就是说,这里有一些使用更少寄存器的想法:

  1. 移位可以预先计算并存储在常量内存中吗?那么每个线程只需要查找shift[i]即可。
  2. a 和 c 必须是 float 吗?
  3. 或者,a 和 c 可以从循环中删除并计算一次吗?从而完全删除?

a 被计算为一个简单的算术序列,因此减少它......(类似这样)

a = ((NEIGHBORS-shift) - (-shift) + 1) * ((NEIGHBORS-shift) + (-shift)) / 2

a = (NEIGHBORS + 1) * ((NEIGHBORS - 2*shift)) / 2

因此,请执行以下操作(您可能可以进一步减少这些表达式):

str[i] = AMP*((NEIGHBORS + 1) * ((NEIGHBORS - 2*shift)) / 2*e-NEIGHBOURS*f)
str[i] /= ((NEIGHBORS + 1) * ((NEIGHBORS - 2*shift)) / 2*(NEIGHBORS + 1) * ((NEIGHBORS - 2*shift)) / 2-NEIGHBOURS*c)
str[i] /= (float)BLOCK_SPACING;

关于cuda寄存器压力,我们在Stack Overflow上找到一个类似的问题: https://stackoverflow.com/questions/4200230/

32 4 0
Copyright 2021 - 2024 cfsdn All Rights Reserved 蜀ICP备2022000587号
广告合作:1813099741@qq.com 6ren.com