Почему графические процессоры NVIDIA Pascal медленно запускают ядра CUDA при использовании cudaMallocManaged ⇐ C++

[Anonymous]

Я тестировал новый CUDA 8 вместе с графическим процессором Pascal Titan X и ожидал ускорения моего кода, но по какой-то причине он оказался медленнее. Я использую Ubuntu 16.04.

Вот минимальный код, который может воспроизвести результат:

CUDASample.cuh

Код: Выделить всё

class CUDASample{
public:
void AddOneToVector(std::vector &in);
};

CUDASample.cu

Код: Выделить всё

__global__ static void CUDAKernelAddOneToVector(int *data)
{
const int x  = blockIdx.x * blockDim.x + threadIdx.x;
const int y  = blockIdx.y * blockDim.y + threadIdx.y;
const int mx = gridDim.x * blockDim.x;

data[y * mx + x] = data[y * mx + x] + 1.0f;
}

void CUDASample::AddOneToVector(std::vector &in){
int *data;
cudaMallocManaged(reinterpret_cast(&data),
in.size() * sizeof(int),
cudaMemAttachGlobal);

for (std::size_t i = 0; i < in.size(); i++){
data[i] = in.at(i);
}

dim3 blks(in.size()/(16*32),1);
dim3 threads(32, 16);

CUDAKernelAddOneToVector(data);

cudaDeviceSynchronize();

for (std::size_t i = 0; i < in.size(); i++){
in.at(i) = data[i];
}

cudaFree(data);
}

Main.cpp

Код: Выделить всё

std::vector v;

for (int i = 0; i < 8192000; i++){
v.push_back(i);
}

CUDASample cudasample;

cudasample.AddOneToVector(v);

Единственное отличие — это флаг NVCC, который для Pascal Titan X:

Код: Выделить всё

-gencode arch=compute_61,code=sm_61-std=c++11;

а для старого Maxwell Titan X это:

Код: Выделить всё

-gencode arch=compute_52,code=sm_52-std=c++11;

РЕДАКТИРОВАТЬ: Вот результаты запуска NVIDIA Visual Profiling.

Для старого Maxwell Titan время передачи памяти составляет около 205 мс, а запуск ядра — около 268 мс.



Для Pascal Titan время передачи памяти составляет около 202 мс, а запуск ядра занимает безумно долгое время 8343 мс, что заставляет меня думать, что что-то не так.



Я дополнительно изолирую проблему, заменив cudaMallocManaged на старый добрый cudaMalloc, выполнил некоторое профилирование и наблюдаю интересный результат.

CUDASample.cu

Код: Выделить всё

__global__ static void CUDAKernelAddOneToVector(int *data)
{
const int x  = blockIdx.x * blockDim.x + threadIdx.x;
const int y  = blockIdx.y * blockDim.y + threadIdx.y;
const int mx = gridDim.x * blockDim.x;

data[y * mx + x] = data[y * mx + x] + 1.0f;
}

void CUDASample::AddOneToVector(std::vector &in){
int *data;
cudaMalloc(reinterpret_cast(&data), in.size() * sizeof(int));
cudaMemcpy(reinterpret_cast(data),reinterpret_cast(in.data()),
in.size() * sizeof(int), cudaMemcpyHostToDevice);

dim3 blks(in.size()/(16*32),1);
dim3 threads(32, 16);

CUDAKernelAddOneToVector(data);

cudaDeviceSynchronize();

cudaMemcpy(reinterpret_cast(in.data()),reinterpret_cast(data),
in.size() * sizeof(int), cudaMemcpyDeviceToHost);

cudaFree(data);
}

Для старого Maxwell Titan время передачи памяти составляет около 5 мс в обе стороны, а запуск ядра — около 264 мкс.



Для Pascal Titan время передачи памяти составляет около 5 мс в обе стороны, а запуск ядра составляет около 194 мс, что фактически приводит к увеличению производительности, которое я надеюсь увидеть...



Почему графический процессор Pascal так медленно запускает ядра CUDA при использовании cudaMallocManaged? Будет пародией, если мне придется вернуть весь существующий код, использующий cudaMallocManaged, в cudaMalloc. Этот эксперимент также показывает, что время передачи памяти при использовании cudaMallocManaged намного медленнее, чем при использовании cudaMalloc, что также создает ощущение, что что-то не так. Если использование этого приводит к замедлению времени выполнения, даже если код проще, это должно быть неприемлемо, поскольку вся цель использования CUDA вместо простого C++ — ускорить работу. Что я делаю не так и почему наблюдаю такой результат?

Подробнее здесь: https://stackoverflow.com/questions/397 ... mallocmana