시뮬레이션 루프 GPU 활용

Question

시뮬레이션 루프를 활용하는 데 어려움을 겪고 있습니다. 매주기마다 3 개의 커널이 시작되었습니다. 다음 시간 단계 크기는 두 번째 커널에 의해 계산됩니다.

while (time < end) 
{ 
    kernel_Flux<<<>>>(...); 
    kernel_Timestep<<<>>>(d_timestep); 
    memcpy(&h_timestep, d_timestep, sizeof(float), ...); 
    kernel_Integrate<<<>>>(d_timestep); 
    time += h_timestep; 
} 

단지 하나의 플로트 만 복사하면됩니다. 불필요한 동기화를 피하는 가장 효율적인 방법은 무엇입니까?

미리 감사드립니다. :-)

Answer 1

이상적인 솔루션은 모든 것을 GPU로 옮기는 것입니다. 그러나 반복 할 때마다 CUDPP 컴팩트를 실행해야하므로 CUDA 스트림이나 동적 병렬 처리를 지원하지 않기 때문에 그렇게 할 수 없습니다. 나는 추력 1.8 라이브러리에 copy_if 메소드가 있다는 것을 알고 있습니다.이 메소드는 동적 병렬 처리와 함께 작동합니다. 문제는 별도의 컴파일로 컴파일되지 않는다는 것입니다.

이

while (time < end) 
{ 
    kernel_Flux<<<gs,bs, 0, stream1>>>(); 
    kernel_Timestep<<<gs,bs, 0, stream1>>>(d_timestep); 
    cudaEventRecord(event, stream1); 
    cudaStreamWaitEvent(mStream2, event, 0); 
    memcpyasync(&h_timestep, d_timestep, sizeof(float), ..., stream2); 
    kernel_Integrate<<<>>>(d_timestep); 
    cudaStreamSynchronize(stream2); 
    time += h_timestep; 
} 

Answer 2

CUDA에서 기본 스트림에서 실행중인 모든 작업은 동기화됩니다. 따라서 코드를 게시하면 커널이 하나씩 차례로 실행됩니다. 나는 커널에게 kernel_integrate를 볼 수있는() 커널 kernel_Timestep(), 동기화를 피하는 그래서 아무 방법의 결과에서 의존에서. 어쨌든 커널 kernel_Flux 독립적 인 데이터에() 및 kernel_Timestep() 작품은 두 개의 서로 다른 스트림에 병렬로 실행하려고 할 수 있습니다.

Answer 3

반복 시간이 많은 경우 h_timestep의 memcpy 전용 새 스트림을 설정할 수 있습니다 (이 경우에는 cudaMemcpyAsync를 사용해야 함). 그런 다음 투기 적 실행과 같은 것을 사용하십시오. 여기서 시간을 계산하기 전에 루프가 진행됩니다. 이렇게하려면 다음 몇 번의 반복을 위해 GPU 메모리 버퍼를 설정해야합니다. 원형 버퍼를 사용하여이 작업을 수행 할 수 있습니다. 또한 다른 스트림을 동기화하려면 cudaEventRecord과 cudaStreamWaitEvent을 사용해야합니다. 그러면 다음 반복이 덮어 쓰려는 버퍼에 해당하는 경우에만 진행될 수 있습니다 (memcpy 스트림이 작업을 완료했습니다). 그렇지 않으면 그 반복에서 상태를 잃을 것이기 때문입니다.

내가 시도하지는 않았지만 작동 할 것으로 예상되는 또 다른 잠재적 인 솔루션은 동적 병렬 처리를 사용하는 것입니다. 카드가이를 지원하면 전체 루프를 GPU에 넣을 수 있습니다.

편집 :

죄송합니다. 세 번째 커널이 있다는 것을 알았습니다. 동기화로 인한 지연은 cudaMemcpyAsync을 수행하지 않았기 때문일 수 있습니다. 세 번째 커널이 memcpy보다 오래 실행될 가능성이 큽니다. 지체없이 진행할 수 있어야합니다. 필요한 유일한 동기화는 각 반복 후에입니다.