0

私は63個のレジスタ/スレッドを使用しているので(最大32768)、約520個のスレッドを使用できます。この例では現在512個のスレッドを使用しています。

(並列処理は、グローバルcomputeEHfields関数関数内の関数 "computeEvec"にあります。)問題は次のとおりです。

1)以下のmemチェックエラー。

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

numPointsRp> 2000を使用すると、「メモリ不足」と表示されます

これで、実際に使用できるコードがいくつかできました。それをコンパイルして、何が起こるかを見てみましょう。CUDA 4.2ツールチェーンを使用RowRsSize=2000RowRpSize=200てコンパイルすると、次のようになります。

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には7SMがあることを示しています。これにより、pyCUDAを使用して入力と出力に1バイトのメモリを割り当てる前に、カーネルを実行するための合計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 に答える
1

R=1000 を使用し、

block=R/2,1,1 および grid=1,1 すべて OK

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 に答える