1

我假设,可以并行运行的线程/块的数量是有限的。我的意思是,如果我有太多线程/块,其中一些将在某些处理单元上按顺序执行。我需要构建下一个示例。比方说,我有一些kernel<<<B, N>>>(). 的执行时间kernel<<<1,1>>>()等于 t0。

第一个任务是找到 B 和 N 的最大值,即执行时间kernel<<<B, N>>>()t ~ t0。然后我想要kernel<<<B, 2*N>>>()(或kernel<<<2*B, N>>>())t1 ~ 2*t 的执行时间。

我有带有 448 个 CUDA 内核(14 个 SM)的 Tesla C2075,并希望建立一个占用率为 1 的示例。

这是否可能,如果是,内核函数应该是什么样子,可能是一些例子?

4

1 回答 1

0

在 CUDA 中,线程被安排在单个 SM 上作为 warp。每个warp最多可以包含32个线程。调度程序将尝试在SM内部以并行方式执行warp。如果特定warp的数据尚未准备好,它由调度程序持有直到可用。现在您的问题很重要,我相信可以使用cudaEvent_t(用于测量内核的执行时间)来做您想要实现的目标。

启动配置kernel<<<B,Tnum>>>(arg1...argn);完全取决于您可以在算法中利用多少并行度。此外,线程数是您必须根据启动内核获得的最佳执行时间来决定的。

在许多情况下,使用线程启动多个块128/256就足以实现最佳加速。举个例子,假设我们想将两个大小数组的单个元素添加1024到第三个数组中,具有 1 个块的内核函数看起来像

__global__ void kadd(int *c,int *a,int *b)
{
  unsigned int tid = threadIdx.x;//Since only one block of 1024 threads suffices
  if(tid < MAXNUM)  //MAXNUM = 1024
    c[tid] = a[tid]+ b[tid];
}

启动配置将是

kadd<<<1,1024>>>(c,a,b);

然而,这只会在你的 GPU 的一个 SM 上执行一个块,这意味着你没有充分利用 GPU 资源。为了从你的 GPU 中获得更多,你可以做的是你可以使用多个块和线程。内核看起来像

__global__ void kadd(int *c,int *a,int *b)
{
  unsigned int tid = blockIDx.x * blockDim.x + threadIdx.x;//Since multiple blocks are used
  if(tid < MAXNUM)  //MAXNUM = 1024
    c[tid] = a[tid]+ b[tid];
}

并且相应的启动配置将是

kadd<<<8,128>>>(c,a,b);

这将启动每个线程8块。128您可以根据您的算法要求使用此启动配置。2D您可以通过启动或3D网格进一步探索这些启动配置,以充分利用您的 GPU。

因此,对内核进行计时将为您提供最适合您要求的配置。这也将根据共享内存的使用、全局内存的合并访问和其他因素而改变。最后,我想提一下,NVIDIA 提供了一个占用率计算器,您可以使用它来找到块和线程的最佳组合,以实现更高的占用率。

于 2012-10-07T03:30:45.430 回答