CUDA - 出的内存(线程和块问题)--address超出范围(cuda — out of me

2019-07-31 05:58发布

我使用63寄存器/线程,所以(32768是最大值),我可以使用约520 threads.I AM在这个例子中使用现在512个线程。

(平行度是在函数“computeEvec”全球computeEHfields功能函数内。)的问题是:

1)下面的MEM检查错误。

2)当我使用numPointsRp> 2000年,告诉我“内存不足”,但(如果我没有做错了)我计算全局内存,它的确定。

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

我运行CUDA-MEMCHECK程序,这让我(只有当numPointsRs> numPointsRp):

=========大小为4的无效全局

=========在computeEHfields 0x00000428

=========通过在块螺纹(2,0,0)(0,0,0)

=========地址0x4001076e0越界

========= =========大小为4的无效全局

=========在computeEHfields 0x00000428

=========通过在块的线程(1,0,0)(0,0,0)

=========地址0x4001076e0越界

========= =========大小为4的无效全局

=========在computeEHfields 0x00000428

=========通过在块的线程(0,0,0)(0,0,0)

=========地址0x4001076e0越界

错误摘要:160级的错误

- - - - - -编辑 - - - - - - - - - - - - - -

此外,一些时间(如果我只使用线程而不是块(我还没有测试它的块)),如果例如我有numPointsRs = 1000和numPointsRp = 100,然后改变numPointsRp = 200,然后再次更改numPointsRp = 100我没有服用第一的成绩!

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


Rs=np.zeros((numPointsRs,3)).astype(np.float32)
for k in range (numPointsRs): 
    Rs[k]=[0,k,0]

Rp=np.zeros((numPointsRp,3)).astype(np.float32)
for k in range (numPointsRp): 
    Rp[k]=[1+k,0,0]


#---- 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>
#define RowRsSize %(numrs)d
#define RowRpSize %(numrp)d


typedef  pycuda::complex<float> cmplx;
extern "C"{


    __device__ void computeEvec(float Rs_mat[][3], int numPointsRs,   
         cmplx J[][3],
         cmplx M[][3],
         float *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 )
    {
        float Rs_mat[RowRsSize][3];
        float Rp_mat[RowRpSize][3];

        cmplx J[RowRsSize][3];
        cmplx M[RowRsSize][3];


    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;


    }

}
}

"""% { "numrs":numPointsRs, "numrp":numPointsRp},no_extern_c=1)


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))

print(" \n")


#----- 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()

-------------------- GPU型号---------------------------- --------------------

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:

当我使用numPointsRp> 2000年,告诉我“内存不足”

现在我们有一些实际的代码一起工作,让我们编译它,看看会发生什么。 使用RowRsSize=2000RowRpSize=200 ,并与CUDA 4.2工具链编译,我得到:

nvcc -arch=sm_21 -Xcompiler="-D RowRsSize=2000 -D RowRpSize=200" -Xptxas="-v" -c -I./ kivekset.cu 
ptxas info    : Compiling entry function '_Z15computeEHfieldsPfiS_iPN6pycuda7complexIfEES3_S2_S2_PA3_S2_S5_S3_' for 'sm_21'
ptxas info    : Function properties for _Z15computeEHfieldsPfiS_iPN6pycuda7complexIfEES3_S2_S2_PA3_S2_S5_S3_
    122432 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info    : Used 57 registers, 84 bytes cmem[0], 168 bytes cmem[2], 76 bytes cmem[16]

密钥号是57级的寄存器和堆栈每线程帧122432个字节。 占有率计算器建议的512个线程块最多将有每个SM 1块的,你的GPU有7个SM。 这给出了一个总的122432 * 512 * 7 = 438796288字节堆栈帧(本地存储器)来运行你的内核,你已经分配的内存字节的单一输入和输出使用pyCUDA为之前。 在有1GB内存一个GPU,这是不难想象的运行内存。 你的内核有一个巨大的本地内存占用。 开始考虑如何减少它。


正如我在评论中指出,这是绝对不清楚为什么每个线程需要在这个内核代码的输入数据的完整副本。 这导致了一个巨大的本地内存占用和似乎是绝对没有理由的代码应该以这种方式来写。 你可以,我怀疑,修改内核是这样的:

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

__global__  
void computeEHfields2(
        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_;
    cp3 * J = (cp3 *)J_;
    cp3 * M = (cp3 *)M_;

    int k=threadIdx.x+blockIdx.x*blockDim.x;
    while (k<numPointsRp)  
    {
        fp3 * Rp_mat = (fp3 *)(Rp_mat_+k);
        computeEvec2( Rs_mat, numPointsRs, J, M, *Rp_mat, kp, eta, E[k], H[k], All );
        k+=blockDim.x*gridDim.x;
    }
}

和主__device__函数调用来是这样的:

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

和消除线程本地存储的每个字节 ,而完全不改变计算代码的功能。



Answer 2:

使用R = 1000,然后

块= R / 2,1,1和网格= 1,1一切OK

如果我尝试R = 10000

块= R / 20,1,1和电网= 20,1,那么它告诉我“内存”

我不熟悉pycuda并没有读入你的代码太深。 然而,你有更多的块,更多的线程,所以它会

  • 本地内存(可能是内核堆栈,它的每个线程分配),

  • 共享存储器(每块中分配),或

  • 该系统会根据分配的全局内存gridgridDim

您可以减少堆栈大小调用

cudeDeviceSetLimit(cudaLimitStackSize, N));

(该代码是C运行时API,但pycuda相当于不应该太难找)。



文章来源: cuda — out of memory (threads and blocks issue) --Address is out of bounds
标签: cuda pycuda