内存管理

在这里插入图片描述

全局内存数据主要存放的地方
共享内存块内数据同步的地方

在这里插入图片描述
返回一个标记符是否成功

在这里插入图片描述
比普通Malloc快一倍

统一内存分配释放

在这里插入图片描述
少了传输过程

同步拷贝

在这里插入图片描述
必须等拷贝完才会执行下面的

异步

在这里插入图片描述
注意拷贝是否完成，如果直接使用，并不知道拷贝是否完成

共享内存

主要是在程序块中定义
编译的时候就确定了，大小无法改变，程序写死了

内存空间相对较小
在这里插入图片描述

两种使用共享内存的方式

一种是在编译的时候知道的
另一种是在启动的时候知道的

内存使用代码解析

在这里插入图片描述


#include <stdio.h>
#include <cuda.h>

typedef double FLOAT;

__global__ void sum(FLOAT *x)
{
    int tid = threadIdx.x;

    x[tid] += 1; 定义核函数 ，找到索引，给他向量加1
}

int main()
{
    int N = 32;
    int nbytes = N * sizeof(FLOAT);

    FLOAT *dx = NULL, *hx = NULL;
    int i;

    /* allocate GPU mem */
    cudaMalloc((void **)&dx, nbytes);

    if (dx == NULL) {
        printf("couldn't allocate GPU memory\n");
        return -1;
    }

    /* alllocate CPU host mem: memory copy is faster than malloc */
    hx = (FLOAT *)malloc(nbytes);

    if (hx == NULL) {
        printf("couldn't allocate CPU memory\n");
        return -2;
    }

    /* init */
    printf("hx original: \n");
    for (i = 0; i < N; i++) {
        hx[i] = i;

        printf("%g\n", hx[i]);
    }

    /* copy data to GPU */
    cudaMemcpy(dx, hx, nbytes, cudaMemcpyHostToDevice);

    /* call GPU */
    sum<<<1, N>>>(dx);

    /* let GPU finish */
    cudaDeviceSynchronize();

    /* copy data from GPU */
    cudaMemcpy(hx, dx, nbytes, cudaMemcpyDeviceToHost);

    printf("\nhx from GPU: \n");
    for (i = 0; i < N; i++) {
        printf("%g\n", hx[i]);
    }

    cudaFree(dx);
    free(hx);

    return 0;
}

1

这是一个简单的 CUDA 内核函数，用于对一个 FLOAT 类型的数组 x 的每个元素进行加 1 的操作。下面是该代码块的解释：

__global__ void sum(FLOAT *x)
{
    int tid = threadIdx.x;

    x[tid] += 1; 
}

__global__ 定义了一个在设备端运行的函数，在 CUDA 中被称为“内核函数”。__global__ 说明符告诉编译器该函数将在设备上执行，并且是由主机CPU启动的。

sum 是函数名，其中包含一个参数 FLOAT *x，它是一个指向 FLOAT 类型数组的指针，表示输入的数据。

int tid = threadIdx.x 定义了一个整型变量 tid，用于存储当前线程的 ID。threadIdx.x 是一个内置变量，它表示当前线程在块内的唯一编号。对于一维线程块，threadIdx.x 的范围是从 0 到块大小减一。

x[tid] += 1 表示对输入数组中当前线程处理的元素位置的值加一。

这个代码块可以作为一个 CUDA 程序的一部分，与主机端的代码配合使用。当程序启动时，它将首先执行主机端的代码；主机端的代码将把需要处理的数据拷贝到设备端，然后调用内核函数 sum 对其进行处理。处理完成后，主机端的代码将把结果从设备端拷贝回主机端进行处理。该过程实现了CPU和GPU之间协同工作，发挥了二者的优势。

3

cudaDeviceSynchronize() 是一个 CUDA 运行时函数，用于同步所有在当前设备上排队的操作。它会阻塞调用线程，直到设备上的所有任务全部完成。

在 CUDA 程序中，通常需要计算内核函数执行的时间或是检查内核函数是否成功执行完毕。为了确保正确性，需要在内核函数调用之后使用 cudaDeviceSynchronize() 阻塞主机线程，等待设备端任务的完成。只有所有设备端任务都执行完毕，才能继续执行主机端代码。否则，会导致结果不准确或程序错误。

例如，以下代码片段展示了如何使用 cudaDeviceSynchronize() 排除执行内核函数的错误：

kernel<<<grid_size, block_size>>>(input, output);
cudaError_t error = cudaGetLastError();
if (error != cudaSuccess)
{
    printf("CUDA error: %s\n", cudaGetErrorString(error));
    exit(-1);
}
cudaDeviceSynchronize();

首先，使用 <<<...>>> 语法启动了一个内核函数，然后使用 cudaGetLastError() 检查是否有错误发生。如果出现错误，则会打印错误消息并退出程序。最后，使用 cudaDeviceSynchronize() 防止主机端线程继续执行，直到所有设备端任务完成。

总之，cudaDeviceSynchronize() 是一个重要的 CUDA 运行时函数，用于同步主机和设备之间的数据交互以及确保内核函数执行的正确性。

代码

与上面唯一的不同在于

 cudaMallocHost((void **)&hx, nbytes);

速度要快


#include <stdio.h>
#include <cuda.h>

typedef double FLOAT;

__global__ void sum(FLOAT *x)
{
    int tid = threadIdx.x;

    x[tid] += 1;
}

int main()
{
    int N = 32;
    int nbytes = N * sizeof(FLOAT);

    FLOAT *dx = NULL, *hx = NULL;
    int i;

    /* allocate GPU mem */
    cudaMalloc((void **)&dx, nbytes);

    if (dx == NULL) {
        printf("couldn't allocate GPU memory\n");
        return -1;
    }

    /* alllocate CPU host mem: memory copy is faster than malloc */
    cudaMallocHost((void **)&hx, nbytes);

    if (hx == NULL) {
        printf("couldn't allocate CPU memory\n");
        return -2;
    }

    /* init */
    printf("hx original: \n");
    for (i = 0; i < N; i++) {
        hx[i] = i;

        printf("%g\n", hx[i]);
    }

    /* copy data to GPU */
    cudaMemcpy(dx, hx, nbytes, cudaMemcpyHostToDevice);

    /* call GPU */
    sum<<<1, N>>>(dx);

    /* let GPU finish */
    cudaDeviceSynchronize();

    /* copy data from GPU */
    cudaMemcpy(hx, dx, nbytes, cudaMemcpyDeviceToHost);

    printf("\nhx from GPU: \n");
    for (i = 0; i < N; i++) {
        printf("%g\n", hx[i]);
    }

    cudaFree(dx);
    cudaFreeHost(hx);

    return 0;
}