CUDA неэффективный шаблон доступа к глобальной памяти

Question

Я кодировал ядро ​​CUDA, которое выполняет добавление массивов из двух массивов arr1 и arr2. Информация о том, какой индекс arr1 должен быть добавлен, с каким индексом arr2 хранится в массиве idx.

Вот пример кода:

__global__ add(float* arr1, float* arr2, int* idx, int length)
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;

    // each thread performs (length) additions,
    // arr2 is (lenght) times larger than arr1
    for (int j=threadIdx.x; j<length*blockDim.x; j+=blockDim.x)
    {
        arr1[i] += arr2[ idx[blockIdx.x*blockDim.x + j] ]; // edited here
    }
}

Код выдает правильный вывод, хотя он чуть быстрее, чем код openmp-параллельный на ЦП с 8 потоками. Я пробовал это для разных размеров блока.

Я подозреваю, что шаблон доступа к arr2 неэффективен, поскольку arr2 находится в глобальной памяти и доступен почти случайно - массив idx содержит уникальные, отсортированные, но несмежные индексы (может быть 2, 3, 57, 103, ...). Поэтому кеш L1 не используется. Кроме того, массивы очень большие и не могут полностью поместиться в общей памяти.

Есть ли способ обойти это препятствие? У вас есть идея, как оптимизировать схему доступа к arr2?

Answer 1

Несколько вещей, которые вы можете попробовать сделать здесь:

• Глобальная память медленная, вы можете попробовать загрузить массивы в общую память для каждого блока. Таким образом, вы должны загрузить массив частично для каждого блока, например.
• Если есть что-то, что определено как константа, вы должны переместить это в постоянную память.
• Попробуйте развернуть петлю.
• Обязательно попробуйте запустить несколько ядер с разными параметрами потока, а также с различными комбинациями блоков / потоков.

Надеюсь, это даст вам представление о вашем пути к оптимизации:)