bi-directional PCI-E transfer overlap

[fishbupt]

I run your code on my new GTX480 fermi card with cuda sdk3.1beta, but still overlap bi-directional data transfer.

as NV announced that fermi should support two DMA transfer engine, which mean PCI-E duplex communication should be enabled.

I have a GTX280 with PCI-E bus.

I intended to hide the bi-directional data transfer time between host and device using two streams.

Here is the test code:

include <cuda.h>

include <stdio.h>

include <sys/time.h>

define N (4096*4096)

int main (void) {

float2 *in, *out, *i_d, *o_d;

struct timeval tim;

double t1,t2;

cudaHostAlloc( (void **)&in, sizeof(float2)*N, cudaHostAllocDefault);

cudaHostAlloc( (void **)&out, sizeof(float2)*N, cudaHostAllocDefault);

cudaMalloc((void**)&i_d, sizeof(float2)*N);

cudaMalloc((void**)&o_d, sizeof(float2)*N);

cudaStream_t stream1,stream2;

cudaStreamCreate(&stream1);

cudaStreamCreate(&stream2);

gettimeofday(&tim, NULL);

t1=tim.tv_sec+(tim.tv_usec/1000000.0);

cudaMemcpyAsync(i_d, in, sizeof(float2)*N, cudaMemcpyHostToDevice,stream1);

cudaMemcpyAsync(out, o_d, sizeof(float2)*N, cudaMemcpyDeviceToHost,stream2);

cudaThreadSynchronize();

gettimeofday(&tim, NULL);

t2=tim.tv_sec+(tim.tv_usec/1000000.0);

printf(" Run time=%f\n",(t2-t1));

cudaFreeHost(in);

cudaFreeHost(out);

cudaFree(i_d);

cudaFree(o_d);

return 0;

}

Steam1 along takes 40ms; stream2 along take 39 ms;

All together they take 79ms, there is no overlap!

PCI-E supports full duplex communications.

Can anybody explain why is this? Is there a way to overlap at all?

[fishbupt]

I run your code on my new GTX480 fermi card with cuda sdk3.1beta, but still overlap bi-directional data transfer.

as NV announced that fermi should support two DMA transfer engine, which mean PCI-E duplex communication should be enabled.

I have a GTX280 with PCI-E bus.

I intended to hide the bi-directional data transfer time between host and device using two streams.

Here is the test code:

include <cuda.h>

include <stdio.h>

include <sys/time.h>

define N (4096*4096)

int main (void) {

float2 *in, *out, *i_d, *o_d;

struct timeval tim;

double t1,t2;

cudaHostAlloc( (void **)&in, sizeof(float2)*N, cudaHostAllocDefault);

cudaHostAlloc( (void **)&out, sizeof(float2)*N, cudaHostAllocDefault);

cudaMalloc((void**)&i_d, sizeof(float2)*N);

cudaMalloc((void**)&o_d, sizeof(float2)*N);

cudaStream_t stream1,stream2;

cudaStreamCreate(&stream1);

cudaStreamCreate(&stream2);

gettimeofday(&tim, NULL);

t1=tim.tv_sec+(tim.tv_usec/1000000.0);

cudaMemcpyAsync(i_d, in, sizeof(float2)*N, cudaMemcpyHostToDevice,stream1);

cudaMemcpyAsync(out, o_d, sizeof(float2)*N, cudaMemcpyDeviceToHost,stream2);

cudaThreadSynchronize();

gettimeofday(&tim, NULL);

t2=tim.tv_sec+(tim.tv_usec/1000000.0);

printf(" Run time=%f\n",(t2-t1));

cudaFreeHost(in);

cudaFreeHost(out);

cudaFree(i_d);

cudaFree(o_d);

return 0;

}

Steam1 along takes 40ms; stream2 along take 39 ms;

All together they take 79ms, there is no overlap!

PCI-E supports full duplex communications.

Can anybody explain why is this? Is there a way to overlap at all?

[Uncle_Joe]

Could someone clarify if bidirectional copies can be done on non-Tesla Fermis?

Tim said it can only be done on Tesla Fermis,

but fishbupt, says above it can be done on GTX 480?

I would like to know so I can decide if it’s worth to get a GTX 470 or 480 so another developer can test & benchmark, without buying another Tesla C2050.

[gshi]

I have seen bi-directional bandwidth in GTX480 in our lab. The result is reported it here

http://forums.nvidia.com/index.php?showtop…t=#entry1073830

But I also tested another system with GTX480 in it in some other lab. In that system I do not see the bi-directional bandwidth.

According to “sumitq” in the above link, Fermi consumer card should not have the bi-directional capability.

- Tesla 20-series has 2 DMA Engines (copy engines). GeForce has 1 DMA Engine. This

means that CUDA applications can overlap computation and communication on Tesla using

bi-directional communication over PCI-e.

Could someone clarify if bidirectional copies can be done on non-Tesla Fermis?

Tim said it can only be done on Tesla Fermis,

but fishbupt, says above it can be done on GTX 480?

I would like to know so I can decide if it’s worth to get a GTX 470 or 480 so another developer can test & benchmark, without buying another Tesla C2050.