2回目の反復クラッシュ - オーダー無関係

Question

グローバルメモリ転送を節約し、コードのすべてのステップが個別に動作するため、すべてのカーネルを1つのカーネルに結合しようとしました。ステップはデバイスとして実行され、グローバルではなくコールされます。 これは最初のステップの後半では失敗しています。

イメージの2つの半分を計算するために2回呼び出す必要がある関数があります。イメージが計算される順序にかかわらず、2回目の反復でクラッシュします。

コードを調べて、私ができることと同じように、異なるリターンポイントで複数回実行したところ、クラッシュする原因が見つかりました。

__device__ 
void IntersectCone(float* ModDistance, 
       float* ModIntensity, 
       float3 ray, 
       int threadID, 
       modParam param) 
{ 

bool ignore = false; 

float3 normal = make_float3(0.0f,0.0f,0.0f); 
float3 result = make_float3(0.0f,0.0f,0.0f); 
float normDist = 0.0f; 
float intensity = 0.0f; 

float check = abs(Dot(param.position, Cross(param.direction,ray))); 
if(check > param.r1 && check > param.r2) 
    ignore = true; 

float tran = param.length/(param.r2/param.r1 - 1); 
float length = tran + param.length; 
float Lsq = length * length; 
float cosSqr = Lsq/(Lsq + param.r2 * param.r2); 

//Changes the centre position? 
float3 position = param.position - tran * param.direction; 

float aDd = Dot(param.direction, ray); 
float3 e = position * -1.0f; 
float aDe = Dot(param.direction, e); 
float dDe = Dot(ray, e); 
float eDe = Dot(e, e); 
float c2 = aDd * aDd - cosSqr; 
float c1 = aDd * aDe - cosSqr * dDe; 
float c0 = aDe * aDe - cosSqr * eDe; 

float discr = c1 * c1 - c0 * c2; 

if(discr <= 0.0f) 
    ignore = true; 

if(!ignore) 
{ 
    float root = sqrt(discr); 
    float sign; 

    if(c1 > 0.0f) 
     sign = 1.0f; 
    else 
     sign = -1.0f; 

    //Try opposite sign....? 
    float3 result = (-c1 + sign * root) * ray/c2; 


    e = result - position; 
    float dot = Dot(e, param.direction);   
    float3 s1 = Cross(e, param.direction);   
    float3 normal = Cross(e, s1); 

    if((dot > tran) || (dot < length)) 
    { 
     if(Dot(normal,ray) <= 0) 
     { 
      normal = Norm(normal); //This stuff (1) 
      normDist = Magnitude(result); 
      intensity = -IntensAt1m * Dot(ray, normal)/(normDist * normDist); 
     } 
    } 
} 
ModDistance[threadID] = normDist; and this stuff (2) 
ModIntensity[threadID] = intensity; 
} 

機能のポイントを否定し、どちらもオフ私は、これはクラッシュしないにするために行うことができる2つのものがあります：私はModDistance []とModIntensity []、または私は場合に書き込みをしようとしない場合normDistと強度に書き込まないでください。

最初のチャンスの例外は、上記のコードによってスローされますが、いずれかのブロックがコメントアウトされている場合は発生しません。 また、プログラムはこのルーチンが呼び出される2回目にクラッシュします。

これを一日中考えようとしているが、どんな助けも素晴らしいだろう。

それが呼び出すコード：

int subrow = threadIdx.y + Mod_Height/2; 
int threadID = subrow * (Mod_Width+1) + threadIdx.x;   
int obsY = windowY + subrow; 
float3 ray = CalculateRay(obsX,obsY); 

if(!IntersectSphere(ModDistance, ModIntensity, ray, threadID, param)) 
{ 
    IntersectCone(ModDistance, ModIntensity, ray, threadID, param); 
} 

subrow = threadIdx.y; 
threadID = subrow * (Mod_Width+1) + threadIdx.x;   
obsY = windowY + subrow; 
ray = CalculateRay(obsX,obsY); 

if(!IntersectSphere(ModDistance, ModIntensity, ray, threadID, param)) 
{ 
    IntersectCone(ModDistance, ModIntensity, ray, threadID, param); 
} 

Answer 1

カーネルは、リソースが不足しています。コメントに掲載されているように、それはエラーCudaErrorLaunchOutOfResourcesを与えていました。

これを避けるには、__launch_bounds__指定子を使用して、カーネルに必要なブロックサイズを指定する必要があります。これにより、コンパイラは十分なリソースがあることを保証します。 __launch_bounds__の詳細については、CUDAプログラミングガイドを参照してください。