cuda runtime-GPU内存模型和数据拷贝

最新推荐文章于 2024-04-08 21:03:03 发布

Attention is all you

最新推荐文章于 2024-04-08 21:03:03 发布

阅读量315

点赞数 1

分类专栏：学习Cuda和Tensorrt的系列文章标签：算法人工智能自动驾驶深度学习 gpu算力

本文链接：https://blog.csdn.net/zwhdldz/article/details/131984812

版权

学习Cuda和Tensorrt的系列专栏收录该内容

8 篇文章 0 订阅

订阅专栏

前言

CUDA Driver 是与 GPU 沟通的驱动级别底层 API，对 Driver API 的理解，有利于理解后续的 Runtime API。CUDA Driver 随显卡驱动发布，需要与 cudatoolkit 分开看，CUDA Driver 对应于 cuda.h 和 libcuda.so。Driver API 主要知识点是 Context 的管理机制以及 CUDA 系列接口的开发习惯（错误检查方法），还有内存模型。

1.内存模型

类型	片上?	作用范围	速度	介绍
Global Memory	不在	所有线程	普通	读写，大小:显存大小
Constant Memory	不在	所有线程	很快	只读，大小:一般64Kb
Shared Memory	在	block内线程	快	读写，大小:2080Ti 48Kb
Local Memory	不在	单个线程	普通	读写，大小:剩余可用显存/sm/sm最大线程数
Register	在	单个线程	最快	读写
Texture Memory	不在	所有线程	快	纹理内存

2.内存实例

测试Local Memory代码如下，其中test_kenel()为Cuda的核函数(具体在后续博客进行介绍)，__isLocal()函数用于返回该类型是否是Local Memory。

// cuda驱动头文件
#include <cuda.h>
// cuda runtime 头文件
#include <cuda_runtime.h>
#include <stdio.h>

using namespace std;

static __global__ void test_kenel(){
    float a=5.0;
    int b[3];
    __shared__ int shared_value;
    printf("test local_memory\n    %s\n", __isLocal(&a)?"is Local Memory":"isnot Local Memory");
    printf("    %s\n", __isLocal(b)?"is Local Memory":"isnot Local Memory");
    printf("    %s\n", __isLocal(&shared_value)?"is Local Memory":"isnot Local Memory");
}

void local_memory(){
    test_kenel<<<1,1>>>();
    cudaDeviceSynchronize();
    return;
}

测试Shared Memory代码如下，代码展示了启动1个block+3个线程。由于Shared Memory在同一个block内是共享的，所以每个线程内打印的结果应该是一样的(除了线程ID)。其中__syncthreads()用于线程同步，需要等所有线程执行到__syncthreads()，才开始执行后面的代码。
Shared Memory还可以使用extern声明外部的动态大小共享内存，大小由核函数的第三个参数指定。

// cuda驱动头文件
#include <cuda.h>
// cuda runtime 头文件
#include <cuda_runtime.h>
#include <stdio.h>

using namespace std;

__shared__ int a;

static __global__ void test_kenel(){
    __shared__ float b;
    __shared__ float c[5];
    extern __shared__ float d[];//使用extern声明外部的动态大小共享内存，大小由核函数的第三个参数指定
    extern __shared__ float e[];//e和d共用一块内存

    if(threadIdx.x==0){
        a = threadIdx.x;
    }
    if(threadIdx.x==1){
        b = threadIdx.x*10;
    }
    if(threadIdx.x==2){
        c[2] = threadIdx.x*100;
        d[2] = threadIdx.x + 100;
    }

    __syncthreads();
    printf("    thread_x = %d\n", threadIdx.x);
    printf("    %s\n", __isLocal(&b)?"is Local Memory":"isnot Local Memory");
    printf("    a = %d  b = %f  c[2] = %f  d[2] = %f  e[2] = %f\n", a,b,c[2],d[2],e[2]);
}

void shared_memory(){
    printf("test shared_memory\n");
    test_kenel<<<1,3,sizeof(float)*5>>>();
    cudaDeviceSynchronize();
    return;
}

测试Global Memory代码如下

// cuda驱动头文件
#include <cuda.h>
// cuda runtime 头文件
#include <cuda_runtime.h>
#include <stdio.h>

using namespace std;

// 方法2: __device__ 定义
__device__ float a[100];

static __global__ void test_kenel(float* device_ptr){
    float b=5.0;
    printf("    %s\n", __isGlobal(a)?"is Global Memory":"isnot Global Memory");
    printf("    %s\n", __isGlobal(&b)?"is Global Memory":"isnot Global Memory");
    printf("    %s\n", __isGlobal(device_ptr)?"is Global Memory":"isnot Global Memory");

}

void global_memory(){
    printf("test global_memory\n");
    // 方法1: 主机分配
    float* device_ptr = nullptr;
    cudaMalloc(&device_ptr,sizeof(float)*100);
    test_kenel<<<1,1>>>(device_ptr);
    cudaDeviceSynchronize();
    return;
}

执行测试代码:


void local_memory();
void shared_memory();
void global_memory();
int main()
{
    // test memory
    local_memory();
    shared_memory();
    global_memory();
    return 0;
}

测试结果如下:

test local_memory
is Local Memory
is Local Memory
isnot Local Memory
test shared_memory
thread_x = 0
thread_x = 1
thread_x = 2
isnot Local Memory
isnot Local Memory
isnot Local Memory
a = 0 b = 10.000000 c[2] = 200.000000 d[2] = 102.000000 e[2] = 102.000000
a = 0 b = 10.000000 c[2] = 200.000000 d[2] = 102.000000 e[2] = 102.000000
a = 0 b = 10.000000 c[2] = 200.000000 d[2] = 102.000000 e[2] = 102.000000
test global_memory
is Global Memory
isnot Global Memory
is Global Memory

3.内存数据拷贝和锁页内存

CPU到GPU之间的传输是PCIe，速度较慢8GB/s，而GPU内存间传输是GDDR5，速度很快约144GB/s，因为尽量避免CPU和GPU之间的数据传输。
对于整个 Host Memory 内存条而言，操作系统区分为两个大类:
Pageable Memory，可分页内存
Page lock Memory(Pinned Memory)，页锁定内存
两者区别如下:
1.Pageable Memory没有锁定特性，对于第三方设备（比如GPU），去访问时，因为无法感知内存是否被交换，可能得不到正确的数据。Pinned Memory具有锁定特性，是稳定不会被交换的。
2.Pageable Memory策略能使用虚拟内存，实际是内存 8GB 但是可以使用外部硬盘的存储空间，以提高程序运行数量。Pinned Memory空间有限(实际是内存 8GB就是8GB)。处理 Host 到 Device 数据传输时用 Pinned Memory 做缓存可以提高性能，但是Pinned Memory使用太多，会导致操作系统整体性能降低。
3.GPU 可以直接访问Pinned Memory(DMA技术)而不能访问Pageable Memory。
4.由 new、malloc 分配的是Pageable Memory，由 cudaMallocHost 分配的是 Pinned Memory（C语言函数分配也行），由 cudaMalloc 分配的是 Global Memory。

不同的 Host Memory 数据传输到 GPU 上的方式不同，具体如下图所示：
在这里插入图片描述
以下代码展示数据从host内存拷贝到GPU的Global Memory，然后拷贝到host的Pinned Memory，实例代码如下:

// cuda驱动头文件
#include <cuda.h>
// cuda runtime 头文件
#include <cuda_runtime.h>
#include <stdio.h>
#include <string.h>

#define checkRuntime(op) __check_cuda_runtime((op), #op, __FILE__, __LINE__)

bool __check_cuda_runtime(cudaError_t code, const char *op, const char *file, int line)
{

    if (code != cudaSuccess)
    {
        const char *err_name = cudaGetErrorName(code);
        const char *err_message = cudaGetErrorString(code);

        printf("%s:%d  %s failed. \n  code = %s, message = %s\n", file, line, op, err_name, err_message);
        return false;
    }
    return true;
}

int main()
{
    int device_id = 0;
    checkRuntime(cudaSetDevice(device_id));

    float *memory_device = nullptr;
    checkRuntime(cudaMalloc(&memory_device, 100 * sizeof(float)));
    printf("memory_device = %p\n", memory_device);

    float *memory_host = new float[100];
    memory_host[2] = 520.25;
    printf("memory_host = %p\n", memory_host);
    printf("memory_host[2] = %f\n", memory_host[2]);
    checkRuntime(cudaMemcpy(memory_device, memory_host, sizeof(float) * 100, cudaMemcpyHostToDevice));
    printf("memory_device = %p\n", memory_device);
    printf("memory_device[2] = %f\n", memory_device[2]);

    float *memory_page_locked = nullptr;
    checkRuntime(cudaMallocHost(&memory_page_locked, 100 * sizeof(float)));
    printf("memory_page_locked = %p\n", memory_page_locked);
    // 同步copy
    checkRuntime(cudaMemcpy(memory_page_locked, memory_device, sizeof(float) * 100, cudaMemcpyDeviceToHost));

    printf("memory_page_locked[2] = %f\n", memory_page_locked[2]);

    checkRuntime(cudaFreeHost(memory_page_locked));
    delete[] memory_host;
    checkRuntime(cudaFree(memory_device));

    return 0;
}

上述代码展示了如下功能，输出结果为

memory_device = 0x7ff396800000
memory_host = 0x5591e0250130
memory_host[2] = 520.250000
memory_device = 0x7ff396800000
memory_page_locked = 0x7ff396a00000
memory_page_locked[2] = 520.250000

Attention is all you

关注

1
点赞
踩
1

收藏

觉得还不错? 一键收藏
打赏
0
评论
cuda runtime-GPU内存模型和数据拷贝

CUDA Driver 是与 GPU 沟通的驱动级别底层 API，对 Driver API 的理解，有利于理解后续的 Runtime API。CUDA Driver 随显卡驱动发布，需要与 cudatoolkit 分开看，CUDA Driver 对应于 cuda.h 和 libcuda.so。Driver API 主要知识点是 Context 的管理机制以及 CUDA 系列接口的开发习惯（错误检查方法），还有内存模型。
复制链接

扫一扫