Карацуба - умножение полиномов с помощью CUDA - PullRequest
Новая
       28

Карацуба - умножение полиномов с помощью CUDA

0 голосов
/ 25 апреля 2018

Я использую CUDA для итеративного алгоритма Карацубы, и я хотел бы спросить, почему одна строка всегда вычисляется по-разному.

Сначала я реализовал эту функцию, которая всегда правильно вычисляла результат: 

__global__ void kernel_res_main(TYPE *A, TYPE *B, TYPE *D, TYPE *result, TYPE size, TYPE resultSize){
    int i = blockDim.x * blockIdx.x + threadIdx.x;

    if( i > 0 && i < resultSize - 1){

        TYPE start = (i >= size) ? (i % size ) + 1 : 0;


        TYPE end = (i + 1) / 2;


        for(TYPE inner = start; inner < end; inner++){
            result[i] += ( A[inner] + A[i - inner] ) * ( B[inner] + B[i - inner] );
            result[i] -= ( D[inner] + D[i-inner] );
        }
    }
}

Теперь я хотел бы использовать 2D-сетку и использовать CUDA для цикла for, поэтому я изменил свою функцию следующим образом: 

__global__ void kernel_res_nested(TYPE *A, TYPE *B, TYPE *D, TYPE *result, TYPE size, TYPE resultSize){

    int i = blockDim.x * blockIdx.x + threadIdx.x;
    int j = blockDim.y * blockIdx.y + threadIdx.y;

    TYPE rtmp = result[i];

    if( i > 0 && i < resultSize - 1){

        TYPE start = (i >= size) ? (i % size ) + 1 : 0;
        TYPE end = (i + 1) >> 1;

        if(j >= start && j <= end ){

           // WRONG 
           rtmp += ( A[j] + A[i - j] ) * ( B[j] + B[i - j] ) - ( D[j] + D[i - j] );
        }
    }

    result[i] = rtmp;
}

Я вызываю эту функцию следующим образом:

dim3 block( 32, 8 );
dim3 grid( (resultSize+1/32) , (resultSize+7/8) );
kernel_res_nested <<<grid, block>>> (devA, devB, devD, devResult, size, resultSize);

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

Спасибо.

1 Ответ

0 голосов
/ 25 апреля 2018

Для такого вопроса вы должны предоставить MCVE.(См. Пункт 1 здесь ) Например, я не знаю, какой тип обозначен TYPE, и это имеет значение для правильности решения, которое я предложу.

ВВаше первое ядро, только один поток во всей вашей сетке считывал и записывал местоположение result[i].Но во втором ядре у вас есть несколько потоков, записывающих в папку result[i].Они конфликтуют друг с другом.CUDA не определяет порядок, в котором будут выполняться потоки, и некоторые могут запускаться до, после или одновременно с другими.В этом случае некоторые потоки могут читать result[i] одновременно с другими.Затем, когда потоки запишут свои результаты, они будут противоречивыми.И это может варьироваться от бега к бегу.У вас есть условие гонки (зависимость от порядка выполнения, а не от данных).

Канонический метод, чтобы разобраться в этом, состоял бы в использовании сокращения техника.

Однако для простоты я предлагаю, чтобы атомика могла помочь вам разобраться с этим.Это легче реализовать на основе того, что вы показали, и поможет подтвердить состояние гонки.После этого, если вы хотите попробовать метод сокращения, для этого есть множество учебных пособий (один из них приведен выше) и множество вопросов по тегу cuda об этом.

Вы можете изменить свое ядро.что-то вроде этого, чтобы разобраться в состоянии гонки: 

__global__ void kernel_res_nested(TYPE *A, TYPE *B, TYPE *D, TYPE *result, TYPE size, TYPE resultSize){

    int i = blockDim.x * blockIdx.x + threadIdx.x;
    int j = blockDim.y * blockIdx.y + threadIdx.y;

    if( i > 0 && i < resultSize - 1){

        TYPE start = (i >= size) ? (i % size ) + 1 : 0;
        TYPE end = (i + 1) >> 1;

        if(j >= start && j < end ){ // see note below

           atomicAdd(result+i, (( A[j] + A[i - j] ) * ( B[j] + B[i - j] ) - ( D[j] + D[i - j] )));
        }
    }

}

Обратите внимание, что в зависимости от типа вашего графического процессора и фактического типа TYPE, который вы используете, это может не сработать (может не скомпилироваться)как есть.Но так как вы ранее использовали TYPE в качестве переменной цикла, я предполагаю, что это целочисленный тип, и необходимые atomicAdd для них должны быть доступны.

Несколько других комментариев:

  1. Это может не дать вам ожидаемый размер сетки: 

    dim3 grid( (resultSize+1/32) , (resultSize+7/8) );

    Я думаю, что обычные вычисления будут: 

    dim3 grid( (resultSize+31)/32, (resultSize+7)/8 );

  2. Я всегда рекомендую правильную проверку ошибок CUDA и запуск ваших кодов с cuda-memcheck, всякий раз, когда у вас возникают проблемы с кодом CUDA, чтобы убедиться, что нет ошибок времени выполнения.

  3. Это также выглядит для меня следующим образом: 

    if(j >= start && j <= end ){

    должно быть таким: 

    if(j >= start && j < end ){

    в соответствии с вашим диапазоном цикла.Я также делаю предположение, что size меньше resultSize (опять же, MCVE поможет).

...