Сбой второй итерации - порядок не имеет значения - PullRequest
Новая
       26

Сбой второй итерации - порядок не имеет значения

1 голос
/ 01 марта 2012

Чтобы сэкономить на глобальных передачах памяти и поскольку все шаги кода работают по отдельности, я попытался объединить все ядра в один ядро, причем первые 2 (из 3) шага выполняются как * 1001.* устройство звонки, а не глобальные звонки.Это терпит неудачу во второй половине первого шага.

Есть функция, которую мне нужно вызвать дважды, чтобы вычислить 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; 
}

Есть две вещи, которые я могу сделать, чтобы сделать это не сбоем, обе из которых сводят на нет суть функции: если я не пытаюсь писать в ModDistance [] и ModIntensity [],или если я не пишу в normDist и интенсивность.

Исключения первого шанса вызываются приведенным выше кодом, но не в том случае, если один из блоков закомментирован.Кроме того, программа аварийно завершает работу только во второй раз, когда вызывается эта подпрограмма.

Пытались выяснить это весь день, любая помощь была бы фантастической.

Код, который ее вызывает: 

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

1 Ответ

2 голосов
/ 06 марта 2012

В ядре заканчиваются ресурсы.Как отмечалось в комментариях, он выдавал ошибку CudaErrorLaunchOutOfResources.

Чтобы избежать этого, вы должны использовать спецификатор __launch_bounds__, чтобы указать размеры блока, которые вы хотите для своего ядра.Это заставит компилятор убедиться, что ресурсов достаточно.Подробнее см. Руководство по программированию CUDA __launch_bounds__.

...