#include <MatrixVectorMul.cuh> __global__ void Ma...

В вашем коде есть несколько потенциальных проблем. Вот некоторые из них:

1. Неверный порядок индексов в ядре CUDA: В строке, где вы выполняете операцию умножения матрицы на вектор:

   cpp
   result[index] += vector[i] * matrix[i + height * index];

   Вам нужно правильно индексировать элементы матрицы. Обычно матрица представлена как одномерный массив в памяти, и индексация строки и столбца будет выглядеть так:

   cpp
   result[index] += vector[i] * matrix[index * width + i];

   Это правильно индексирует элементы матрицы с учётом ее ширины.

2. Множество блоков и нитей: Ваш код использует два блока с по 256 нитей каждый:

   cpp
   MatrixVectorMul<<<2, 256>>>(512, 256, matrix_gpu, vector_gpu, result);

   Однако, у вас есть 512 элементов в результате, что требует 512 нитей для полного выполнения. Лучше использовать 512 нитей (или больше блоков с меньшим количеством нитей), чтобы покрыть весь диапазон индексов. Например, вы можете использовать:

   cpp
   MatrixVectorMul<<<1, 512>>>(512, 256, matrix_gpu, vector_gpu, result);

3. Ошибки при тестировании: В вашем тесте вы вызываете функцию MatrixVectorMul, но не проверяете, произошли ли ошибки CUDA после вызова ядра (вы проверяете только после cudaPeekAtLastError()), это может привести к незаметным проблемам, если вызов ядра не завершится с ошибкой. Чтобы упростить проверку ошибок, используйте cudaDeviceSynchronize() перед проверкой ошибок:

   cpp
   MatrixVectorMul<<<1, 512>>>(512, 256, matrix_gpu, vector_gpu, result);
   cudaDeviceSynchronize();
   gpuErrchkTest(cudaPeekAtLastError());

4. Ошибка при выделении памяти для результатов: Вы используете функцию Alloc(512) для выделения памяти, но не указано, что это за функция. Если она не предоставляет правильный способ выделения памяти на GPU, это может привести к ошибкам. Убедитесь, что Alloc правильно выделяет память на устройстве.

Итак, исправленный код ядра и теста будет выглядеть следующим образом:

cpp
__global__
void MatrixVectorMul(int height, int width, float* matrix, float* vector, float* result) {
  int index = blockIdx.x * blockDim.x + threadIdx.x;
  if (index < height) {
    result[index] = 0;
    for (int i = 0; i < width; ++i) {
      result[index] += vector[i] * matrix[index * width + i];
    }
  }
}

TEST_F(MatrixVectorMulTests, RandomMul) {
    auto matrix = nc::random::randN<float>(nc::Shape(512, 256));
    auto vector = nc::random::randN<float>(nc::Shape(1, 256));

    auto matrix_gpu = Gpu(matrix);
    auto vector_gpu = Gpu(vector);

    auto result = Alloc(512);

    MatrixVectorMul<<<1, 512>>>(512, 256, matrix_gpu, vector_gpu, result);
    cudaDeviceSynchronize();
    gpuErrchkTest(cudaPeekAtLastError());
    float* actual_result = Cpu(result, 512);

    auto expected_result = matrix.dot(vector.transpose());
    auto np_actual = nc::NdArray<float>(actual_result, 512).transpose();

    ASSERT_TRUE(allclose(expected_result, np_actual, 0.01));
}

Теперь ядро правильно индексирует элементы и используется правильное количество нитей для выполнения умножения.