0

我正在使用 63 个寄存器/线程，所以（最多 32768 个）我可以使用大约 520 个线程。在这个例子中我现在使用 512 个线程。

（并行性在全局 computeEHfields 函数函数内的函数“computeEvec”中。）问题是：

1）下面的内存检查错误。

2）当我使用 numPointsRp>2000 时，它显示“内存不足”，但是（如果我没有做错）我计算全局内存并且没关系。

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

我用 cuda-memcheck 运行程序，它给了我（仅当 numPointsRs>numPointsRp 时）：

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

========= 在 computeEHfields 中的 0x00000428

========= 块 (0,0,0) 中的线程 (2,0,0)

========= 地址 0x4001076e0 超出范围

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

========= 在 computeEHfields 中的 0x00000428

========= 块 (0,0,0) 中的线程 (1,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
4

2 回答 2

1

使用 R=1000 然后

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

如果我尝试 R=10000 并且

block=R/20,1,1 和 grid=20,1 ，然后它显示“内存不足”

我对 pycuda 不熟悉，也没有深入阅读您的代码。但是你有更多的块和更多的线程，所以它会

  • 本地内存（可能是内核的堆栈，它是按线程分配的），

  • 共享内存（按块分配），或

  • 基于grid或分配的全局内存gridDim。

您可以减少堆栈大小调用

cudeDeviceSetLimit(cudaLimitStackSize, N));

（代码用于 C 运行时 API，但 pycuda 等价物不应该太难找到）。

于 2012-08-31T17:14:01.190 回答
1

当我使用 numPointsRp>2000 时，它显示“内存不足”

现在我们有一些真实的代码可以使用，让我们编译它看看会发生什么。使用RowRsSize=2000和RowRpSize=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。在您使用 pyCUDA 为输入和输出分配一个字节的内存之前，这总共提供了 122432 * 512 * 7 = 438796288 字节的堆栈帧（本地内存）来运行您的内核。在具有 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)
{
 ....
}

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

于 2012-09-01T20:03:47.423 回答