Для начала несколько общих комментариев. Ваша петля
for (i=0;i<n;i++)
for (j=0;j<n;j++)
v[i]=v[i]+(B[i*n+j]*d[j]);
является эквивалентом стандартной BLAS gemv операции
![enter image description here](https://i.stack.imgur.com/BlQ7H.gif)
где матрица хранится в главном порядке строк. Оптимальным способом сделать это на устройстве было бы использование CUBLAS, а не чего-то созданного из примитивов тяги.
Сказав это, абсолютно невозможно, чтобы код направления, который вы опубликовали, когда-либо делал то, что делает ваш серийный код. Ошибки, которые вы видите, не являются результатом ассоциативности с плавающей точкой. Фундаментально thrust::transform
применяет функтор, предоставляемый каждому элементу входного итератора, и сохраняет результат в выходном итераторе. Чтобы получить тот же результат, что и опубликованный вами цикл, вызов thrust::transform
должен будет выполнить (n * n) операций с объявленным вами функтором fmad. Ясно, что это не так. Кроме того, нет никакой гарантии, что thrust::transform
выполнит операцию суммирования / сокращения таким образом, который будет безопасен для гонок памяти.
Правильное решение, вероятно, будет примерно таким:
- Используйте команду thrust :: transform для вычисления (n * n) произведений элементов B и d
- Используйте команду thrust :: redu_by_key, чтобы преобразовать продукты в частичные суммы, получив Bd
- Используйте thrust :: transform, чтобы добавить полученный матрично-векторный продукт к v , чтобы получить окончательный результат.
В коде сначала определите функтор следующим образом:
struct functor
{
template <typename Tuple>
__host__ __device__
double operator()(Tuple v)
{
return thrust::get<0>(v) * thrust::get<1>(v);
}
};
Затем выполните следующие действия для вычисления умножения матрицы на вектор
typedef thrust::device_vector<int> iVec;
typedef thrust::device_vector<double> dVec;
typedef thrust::counting_iterator<int> countIt;
typedef thrust::transform_iterator<IndexDivFunctor, countIt> columnIt;
typedef thrust::transform_iterator<IndexModFunctor, countIt> rowIt;
// Assuming the following allocations on the device
dVec B(n*n), v(n), d(n);
// transformation iterators mapping to vector rows and columns
columnIt cv_begin = thrust::make_transform_iterator(thrust::make_counting_iterator(0), IndexDivFunctor(n));
columnIt cv_end = cv_begin + (n*n);
rowIt rv_begin = thrust::make_transform_iterator(thrust::make_counting_iterator(0), IndexModFunctor(n));
rowIt rv_end = rv_begin + (n*n);
dVec temp(n*n);
thrust::transform(make_zip_iterator(
make_tuple(
B.begin(),
thrust::make_permutation_iterator(d.begin(),rv_begin) ) ),
make_zip_iterator(
make_tuple(
B.end(),
thrust::make_permutation_iterator(d.end(),rv_end) ) ),
temp.begin(),
functor());
iVec outkey(n);
dVec Bd(n);
thrust::reduce_by_key(cv_begin, cv_end, temp.begin(), outkey.begin(), Bd.begin());
thrust::transform(v.begin(), v.end(), Bd.begin(), v.begin(), thrust::plus<double>());
Конечно, это ужасно неэффективный способ выполнения вычислений по сравнению с использованием специально разработанного кода умножения матрицы на вектор, такого как dgemv
из CUBLAS.