cuda学习笔记(6)

前面基础部分说cpu内存跟gpu内存在各自的函数中不能相互操作。申请主机内存用malloc，释放用free。但其实cuda可以通过DMA(Direct Memory Access)把主机内存地址直接映射到cuda地址，称为锁页内存（pinned memory or page locked memory）

锁页内存

主机端存在虚拟内存，主机内存不足是会将内存数据交换到虚拟内存中，虚拟内存就是主机中的磁盘空间，需要该页时再重新从磁盘加载回来。这样做可以使用比实际内存更大的内存空间。

锁页内存允许GPU上的MDA控制器在使用主机内存时不用CPU参与。GPU上的显存都是锁页的，因为GPU上的内存时不支持交换到磁盘的。锁页内存就是分配主机内存时锁定该页，让其不与磁盘交换。

CUDA中锁页内存的使用可以使用CUDA驱动API（ driver API’s）cuMemAllocHost()或者使用CUDA的运行时API（runtime API）中的cudaMallocHost()，同时释放要用cudaFreeHost()。除此之外还可以直接用主机上Malloc()分配的空间，然后将其注册为锁页内存（使用cudaHostRegister()函数完成注册）。

使用锁页内存的好处有以下几点：

• 设备内存与锁页内存之间的数据传输可以与内核执行并行处理。
• 锁页内存可以映射到设备内存，减少设备与主机的数据传输。
• 在前端总线的主机系统锁页内存与设备内存之间的数据交换会比较快；并且可以是write-combining的，此时带宽会跟大。

在多线程的不同线程里调用kernel函数，就是cuda多线程。 如果要所有的线程都可以使用锁页内存的好处，需要在分配时将cudaHostAllocPortable标志传给cudaMallocHost()，或者将cudaHostRegisterPortable标志传给函数cudaHostRegister()

默认情况下锁页内存时可缓存的，可以再使用cudaMallocHost()函数时使用cudaHostAllocWriteCombined标志声明为write-combining的，write-combining内存没有一二级缓存，这样其他的应用可拥有更多的缓存资源。此外write-combining在PCI总线的系统中没有snooped过程，可以获得高达40%的传输加速。但是从主机读取write-combining内存速度很慢，因此应该用于主机端只写的数据。

要将锁页内存映射到设备内存的地址空间还需要在cudaMalloHost()中使用cudaHostAllocMapped标志，或者在用cudaHostRegister()函数注册时使用标志cudaHostRegisterMapped。

用来分配一块被映射到设备内存空间的锁页内存。这样的锁页内存会有两个内存地址：主机上的内存地址和设备上的内存地址。主机内存地址直接由函数cudaMallocHost()或Malloc()返回，设备内存地址则由函数cudaHostGetDevicePointer()查询，用以在kernel中访问锁页内存。

多线程、页锁存的点乘示例

#include "../common/book.h"


#define imin(a,b) (a<b?a:b)

#define     N    (33*1024*1024)
const int threadsPerBlock = 256;
const int blocksPerGrid =
            imin( 32, (N/2+threadsPerBlock-1) / threadsPerBlock );


__global__ void dot( int size, float *a, float *b, float *c ) {
    __shared__ float cache[threadsPerBlock];
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    int cacheIndex = threadIdx.x;

    float   temp = 0;
    while (tid < size) {
        temp += a[tid] * b[tid];
        tid += blockDim.x * gridDim.x;
    }
    
    // set the cache values
    cache[cacheIndex] = temp;
    
    // synchronize threads in this block
    __syncthreads();

    // for reductions, threadsPerBlock must be a power of 2
    // because of the following code
    int i = blockDim.x/2;
    while (i != 0) {
        if (cacheIndex < i)
            cache[cacheIndex] += cache[cacheIndex + i];
        __syncthreads();
        i /= 2;
    }

    if (cacheIndex == 0)
        c[blockIdx.x] = cache[0];
}


struct DataStruct {
    int     deviceID;
    int     size;
    int     offset;
    float   *a;
    float   *b;
    float   returnValue;
};


void* routine( void *pvoidData ) {
    DataStruct  *data = (DataStruct*)pvoidData;
    if (data->deviceID != 0) {
        HANDLE_ERROR( cudaSetDevice( data->deviceID ) );
        HANDLE_ERROR( cudaSetDeviceFlags( cudaDeviceMapHost ) );
    }

    int     size = data->size;
    float   *a, *b, c, *partial_c;
    float   *dev_a, *dev_b, *dev_partial_c;

    // allocate memory on the CPU side
    a = data->a;
    b = data->b;
    partial_c = (float*)malloc( blocksPerGrid*sizeof(float) );

    // allocate the memory on the GPU
    HANDLE_ERROR( cudaHostGetDevicePointer( &dev_a, a, 0 ) );
    HANDLE_ERROR( cudaHostGetDevicePointer( &dev_b, b, 0 ) );
    HANDLE_ERROR( cudaMalloc( (void**)&dev_partial_c,
                              blocksPerGrid*sizeof(float) ) );

    // offset 'a' and 'b' to where this GPU is gets it data
    dev_a += data->offset;
    dev_b += data->offset;

    dot<<<blocksPerGrid,threadsPerBlock>>>( size, dev_a, dev_b,
                                            dev_partial_c );
    // copy the array 'c' back from the GPU to the CPU
    HANDLE_ERROR( cudaMemcpy( partial_c, dev_partial_c,
                              blocksPerGrid*sizeof(float),
                              cudaMemcpyDeviceToHost ) );

    // finish up on the CPU side
    c = 0;
    for (int i=0; i<blocksPerGrid; i++) {
        c += partial_c[i];
    }

    HANDLE_ERROR( cudaFree( dev_partial_c ) );

    // free memory on the CPU side
    free( partial_c );

    data->returnValue = c;
    return 0;
}


int main( void ) {
    int deviceCount;
    HANDLE_ERROR( cudaGetDeviceCount( &deviceCount ) );
    if (deviceCount < 2) {
        printf( "We need at least two compute 1.0 or greater "
                "devices, but only found %d\n", deviceCount );
        return 0;
    }

    cudaDeviceProp  prop;
    for (int i=0; i<2; i++) {
        HANDLE_ERROR( cudaGetDeviceProperties( &prop, i ) );
        if (prop.canMapHostMemory != 1) {
            printf( "Device %d can not map memory.\n", i );
            return 0;
        }
    }

    float *a, *b;
    HANDLE_ERROR( cudaSetDevice( 0 ) );
    HANDLE_ERROR( cudaSetDeviceFlags( cudaDeviceMapHost ) );
    HANDLE_ERROR( cudaHostAlloc( (void**)&a, N*sizeof(float),
                              cudaHostAllocWriteCombined |
                              cudaHostAllocPortable |
                              cudaHostAllocMapped ) );
    HANDLE_ERROR( cudaHostAlloc( (void**)&b, N*sizeof(float),
                              cudaHostAllocWriteCombined |
                              cudaHostAllocPortable      |
                              cudaHostAllocMapped ) );

    // fill in the host memory with data
    for (int i=0; i<N; i++) {
        a[i] = i;
        b[i] = i*2;
    }

    // prepare for multithread
    DataStruct  data[2];
    data[0].deviceID = 0;
    data[0].offset = 0;
    data[0].size = N/2;
    data[0].a = a;
    data[0].b = b;

    data[1].deviceID = 1;
    data[1].offset = N/2;
    data[1].size = N/2;
    data[1].a = a;
    data[1].b = b;

    CUTThread   thread = start_thread( routine, &(data[1]) );
    routine( &(data[0]) );
    end_thread( thread );


    // free memory on the CPU side
    HANDLE_ERROR( cudaFreeHost( a ) );
    HANDLE_ERROR( cudaFreeHost( b ) );

    printf( "Value calculated:  %f\n",
            data[0].returnValue + data[1].returnValue );

    return 0;
}