可移动固定内存测试

可移动固定内存测试，项目打包下载

#include "../common/book.h"
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include "device_functions.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;

    //块内线程同步
    __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;
};

unsigned WINAPI routine(void *pvoidData)
//void* routine(void *pvoidData) 
{
    DataStruct  *data = (DataStruct*)pvoidData;
    //device0上已经调用了这个代码，这里是device为非0才调用
    if (data->deviceID != 0) {
        HANDLE_ERROR(cudaSetDevice(data->deviceID));
        //告诉运行时希望在和这个设备上分配零拷贝内存，不用在设定是否为可移动的，因为在device0中已经设定
        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
", 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.
", i);
            return 0;
        }
    }

    float *a, *b;
    HANDLE_ERROR(cudaSetDevice(0));
    HANDLE_ERROR(cudaSetDeviceFlags(cudaDeviceMapHost));
    /*
    在设置了设备0后，设置了分配内存的类型为cudaHostAllocPortable，
    否则只有设备0会将这些分配的内存视为固定内存
    只在device0中设定为可移动的
    */
    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
",
        data[0].returnValue + data[1].returnValue);

    return 0;
}