CUDA: Почему Thrust так медленно загружает данные на GPU?

Question

Я новичок в мире GPU и только что установил CUDA для написания какой-то программы. Я играл с платформой, но обнаружил, что она очень медленная при загрузке данных на GPU. Примерно 35 Мбайт/с в части «устройство-устройство» на моем не-плохом рабочем столе. Как получилось?

Окружающая среда: Visual Studio 2012, CUDA 5.0, GTX760, Intel-i7, Windows 7 x64

GPU тест пропускной способности:

Предполагается иметь, по крайней мере 11GB/с скорости передачи для хост к устройству или наоборот! Но это не так!

Вот тестовая программа:

#include <iostream> 
#include <ctime> 
#include <thrust/device_vector.h> 
#include <thrust/host_vector.h> 

#define N 32<<22 

int main(void) 
{ 
    using namespace std; 

    cout<<"GPU bandwidth test via thrust, data size: "<< (sizeof(double)*N)/1000000000.0 <<" Gbytes"<<endl; 
    cout<<"============program start=========="<<endl; 

    int now = time(0); 
    cout<<"Initializing h_vec..."; 
    thrust::host_vector<double> h_vec(N,0.0f); 
    cout<<"time spent: "<<time(0)-now<<"secs"<<endl; 

    now = time(0); 
    cout<<"Uploading data to GPU..."; 
    thrust::device_vector<double> d_vec = h_vec; 
    cout<<"time spent: "<<time(0)-now<<"secs"<<endl; 

    now = time(0); 
    cout<<"Downloading data to h_vec..."; 
    thrust::copy(d_vec.begin(), d_vec.end(), h_vec.begin()); 
    cout<<"time spent: "<<time(0)-now<<"secs"<<endl<<endl; 

    system("PAUSE"); 
    return 0; 
} 

Программа из пут:

• Скорость загрузки: менее 1 сек, довольно имеет смысл сравнить с номинальным 11GB/с.

• Скорость загрузки: 1.07374 ГБ/32 с составляет около 33,5 МБ/с, что совершенно не имеет смысла.

Кто-нибудь знает причину? Или это просто способ толчка?

Спасибо!

Answer 1

У вас есть несколько недостатков, некоторые из которых описаны в комментариях.

1. Вам необходимо устранить любые эффекты распределения. Вы можете сделать это, сделав сначала «разминки».
2. Вам необходимо устранить любые «пусковые» эффекты. Вы можете сделать это, сделав сначала «разминки».
3. При сравнении данных помните, что bandwidthTest использует выделение памяти PINNED, которое не используется. Поэтому скорость передачи данных тяги будет медленнее. Это обычно приводит к коэффициенту 2x (т. Е. Передача фиксированной памяти, как правило, примерно в 2 раза быстрее, чем переносимая память в память. Если вы хотите получить лучшее сравнение с bandwidthTest, запустите его с помощью переключателя --memory=pageable.
4. Ваш выбор функций синхронизации может быть не лучшим . cudaEvents является довольно надежным для синхронизации операций CUDA

Вот код, который делает правильную синхронизацию:.

$ cat t213.cu 
#include <iostream> 
#include <thrust/device_vector.h> 
#include <thrust/host_vector.h> 
#include <thrust/copy.h> 
#include <thrust/fill.h> 

#define DSIZE ((1UL<<20)*32) 

int main(){ 

    thrust::device_vector<int> d_data(DSIZE); 
    thrust::host_vector<int> h_data(DSIZE); 
    float et; 
    cudaEvent_t start, stop; 
    cudaEventCreate(&start); 
    cudaEventCreate(&stop); 

    thrust::fill(h_data.begin(), h_data.end(), 1); 
    thrust::copy(h_data.begin(), h_data.end(), d_data.begin()); 

    std::cout<< "warm up iteration " << d_data[0] << std::endl; 
    thrust::fill(d_data.begin(), d_data.end(), 2); 
    thrust::copy(d_data.begin(), d_data.end(), h_data.begin()); 
    std::cout<< "warm up iteration " << h_data[0] << std::endl; 
    thrust::fill(h_data.begin(), h_data.end(), 3); 
    cudaEventRecord(start); 
    thrust::copy(h_data.begin(), h_data.end(), d_data.begin()); 
    cudaEventRecord(stop); 
    cudaEventSynchronize(stop); 
    cudaEventElapsedTime(&et, start, stop); 
    std::cout<<"host to device iteration " << d_data[0] << " elapsed time: " << (et/(float)1000) << std::endl; 
    std::cout<<"apparent bandwidth: " << (((DSIZE*sizeof(int))/(et/(float)1000))/((float)1048576)) << " MB/s" << std::endl; 
    thrust::fill(d_data.begin(), d_data.end(), 4); 
    cudaEventRecord(start); 
    thrust::copy(d_data.begin(), d_data.end(), h_data.begin()); 
    cudaEventRecord(stop); 
    cudaEventSynchronize(stop); 
    cudaEventElapsedTime(&et, start, stop); 
    std::cout<<"device to host iteration " << h_data[0] << " elapsed time: " << (et/(float)1000) << std::endl; 
    std::cout<<"apparent bandwidth: " << (((DSIZE*sizeof(int))/(et/(float)1000))/((float)1048576)) << " MB/s" << std::endl; 

    std::cout << "finished" << std::endl; 
    return 0; 
} 

Я компилировать с (у меня есть PCIE система Gen2 с cc2.0 устройством)

$ nvcc -O3 -arch=sm_20 -o t213 t213.cu 

Когда я запускаю его я получаю следующие результаты:

$ ./t213 
warm up iteration 1 
warm up iteration 2 
host to device iteration 3 elapsed time: 0.0476644 
apparent bandwidth: 2685.44 MB/s 
device to host iteration 4 elapsed time: 0.0500736 
apparent bandwidth: 2556.24 MB/s 
finished 
$ 

Это выглядит правильно для меня, потому что bandwidthTest в моей системе будет сообщать о 6GB/s в обоих направлениях, как у меня есть PCIE системы Gen2 ,Поскольку тяга использует доступную для скачивания, а не закрепленную память, я получаю примерно половину этой полосы пропускания, т. Е. 3 ГБ/с, а толчок сообщает о 2,5 ГБ/с.

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

$ /usr/local/cuda/samples/bin/linux/release/bandwidthTest --memory=pageable 
[CUDA Bandwidth Test] - Starting... 
Running on... 

Device 0: Quadro 5000 
Quick Mode 

Host to Device Bandwidth, 1 Device(s) 
PAGEABLE Memory Transfers 
    Transfer Size (Bytes)  Bandwidth(MB/s) 
    33554432      2718.2 

Device to Host Bandwidth, 1 Device(s) 
PAGEABLE Memory Transfers 
    Transfer Size (Bytes)  Bandwidth(MB/s) 
    33554432      2428.2 

Device to Device Bandwidth, 1 Device(s) 
PAGEABLE Memory Transfers 
    Transfer Size (Bytes)  Bandwidth(MB/s) 
    33554432      99219.1 

$