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

Question

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

Я использую 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);

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

Спасибо.

0

c++cuda karatsuba nvidia parallel-processing

Решение

Другие решения

Других решений пока нет …

Источник

Accepted Answer

Для такого вопроса вы должны предоставить 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 для тех, кто должен быть доступен.

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

Это может не дать вам ожидаемый размер сетки:
```
dim3 grid( (resultSize+1/32) , (resultSize+7/8) );
```
Я думаю, что обычные расчеты там будут:
```
dim3 grid( (resultSize+31)/32, (resultSize+7)/8 );
```
Я всегда рекомендую правильная проверка ошибок CUDA и запустить ваши коды с cuda-memcheckвсякий раз, когда у вас возникают проблемы с кодом CUDA, чтобы убедиться в отсутствии ошибок во время выполнения.
Это также выглядит для меня так:
```
if(j >= start && j <= end ){
```
должно быть так:
```
if(j >= start && j < end ){
```
чтобы соответствовать вашему диапазону цикла. Я также делаю предположение, что size меньше чем resultSize (опять же, MCVE поможет).

0