driver.Context.synchronize()- 还有什么需要考虑的 -- -- 清理操作失败

Question

我在这里有 这个代码（由于答案而修改）。

信息

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

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

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

我的计算：

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

我的卡有限制 32768 个寄存器，所以 32768/5888=5 +一些 => 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 分钟或更长时间）。我应该使用它还是不使用它？

--------------UPDATED2----------------------------------- -----------

现在，如果不做任何事情，代码将无法再次运行！

代码截图：

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

Answer 1

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

1. PTX输出的46个寄存器不是你的内核使用的寄存器数量。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 block/SM （我的卡有限制 6）。

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

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

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

A。在代码开头添加一行

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

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

b. 如果你有太多的波浪，那么就必须减少发射配置。鉴于您通过 gridDim.x 和 blockDim.x 进行索引，您可以通过将这些维度作为参数传递给内核来完成此操作。这将要求您最少地更改索引代码。您还必须传递 blockIdx.x 偏移量。更改主机代码以连续启动多个内核。由于不应该存在冲突，因此您可以在多个流中启动它们，以便从每波结束时的重叠中受益。