CUDA 4.0 - cudaHostUnregister가 느립니다.

Question

저는 cudaHostRegister 및 cudaHostUnregister 함수로 작업 해 왔으며 후자가 매우 오래 걸리는 것으로 나타났습니다. 동일한 데이터에서 cudaMemcpy와 cudaHostUnregister를 비교하는 경우에도 memcpy에 페이지 잠금 메모리를 사용하지 않아도 매우 오랜 시간이 걸립니다.

나는 다음과 같은 쇼트 프로그램했습니다 :

#include <stdio.h> 
#include <time.h> 
#include <assert.h> 
#include <stdlib.h> 

static struct timespec tp; 
static clockid_t clk = CLOCK_REALTIME; 

static void tu_timer_start(void) 
{ 
    int res = clock_gettime(clk, &tp); 
    assert(!res); 
} 

static long long tu_timer_stop(void) 
{ 
    struct timespec tp_new; 
    long long elapsed; 
    int res = clock_gettime(clk, &tp_new); 

    assert(!res); 

    elapsed = 1000000000LL * (tp_new.tv_sec - tp.tv_sec) + tp_new.tv_nsec - tp.tv_nsec; 
    tp = tp_new; 

    return elapsed; 
} 

int main() { 
    const int length = 999424; 
    const int pagesize = 4096; 

    // Allocating page-aligned host data and filling it with zeroes. 
    int *paged, *locked; 
    posix_memalign((void**) &paged, pagesize, length * sizeof(int)); 
    posix_memalign((void**) &locked, pagesize, length * sizeof(int)); 
    memset(paged, 0, length * sizeof(int)); 
    memset(locked, 0, length * sizeof(int)); 

    // Allocating device data. 
    int *devPaged, *devLocked; 
    tu_timer_start(); 
    printf("%20d\n", cudaMalloc(&devPaged, length * sizeof(int))); 
    printf("%20d\n", cudaMalloc(&devLocked, length * sizeof(int))); 
    printf("Initialization: %12lld ns\n", tu_timer_stop()); 

    // Measure copy time with pageable data. 
    tu_timer_start(); 
    printf("%20d\n", cudaMemcpy(devPaged, paged, length * sizeof(int), cudaMemcpyHostToDevice)); 
    printf("Copy pageable: %12lld ns\n", tu_timer_stop()); 

    // Measure time to page-lock host data. 
    tu_timer_start(); 
    printf("%20d\n", cudaHostRegister(locked, length * sizeof(int), 0)); 
    printf("Host register: %12lld ns\n", tu_timer_stop()); 

    // Measure copy time with page-locked data. 
    tu_timer_start(); 
    printf("%20d\n", cudaMemcpy(devLocked, locked, length * sizeof(int), cudaMemcpyHostToDevice)); 
    printf("Copy page-locked: %12lld ns\n", tu_timer_stop()); 

    // Measure time to release page-lock on host data. 
    tu_timer_start(); 
    cudaHostUnregister(locked); 
    printf("Host unregister: %12lld ns\n", tu_timer_stop()); 

    return 0; 
} 

이 인텔 I5 760 CUDA 리턴 코드와 테슬라 C2050 (와 (코어 당 2.80 GHz의)이 인쇄되지 쿼드 코어에 다음과 같은 출력을 제공을 여기) :

Initialization:  81027005 ns 
Copy pageable:   1263236 ns 
Host register:   436132 ns 
Copy page-locked:  706051 ns 
Host unregister:  2139736 ns 

이것은 내 문제를 보여줍니다. 내 실제 프로그램에서 그것은 더 나쁘다. 나는 종종 cudaHostUnregister를 약 3460000 ns 측정한다. 이것은 천천히 진행되는 것 외에 동시 비동기식 memcopies 또는 커널 실행에서 잘 작동하지 않는다는 것을 암시합니다.

왜이 기능이 오래 걸리고 속도를 높이는 방법이 있습니까? 그리고 실제로 memcopies와 커널과 병렬로 작동하지 않습니까? 그렇다면 왜 그렇게하지 않습니까?

아니면 단순히 memcopies와 커널을 병렬 처리하는 더 좋은 방법이 있습니까?

Answer 1

플랫폼에 따라 다르지만 cuMemHostUnregister()/cudaHostUnregister()를 호출 할 때 드라이버에게 수행 할 작업을 묻는 메시지가 나오지 않습니다. GPU 메모리를 매핑 해제하고 페이지별로 표시 할 수 있습니다. 호스트 운영 체제를 다시 시작하십시오. 이러한 작업은 다음을 수반 할 수 있습니다.

1) 드라이버가 GPU 명령 보류로 인해 메모리가 필요한지 여부를 알기가 매우 어렵 기 때문에 GPU와 동기화하십시오. 2) GPU 페이지 테이블은 커널 모드에서만 편집 할 수 있으므로 커널 썽크를 수행하십시오. 3) 하드웨어 레지스터를 업데이트하여 메모리를 매핑 해제합니다.

메모리가 GPU 용으로 더 이상 매핑되지 않으면 드라이버는 페이지 잠금을 해제 할 수 있습니다. 이는 또한 성능이 플랫폼에 따라 좌우되는 값 비싼 작업 일 수 있습니다.

내 제안은 CUDA에 대해 '등록'된 메모리를 남겨두고 휴리스틱에 따라 등록을 취소하는 것입니다 (예 : 등록을 가비지 수집하거나 등록이 실패하면 "룸 만들기").

GPU가 여러 개 있고 Unified Virtual Addressing이 적용되는 경우 드라이버는 시스템의 모든 GPU에 대해 이러한 작업을 수행해야합니다.