5

我在这里有 这段代码(由于答案而修改)。

信息

32 字节堆栈帧,0 字节溢出存储,0 字节溢出加载
ptxas 信息:使用 46 个寄存器,120 字节 cmem[0],176 字节 cmem[2],76 字节 cmem[16]

我不知道还要考虑什么才能使其适用于点“numPointsRs”和“numPointsRp”的不同组合

例如,当我使用 Rs=10000 和 Rp=100000 运行代码时,block=(128,1,1),grid=(200,1) 很好。

我的计算:

46 个寄存器*128 线程=5888 个寄存器。

我的卡有限制 32768 个寄存器,所以 32768/5888=5 +some => 5 块/SM
(我的卡有限制 6)。

使用占用率计算器,我发现使用 128 个线程/块给了我 42% 并且在我的卡的限制范围内。

此外,每个 MP 的线程数为 640(限制为 1536)

现在,如果我尝试使用 Rs=100000 和 Rp=100000(对于相同的线程和块)它会给我标题中的消息,其中:

cuEventDestroy 失败:启动超时

cuModuleUnload 失败:启动超时

1)我不知道/理解还需要计算什么。

2)我不明白我们如何使用/查找块的数量。我可以看到,大多数情况下,有人放(threads-1+points)/threads,但这仍然不起作用。

- - - - - - - 更新 - - - - - - - - - - - - - - - - - - ------------

使用 driver.Context.synchronize() 后,代码适用于许多点(1000000)!

但是,这个添加对代码有什么影响?(对于许多点,屏幕冻结 1 分钟或更长时间)。我应该使用它还是不使用它?

--------------已更新2----------------------------------- ------------

现在,代码在不做任何事情的情况下无法再次工作!

代码快照:

import pycuda.gpuarray as gpuarray
import pycuda.autoinit
from pycuda.compiler import SourceModule
import numpy as np
import cmath
import pycuda.driver as drv
import pycuda.tools as t






#---- Initialization and passing(allocate memory and transfer data) to GPU -------------------------
Rs_gpu=gpuarray.to_gpu(Rs)
Rp_gpu=gpuarray.to_gpu(Rp)

J_gpu=gpuarray.to_gpu(np.ones((numPointsRs,3)).astype(np.complex64))
M_gpu=gpuarray.to_gpu(np.ones((numPointsRs,3)).astype(np.complex64))

Evec_gpu=gpuarray.to_gpu(np.zeros((numPointsRp,3)).astype(np.complex64))
Hvec_gpu=gpuarray.to_gpu(np.zeros((numPointsRp,3)).astype(np.complex64))
All_gpu=gpuarray.to_gpu(np.ones(numPointsRp).astype(np.complex64))

#-----------------------------------------------------------------------------------    
mod =SourceModule("""
#include <pycuda-complex.hpp>
#include <cmath>
#include <vector>

typedef  pycuda::complex<float> cmplx;
typedef float fp3[3];
typedef cmplx cp3[3];

__device__ __constant__ float Pi;

extern "C"{  


    __device__ void computeEvec(fp3 Rs_mat[], int numPointsRs,   
         cp3 J[],
         cp3 M[],
         fp3 Rp,
         cmplx kp, 
         cmplx eta,
         cmplx *Evec,
         cmplx *Hvec, cmplx *All)

{

            while (c<numPointsRs){

        ...                      
                c++;

                }        
        }


__global__  void computeEHfields(float *Rs_mat_, int numPointsRs,     
        float *Rp_mat_, int numPointsRp,     
    cmplx *J_,
    cmplx *M_,
    cmplx  kp, 
    cmplx  eta,
    cmplx E[][3],
    cmplx H[][3], cmplx *All )
    {

        fp3 * Rs_mat=(fp3 *)Rs_mat_;
        fp3 * Rp_mat=(fp3 *)Rp_mat_;
        cp3 * J=(cp3 *)J_;
        cp3 * M=(cp3 *)M_;


    int k=threadIdx.x+blockIdx.x*blockDim.x;

      while (k<numPointsRp)  
     {

        computeEvec( Rs_mat, numPointsRs,  J, M, Rp_mat[k], kp, eta, E[k], H[k], All );
        k+=blockDim.x*gridDim.x;

    }

}
}

""" ,no_extern_c=1,options=['--ptxas-options=-v'])


#call the function(kernel)
func = mod.get_function("computeEHfields")

func(Rs_gpu,np.int32(numPointsRs),Rp_gpu,np.int32(numPointsRp),J_gpu, M_gpu, np.complex64(kp), np.complex64(eta),Evec_gpu,Hvec_gpu, All_gpu, block=(128,1,1),grid=(200,1))


#----- get data back from GPU-----
Rs=Rs_gpu.get()
Rp=Rp_gpu.get()
J=J_gpu.get()
M=M_gpu.get()
Evec=Evec_gpu.get()
Hvec=Hvec_gpu.get()
All=All_gpu.get()

我的卡:

Device 0: "GeForce GTX 560"
  CUDA Driver Version / Runtime Version          4.20 / 4.10
  CUDA Capability Major/Minor version number:    2.1
  Total amount of global memory:                 1024 MBytes (1073283072 bytes)
  ( 0) Multiprocessors x (48) CUDA Cores/MP:     0 CUDA Cores   //CUDA Cores    336 => 7 MP and 48 Cores/MP
4

3 回答 3

3

“启动超时”似乎表明内核运行时间过长并被看门狗定时器杀死。这可能发生在也用于图形输出(例如图形桌面)的 GPU 上,其中看门狗计时器的任务是防止桌面锁定超过几秒钟。最好我记得看门狗的时间限制大约是 5 秒左右。

在任何给定时刻,GPU 都可以运行图形或 CUDA,因此在运行 GUI 时需要看门狗定时器,以防止 GUI 长时间锁定,从而导致机器无法通过 GUI 操作。

如果可能,请避免将此 GPU 用于桌面和/或其他图形(例如,如果您在 Linux 上,请不要运行 X)。如果在没有图形的情况下运行不是一个选项,为了减少内核执行时间以避免遇到看门狗定时器内核终止,您将不得不在每次内核启动时做更少的工作,优化代码以便内核在相同的工作量下运行得更快,或者部署更快的 GPU。

于 2012-09-14T00:56:31.933 回答
3

有很多问题需要你处理。@njuffa 提供的答案 1 是最好的通用解决方案。我将根据您提供的有限数据提供更多反馈。

  1. 46 个寄存器的 PTX 输出不是内核使用的寄存器数。PTX 是一种中间表示。离线或 JIT 编译器会将其转换为设备代码。设备代码可以使用更多或更少的寄存器。Nsight Visual Studio Edition、Visual Profiler 和 CUDA 命令行分析器都可以为您提供正确的寄存器计数。

  2. 占用计算不仅仅是RegistersPerSM / RegistersPerThread。寄存器是基于粒度分配的。对于 CC 2.1,粒度是每个线程每个线程 4 个寄存器(128 个寄存器)。2.x 设备实际上可以以 2 个寄存器粒度进行分配,但这可能会导致稍后在内核中出现碎片。

  3. 在您的入住计算中,您声明

我的卡有限制 32768 个寄存器,所以 32768/5888=5 +some => 5 块/SM(我的卡有限制 6)。

我不确定 6 是什么意思。您的设备有 7 个 SM。对于 2.x 设备,每个 SM 的最大块数是每个 SM 8 个块。

  1. 您提供的代码数量不足。如果您提供代码片段,请提供所有输入的大小、每个循环将执行的次数以及每个函数的操作描述。查看代码,您可能每个线程执行了太多循环。在不知道外循环的数量级的情况下,我们只能猜测。

  2. 鉴于启动超时,您可能应该按如下方式进行调试:

一种。在代码开头添加一行

if (blockIdx.x > 0) { return; }

运行前面提到的分析器之一中的确切代码,以估计单个块的持续时间。使用分析器提供的启动信息:每个线程的寄存器、共享内存……使用分析器中的占用计算器或 xls 来确定可以同时运行的最大块数。例如,如果理论块占用率是每个 SM 3 个块,并且 SM 的数量是 7 个,那么您一次可以运行 21 个块,对您来说发射 9 个波。注意:这假设每个线程的工作量相等。更改提前退出代码以允许 1 波(21 个块)。如果此启动超时,那么您需要减少每个线程的工作量。如果通过了,那么计算你有多少波并估计你什么时候超时(在 Windows 上为 2 秒,在 Linux 上为 ?)。

湾。如果你有太多的波浪然后减少你必须减少启动配置。鉴于您按 gridDim.x 和 blockDim.x 进行索引,您可以通过将这些维度作为参数传递给内核来做到这一点。这将要求您对您的索引代码进行最低限度的更改。您还必须传递一个 blockIdx.x 偏移量。更改您的主机代码以背靠背启动多个内核。由于不应该有冲突,您可以在多个流中启动这些,以从每个波结束时的重叠中受益。

于 2012-09-14T03:33:38.057 回答
0

要为@njuffa 的答案提供更多输入,在 Windows 系统中,您可以按照以下步骤增加启动超时或TDR (超时检测和恢复):

1:打开Nsight Monitor中的选项

在此处输入图像描述

2:为WDDM TDR 延迟设置适当的值

在此处输入图像描述

注意:如果此值很小,您可能会收到超时错误,并且对于更高的值,您的屏幕将保持冻结状态,直到内核完成它的工作。

来源

于 2019-02-18T15:26:16.003 回答