DispositivoAHost y HostADispositivo tiempos en caso de Memoria Unificada CUDA.

Estoy tratando de comparar los tiempos totales de ejecución tomados por las versiones administradas y no administradas de la gestión de memoria de CUDA.

En el siguiente código de muestra, tengo dos funciones que hacen exactamente lo mismo. La única diferencia está en su gestión de memoria. Una función utiliza cudaMalloc()/cudaMemcpy() y el otro método utiliza solo cudaMallocManaged().

Usé nvprof para calcular diferentes tiempos y obtuve las siguientes salidas:

Salida de nvprof versión Administrada:

== 29028 == Resultado de perfilado:
Tipo Tiempo (%) Tiempo Llamadas Promedio Mínimo Máximo Nombre
GPU actividades: 100.00 % 59.425us 1 59.425us 59.425us 59.425us add(int, float*, float*)
Llamadas API: 78.08 % 296,49ms 2 148,24ms 1.7127ms 294.78ms cudaMallocManaged
19,61 % 74.451ms 1 74.451ms 74.451ms 74.451ms cuDevicePrimaryCtxRelease
1.55 % 5.8705ms 1 5.8705ms 5.8705ms 5.8705ms cudaLaunchKernel
0.67 % 2.5547ms 2 1.2774ms 974.40us 1.5803ms cudaFree
0.07 % 280.60us 1 280.60us 280.60us 280.60us cudaDeviceSynchronize
0.01 % 28.300us 3 9.4330us 3.0000us 13.300us cuModuleUnload
0.01 % 26.800us 1 26.800us 26.800us 26.800us cuDeviceTotalMem
0.00 % 17.700us 101 175ns 100ns 900ns cuDeviceGetAttribute
0.00 % 10.100us 3 3.3660us 300ns 8.8000us cuDeviceGetCount
0.00 % 3.2000us 1 3.2000us 3.2000us 3.2000us cuDeviceGetName
0.00 % 3.0000us 2 1.5000us 300ns 2.7000us cuDeviceGet
0.00 % 500ns 1 500ns 500ns 500ns cuDeviceGetLuid
0.00 % 200ns 1 200ns 200ns 200ns cuDeviceGetUuid

== 29028 == Resultado de la perfilación de memoria unificada:
Dispositivo “GeForce GTX 1070 (0)”
Cuenta Tamaño promedio Tamaño mínimo Tamaño máximo Tamaño total Tiempo total Nombre
64 128.00KB 128.00KB 128.00KB 8.000000MB 3.279000ms Host A Device
146 84.164KB 32.000KB 1.0000MB 12.00000MB 64.50870ms Dispositivo a Host

Salida de nvprof versión no administrada:

== 23864 == Resultado de perfilado:
Tipo Tiempo (%) Tiempo Llamadas Promedio Mínimo Máximo Nombre
GPU actividades: 56.30 % 1.5032ms 2 751.60us 751.44us 751.76us[CUDA memcpy HtoD]
41.48 % 1.1075ms 1 1.1075ms 1.1075ms 1.1075ms[CUDA memcpy DtoH]
2.23 % 59.457us 1 59.457us 59.457us 59.457us add(int, float*, float*)
Llamadas API: 78.92 % 270.08ms 2 135.04ms 656.40us 269.43ms cudaMalloc
19.79 % 67.730ms 1 67.730ms 67.730ms 67.730ms cuDevicePrimaryCtxRelease
1.05 % 3.5796ms 3 1.1932ms 1.0106ms 1.4341ms cudaMemcpy
0.10 % 346.20us 2 173.10us 3.4000us 342.80us cudaFree
0.09 % 314.30us 1 314.30us 314.30us 314.30us cudaDeviceSynchronize
0.02 % 74.200us 1 74.200us 74.200us 74.200us cudaLaunchKernel
0.01 % 34.700us 3 11.566us 2.5000us 29.100us cuModuleUnload
0.01 % 24.100us 1 24.100us 24.100us 24.100us cuDeviceTotalMem
0.00 % 17.100us 101 169ns 100ns 900ns cuDeviceGetAttribute
0.00 % 9.0000us 3 3.0000us 300ns 8.0000us cuDeviceGetCount
0.00 % 3.2000us 1 3.2000us 3.2000us 3.2000us cuDeviceGetName
0.00 % 1.5000us 2 750ns 200ns 1.3000us cuDeviceGet
0.00 % 300ns 1 300ns 300ns 300ns cuDeviceGetUuid
0.00 % 300ns 1 300ns 300ns 300ns cuDeviceGetLuid

Mi código:

int RunManagedVersion()
{
int N = 1 << 20;
float* x, * y;

// Asignar memoria unificada - accesible desde CPU o GPU
cudaMallocManaged(&x, N * sizeof(float));
cudaMallocManaged(&y, N * sizeof(float));

// inicializar los arreglos x e y en el host
for (int i = 0; i < N; i++) {
    x[i] = 1.0f;
    y[i] = 2.0f;
}

// Lanzar kernel en 1M elementos en la GPU
int blockSize = 256;
int numBlocks = (N + blockSize - 1) / blockSize;
add << <numBlocks, blockSize >> > (N, x, y);

// Esperar a que la GPU termine antes de acceder en el host
cudaDeviceSynchronize();

// Verificar errores (todos los valores deben ser 3.0f)
float maxError = 0.0f;
for (int i = 0; i < N; i++)
    maxError = fmax(maxError, fabs(y[i] - 3.0f));

std::cout << &