▶ 本章介绍了页锁定内存和流的使用方法，给出了测试内存拷贝、（单 / 双）流控制下的内存拷贝的例子。
● 测试内存拷贝
#include <stdio.h>
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include "D:\Code\CUDA\book\common\book.h" #define SIZE (64*1024*1024)
#define TEST_TIMES (100) float cuda_malloc_test(int size, bool up)
{
cudaEvent_t start, stop;
int *a, *dev_a;
float elapsedTime; cudaEventCreate(&start);
cudaEventCreate(&stop); a = (int*)malloc(size * sizeof(int));
cudaMalloc((void**)&dev_a,size * sizeof(*dev_a)); cudaEventRecord(start, );
if (up)
{
for (int i = ; i < TEST_TIMES; i++)
cudaMemcpy(dev_a, a, size * sizeof(*dev_a), cudaMemcpyHostToDevice);
}
else
{
for (int i = ; i < TEST_TIMES; i++)
cudaMemcpy(a, dev_a, size * sizeof(*dev_a), cudaMemcpyDeviceToHost);
}
cudaEventRecord(stop, );
cudaEventSynchronize(stop);
cudaEventElapsedTime(&elapsedTime, start, stop); free(a);
cudaFree(dev_a);
cudaEventDestroy(start);
cudaEventDestroy(stop);
return elapsedTime;
} float cuda_host_alloc_test(int size, bool up)
{
cudaEvent_t start, stop;
int *a, *dev_a;
float elapsedTime;
cudaEventCreate(&start);
cudaEventCreate(&stop); cudaHostAlloc((void**)&a, size * sizeof(int), cudaHostAllocDefault);
cudaMalloc((void**)&dev_a, size * sizeof(int)); cudaEventRecord(start, );
if (up)
{
for (int i = ; i < TEST_TIMES; i++)
cudaMemcpy(dev_a, a, size * sizeof(int), cudaMemcpyHostToDevice);
}
else
{
for (int i = ; i < TEST_TIMES; i++)
cudaMemcpy(a, dev_a, size * sizeof(int), cudaMemcpyDeviceToHost);
}
cudaEventRecord(stop, );
cudaEventSynchronize(stop);
cudaEventElapsedTime(&elapsedTime, start, stop); cudaFreeHost(a);
cudaFree(dev_a);
cudaEventDestroy(start);
cudaEventDestroy(stop);
return elapsedTime;
} int main(void)
{
float elapsedTime;
float testSizeByte = (float)SIZE * sizeof(int) * TEST_TIMES / / ; elapsedTime = cuda_malloc_test(SIZE, true);
printf("\n\tcudaMalloc Upwards:\t\t%3.1f ms\t%3.1f MB/s",elapsedTime, testSizeByte / (elapsedTime / ));
elapsedTime = cuda_malloc_test(SIZE, false);
printf("\n\tcudaMalloc Downwards:\t\t%3.1f ms\t%3.1f MB/s", elapsedTime, testSizeByte / (elapsedTime / )); elapsedTime = cuda_host_alloc_test(SIZE, true);
printf("\n\tcudaHostAlloc Upwards:\t\t%3.1f ms\t%3.1f MB/s", elapsedTime, testSizeByte / (elapsedTime / ));
elapsedTime = cuda_host_alloc_test(SIZE, false);
printf("\n\tcudaHostAlloc Downwards:\t%3.1f ms\t%3.1f MB/s", elapsedTime, testSizeByte / (elapsedTime / )); getchar();
return;
}
● 程序输出如下图，可见页锁定内存的读取速度要比内存快一些。
● 页锁定内存的使用方法
int *a, *dev_a, size; cudaHostAlloc((void**)&a, sizeof(int) * size, cudaHostAllocDefault);// 申请页锁定内存
cudaMalloc((void**)&dev_a, sizeof(int) * size); cudaMemcpy(dev_a, a, sizeof(int) * size, cudaMemcpyHostToDevice);// 使用普通的内存拷贝
cudaMemcpy(a, dev_a, sizeof(int) * size, cudaMemcpyDeviceToHost); cudaFreeHost(a);// 释放内存
cudaFree(dev_a);
● 单流内存拷贝
#include <stdio.h>
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include "D:\Code\CUDA\book\common\book.h" #define N (1024*1024)
#define LOOP_TIMES (100) __global__ void kernel(int *a, int *b, int *c)
{
int idx = threadIdx.x + blockIdx.x * blockDim.x;
if (idx < N)
{
int idx1 = (idx + ) % ;
int idx2 = (idx + ) % ;
float as = (a[idx] + a[idx1] + a[idx2]) / 3.0f;
float bs = (b[idx] + b[idx1] + b[idx2]) / 3.0f;
c[idx] = (as + bs) / ;
}
} int main(void)
{
cudaEvent_t start, stop;
float elapsedTime;
cudaStream_t stream;
int *host_a, *host_b, *host_c;
int *dev_a, *dev_b, *dev_c; cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaStreamCreate(&stream); cudaMalloc((void**)&dev_a, N * sizeof(int));
cudaMalloc((void**)&dev_b, N * sizeof(int));
cudaMalloc((void**)&dev_c, N * sizeof(int)); cudaHostAlloc((void**)&host_a, N * LOOP_TIMES * sizeof(int), cudaHostAllocDefault);
cudaHostAlloc((void**)&host_b, N * LOOP_TIMES * sizeof(int), cudaHostAllocDefault);
cudaHostAlloc((void**)&host_c, N * LOOP_TIMES * sizeof(int), cudaHostAllocDefault); for (int i = ; i < N * LOOP_TIMES; i++)
{
host_a[i] = rand();
host_b[i] = rand();
} cudaEventRecord(start, );
for (int i = ; i < N * LOOP_TIMES; i += N)
{
cudaMemcpyAsync(dev_a, host_a + i, N * sizeof(int), cudaMemcpyHostToDevice, stream);
cudaMemcpyAsync(dev_b, host_b + i, N * sizeof(int), cudaMemcpyHostToDevice, stream); kernel << <N / , , , stream >> >(dev_a, dev_b, dev_c); cudaMemcpyAsync(host_c + i, dev_c, N * sizeof(int), cudaMemcpyDeviceToHost, stream);
}
cudaStreamSynchronize(stream); cudaEventRecord(stop, ); cudaEventSynchronize(stop);
cudaEventElapsedTime(&elapsedTime, start, stop);
printf("\n\tTest %d times, spending time:\t%3.1f ms\n", LOOP_TIMES, elapsedTime); cudaFree(dev_a);
cudaFree(dev_b);
cudaFree(dev_c);
cudaFreeHost(host_a);
cudaFreeHost(host_b);
cudaFreeHost(host_c);
cudaStreamDestroy(stream);
cudaEventDestroy(start);
cudaEventDestroy(stop);
getchar();
return ;
}
● 程序输出
● 限定流作为内存拷贝工作时要使用函数cudaMemcpyAsync()，其简单声明和使用为：
cudaError_t cudaMemcpyAsync(void *dst, const void *src, size_t count, enum cudaMemcpyKind kind, cudaStream_t stream __dv()); cudaMemcpyAsync(dev_a, host_a, size, cudaMemcpyHostToDevice, stream);
cudaMemcpyAsync(host_a, dev_a, size, cudaMemcpyDeviceToHost, stream);
● 使用流的基本过程
cudaStream_t stream;// 创建流变量
int *host_a;
int *dev_a;
host_a = (int *)malloc(sizeof(int)*N);
cudaMalloc((void**)&dev_a, N * sizeof(int)); cudaMemcpyAsync(dev_a, host_a + i, N * sizeof(int), cudaMemcpyHostToDevice, stream);// 采用异步内存拷贝 kernel << <blocksize, threadsize, stream >> > (dev_a);// 执行核函数，注意标记流编号 cudaMemcpyAsync(host_c + i, dev_c, N * sizeof(int), cudaMemcpyDeviceToHost, stream);/ cudaStreamSynchronize(stream);// 流同步，保证该留内的工作全部完成 free(a);// 释放内存和销毁流变量
cudaFree(dev_a);
cudaStreamDestroy(stream);
● 双流内存拷贝
#include <stdio.h>
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include "D:\Code\CUDA\book\common\book.h" #define N (1024*1024)
#define LOOP_TIMES (100)
#define NAIVE true __global__ void kernel(int *a, int *b, int *c)
{
int idx = threadIdx.x + blockIdx.x * blockDim.x;
if (idx < N)
{
int idx1 = (idx + ) % ;
int idx2 = (idx + ) % ;
float as = (a[idx] + a[idx1] + a[idx2]) / 3.0f;
float bs = (b[idx] + b[idx1] + b[idx2]) / 3.0f;
c[idx] = (as + bs) / ;
}
} int main(void)
{
cudaEvent_t start, stop;
float elapsedTime;
cudaStream_t stream0, stream1;
int *host_a, *host_b, *host_c;
int *dev_a0, *dev_b0, *dev_c0;
int *dev_a1, *dev_b1, *dev_c1; cudaEventCreate(&start);
cudaEventCreate(&stop); cudaStreamCreate(&stream0);
cudaStreamCreate(&stream1); cudaMalloc((void**)&dev_a0,N * sizeof(int));
cudaMalloc((void**)&dev_b0,N * sizeof(int));
cudaMalloc((void**)&dev_c0,N * sizeof(int));
cudaMalloc((void**)&dev_a1,N * sizeof(int));
cudaMalloc((void**)&dev_b1,N * sizeof(int));
cudaMalloc((void**)&dev_c1,N * sizeof(int)); cudaHostAlloc((void**)&host_a, N * LOOP_TIMES * sizeof(int),cudaHostAllocDefault);
cudaHostAlloc((void**)&host_b, N * LOOP_TIMES * sizeof(int),cudaHostAllocDefault);
cudaHostAlloc((void**)&host_c, N * LOOP_TIMES * sizeof(int),cudaHostAllocDefault); for (int i = ; i < N * LOOP_TIMES; i++)
{
host_a[i] = rand();
host_b[i] = rand();
}
printf("\n\tArray initialized!"); cudaEventRecord(start, );
for (int i = ; i < N * LOOP_TIMES; i += N * )// 每次吃两个N的长度
{
#if NAIVE
cudaMemcpyAsync(dev_a0, host_a + i, N * sizeof(int),cudaMemcpyHostToDevice,stream0);// 前半段给dev_a0和dev_b0
cudaMemcpyAsync(dev_b0, host_b + i, N * sizeof(int),cudaMemcpyHostToDevice,stream0); kernel << <N / , , , stream0 >> >(dev_a0, dev_b0, dev_c0); cudaMemcpyAsync(host_c + i, dev_c0, N * sizeof(int),cudaMemcpyDeviceToHost,stream0); cudaMemcpyAsync(dev_a1, host_a + i + N,N * sizeof(int),cudaMemcpyHostToDevice,stream1);//后半段给dev_a1和dev_b1
cudaMemcpyAsync(dev_b1, host_b + i + N,N * sizeof(int),cudaMemcpyHostToDevice,stream1); kernel << <N / , , , stream1 >> >(dev_a1, dev_b1, dev_c1); cudaMemcpyAsync(host_c + i + N, dev_c1, N * sizeof(int),cudaMemcpyDeviceToHost,stream1);
#else
cudaMemcpyAsync(dev_a0, host_a + i, N * sizeof(int), cudaMemcpyHostToDevice, stream0);// 两个半段拷贝任务
cudaMemcpyAsync(dev_a1, host_a + i + N, N * sizeof(int), cudaMemcpyHostToDevice, stream1); cudaMemcpyAsync(dev_b0, host_b + i, N * sizeof(int), cudaMemcpyHostToDevice, stream0);// 两个半段拷贝任务
cudaMemcpyAsync(dev_b1, host_b + i + N, N * sizeof(int), cudaMemcpyHostToDevice, stream1); kernel << <N / , , , stream0 >> >(dev_a0, dev_b0, dev_c0);// 两个计算执行
kernel << <N / , , , stream1 >> >(dev_a1, dev_b1, dev_c1); cudaMemcpyAsync(host_c + i, dev_c0, N * sizeof(int), cudaMemcpyDeviceToHost, stream0);// 两个半段拷贝任务
cudaMemcpyAsync(host_c + i + N, dev_c1, N * sizeof(int), cudaMemcpyDeviceToHost, stream1);
#endif
}
cudaStreamSynchronize(stream0);
cudaStreamSynchronize(stream1); cudaEventRecord(stop, ); cudaEventSynchronize(stop);
cudaEventElapsedTime(&elapsedTime,start, stop);
printf("\n\tTest %d times, spending time:\t%3.1f ms\n", LOOP_TIMES, elapsedTime); cudaFree(dev_a0);
cudaFree(dev_b0);
cudaFree(dev_c0);
cudaFree(dev_a1);
cudaFree(dev_b1);
cudaFree(dev_c1);
cudaFreeHost(host_a);
cudaFreeHost(host_b);
cudaFreeHost(host_c);
cudaStreamDestroy(stream0);
cudaStreamDestroy(stream1);
getchar();
return ;
}
● 书上简单双流版本输出结果如下左图，异步双流版本输出结果如下右图。发现异步以后速度反而下降了。
