首页 文章

CUDA . 示例表示性能如何取决于线程数

提问于
浏览
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 .

我有特斯拉C2075与448 CUDA核心(14 SM),并希望 Build 一个占用1的例子 .

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

1 回答

  • 0

    在CUDA中,线程被安排在单独的SM 's as warps. Each warp can contain at max 32 threads.The scheduler will try to execute the warps in parallel fashion inside the SM'上 . 如果特定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);
    

    这将分别启动 8128 个线程 . 您可以根据算法要求使用此启动配置 . 您可以通过启动 2D3D 网格来进一步探索这些启动配置,以充分利用您的GPU .

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

相关问题