🍪
cookielau
  • Introduction
  • Machine Learning
    • Distributed
      • Bookmarks
    • NLP
      • Transformers
    • MLC
      • Tensor Program Abstraction
      • End-to-End Module Execution
  • Framework
    • PyTorch
      • Bookmarks
      • Model
      • Shared
      • Miscellaneous
    • Tensorflow
      • Bookmarks
      • Model
      • Shared
      • Miscellaneous
    • CUDA
      • Bookmarks
    • DeepSpeed
    • Bagua
      • Model
      • Optimizer
    • Others
      • Bookmarks
  • About Me
    • 2022-04-28
  • Random Thoughts
  • Archives
    • CPP
      • Bookmarks
      • Container
      • Algorithm
      • FILE CONTROL
      • Virtual Table
      • Assembly
      • Key Words
      • Problems
      • Others
    • JAVA
      • String Container
      • Maps
    • PYTHON
      • Bookmarks
      • Python Tools
        • Batch Rename
        • Combine Excel
        • Excel Oprations
        • Read Write Excel
        • Rotate PDF
      • Library
        • Pandas Notes
        • Numpy Notes
        • Json Notes
      • Spider
        • Selenium Install
        • Selenium Locating
        • Selenium Errors
        • Selenium Basics
      • Django
        • Start Up
      • Others
    • LINUX
      • Installation
      • Cli Tools
      • WSL
      • Bugs
    • JUNIOR2
      • Economics
        • Chapter 0x01 经济管理概述
        • Chapter 0x02 微观市场机制分析
        • Chapter 0x03 生产决策与市场结构
        • Chapter 0x04 宏观经济市场分析
        • Chapter 0x05 管理的职能
        • Chapter 0x06 生产系统结构与战略
        • Chapter 0x0b 投资项目经济评价
        • Chapter 0x0f 投资项目经济评价
      • Computer Network
        • 概述
        • 分层模型
        • 物理层
        • 数据链路层
        • 网络层
        • 传输层
        • 应用层
        • HTTP(s)实验
        • [Practice]
      • Software Engineering
        • Introduction
        • Demand Analysis
        • Task Estimation
        • Presentation
      • Network Security
        • Chapter 0x01 概述
        • Chapter 0x02 密码学
        • Chapter 0x03 公钥体制
        • Chapter 0x04 消息认证
        • Chapter 0x05 密钥管理
        • Chapter 0x06 访问控制
        • Assignments
      • x86 Programming
        • Basic Knowledge
        • Program Design
        • System Interruption
        • Frequently used functions
    • MD&LaTex
      • Markdown
      • LaTex
    • NPM
      • NPM LINK
    • MyBlogs
      • 2020BUAA软工——“停下来,回头看”
      • 2020BUAA软工——“初窥构建之法”
      • 2020BUAA软工——“上手软件工程,PSP初体验!”
      • 2020BUAA软工——“深度评测官”
      • 2020BUAA软工——“并肩作战,平面交点Pro”
    • SC
      • PAC 2022
        • Lectures
      • OpenMP & MPI
        • MPI Overview
        • Message Passing Programming
        • OpenMP Overview
        • Work Sharing Directives
        • Annual Challenge
        • Future Topics in OpenMP
        • Tasks
        • OpenMP & MPI
    • Hardware
      • Nvidia GPU
        • Frequent Error
        • Memory Classification
        • CUDA_7_Streams_Simplify_Concurrency
        • Optimize_Data_Transfers_in_CUDA
        • Overlap_Data_Transfers_in_CUDA
        • Write_Flexible_Kernels_with_Grid-Stride_Loops
        • How_to_Access_Global_Memory_Efficiently
        • Using_Shared_Memory
      • Intel CPU
        • Construction
        • Optimization
        • Compilation
        • OpenMP
    • English
      • Vocab
      • Composition
    • Interview
      • Computer Network
Powered by GitBook
On this page
  • nvprof
  • Pinned Host memory (固定内存,锁页内存)
  • Example 十分标准的一个测试传输带宽的程序
  • Batch Small Data Transfer
  • Example
  • Reference

Was this helpful?

  1. Archives
  2. Hardware
  3. Nvidia GPU

Optimize_Data_Transfers_in_CUDA

PreviousCUDA_7_Streams_Simplify_ConcurrencyNextOverlap_Data_Transfers_in_CUDA

Last updated 5 years ago

Was this helpful?

Optimize Data Transfers in CUDA C/C++

2020 Jan 27th CookieLau

Source:

[toc]

nvprof

运行 cuda 应用时,可以加上 nvprof 来查看每个 device 的指令的运行情况以找到加速的瓶颈。

$ nvprof ./a.out   

----------------------------------------------------  

======== NVPROF is profiling a.out...   
======== Command: a.out  
======== Profiling result:  
Time(%)     Time  Calls      Avg      Min      Max Name   
  50.08 718.11us      1 718.11us 718.11us 718.11us [CUDA memcpy DtoH]  
  49.92 715.94us      1 715.94us 715.94us 715.94us [CUDA memcpy HtoD]

nvprof 有很多 option 运行选项,可以通过 nvprof --help 查看,主要的有: 1. --print-gpu-trace 2. --print-api-trace 3. --events 记录指定的事件,如 warps_launched(Number of warps launched),如有多个用英文逗号分隔,查看全部则接 all;可以用 nvprof --query-events 查看所有可以记录的事件 4. --metrics 记录指定的指标(metrics),如 ipc(Instructions executed per cycle,如有多个用英文逗号分隔,查看全部则接 all;可以用 nvprof --query-metrics 查看所有可以记录的指标

详细信息可查阅官方文档:

Pinned Host memory (固定内存,锁页内存)

所以如果我们需要在 host 和 device 端来回传输数据时,pageable memory 的host端拷贝所带来的开销是不可忽视的,为了减少这部分的开销,我们可以通过 在host端分配数据时直接分配pinned memory。

具体方法为: 1. 创建:cudaMallocHost or cudaHostAlloc 2. 销毁:cudaFreeHost

由于直接分配 pinned memory 是可能失效的,所以我们需要进行 cudaError 的检查:

cudaError_t status = cudaMallocHost((void**)&h_aPinned, bytes);  
if (status != cudaSuccess)  
  printf("Error allocating pinned host memory\n");

Example 十分标准的一个测试传输带宽的程序

#include <stdio.h>  
#include <assert.h>  

// Convenience function for checking CUDA runtime API results
// can be wrapped around any runtime API call. No-op in release builds.
// Cuda 函数调用正确性检查的规范形式
inline
cudaError_t checkCuda(cudaError_t result)
{
#if defined(DEBUG) || defined(_DEBUG)
  if (result != cudaSuccess) {
    fprintf(stderr, "CUDA Runtime Error: %s\n", 
            cudaGetErrorString(result));
    assert(result == cudaSuccess);
  }
#endif
  return result;
}

void profileCopies(float        *h_a, 
                   float        *h_b, 
                   float        *d, 
                   unsigned int  n,
                   char         *desc)
{
  printf("\n%s transfers\n", desc);

  unsigned int bytes = n * sizeof(float);

  // events for timing
  cudaEvent_t startEvent, stopEvent; 

  checkCuda( cudaEventCreate(&startEvent) );
  checkCuda( cudaEventCreate(&stopEvent) );

  checkCuda( cudaEventRecord(startEvent, 0) );
  checkCuda( cudaMemcpy(d, h_a, bytes, cudaMemcpyHostToDevice) );
  checkCuda( cudaEventRecord(stopEvent, 0) );
  checkCuda( cudaEventSynchronize(stopEvent) ); 
  // cudaEventSync 要让CPU在此处等待直到memcpy完成,否则CPU会继续执行下面的代码

  float time;
  checkCuda( cudaEventElapsedTime(&time, startEvent, stopEvent) );
  printf("  Host to Device bandwidth (GB/s): %f\n", bytes * 1e-6 / time);

  checkCuda( cudaEventRecord(startEvent, 0) );
  checkCuda( cudaMemcpy(h_b, d, bytes, cudaMemcpyDeviceToHost) );
  checkCuda( cudaEventRecord(stopEvent, 0) );
  checkCuda( cudaEventSynchronize(stopEvent) );

  checkCuda( cudaEventElapsedTime(&time, startEvent, stopEvent) );
  printf("  Device to Host bandwidth (GB/s): %f\n", bytes * 1e-6 / time);
  // 之所以 *1e-6 是因为后面的 time 是 ms,所以转换成为 s 需要 1e3,合起来就是 1e9 = 1G

  for (int i = 0; i < n; ++i) {
    if (h_a[i] != h_b[i]) {
      printf("*** %s transfers failed ***\n", desc);
      break;
    }
  }

  // clean up events
  checkCuda( cudaEventDestroy(startEvent) );
  checkCuda( cudaEventDestroy(stopEvent) );
}

int main()
{
  unsigned int nElements = 4*1024*1024;
  const unsigned int bytes = nElements * sizeof(float);

  // host arrays
  float *h_aPageable, *h_bPageable;   
  float *h_aPinned, *h_bPinned;

  // device array
  float *d_a;

  // allocate and initialize
  h_aPageable = (float*)malloc(bytes);                    // host pageable
  h_bPageable = (float*)malloc(bytes);                    // host pageable
  checkCuda( cudaMallocHost((void**)&h_aPinned, bytes) ); // host pinned
  checkCuda( cudaMallocHost((void**)&h_bPinned, bytes) ); // host pinned
  checkCuda( cudaMalloc((void**)&d_a, bytes) );           // device

  for (int i = 0; i < nElements; ++i) h_aPageable[i] = i;      
  memcpy(h_aPinned, h_aPageable, bytes);
  memset(h_bPageable, 0, bytes);
  memset(h_bPinned, 0, bytes);

  // output device info and transfer size
  cudaDeviceProp prop;
  checkCuda( cudaGetDeviceProperties(&prop, 0) );

  printf("\nDevice: %s\n", prop.name);
  printf("Transfer size (MB): %d\n", bytes / (1024 * 1024));

  // perform copies and report bandwidth
  profileCopies(h_aPageable, h_bPageable, d_a, nElements, "Pageable");
  profileCopies(h_aPinned, h_bPinned, d_a, nElements, "Pinned");

  printf("n");

  // cleanup
  cudaFree(d_a);
  cudaFreeHost(h_aPinned);
  cudaFreeHost(h_bPinned);
  free(h_aPageable);
  free(h_bPageable);

  return 0;
}

原文中作者用了两种 CPU 做测试,在性能较弱的 CPU 上 pageable memory 的带宽是 pinned memory 的一半不到,而在好的 CPU 上测试时两者带宽近似。

注意: 这里不是无脑推荐只要在 host 上分配内存都用 pinned memory,因为大量的 pinned memory 的分配会导致整体性能的下降,因为连续分配固定内存会使得其他应用程序的可以使用的内存减少,导致性能的降低。具体什么时候分配需要经过自己的测试得到。

Batch Small Data Transfer

建议通过每次传输数据多一些来 尽量减少传输的次数。这就很 是具体情况而定了。

  • host​ cudaError_t cudaMallocPitch(void* devPtr, size_t pitch, size_t widthInBytes, size_t height)

    • devPtr:新开辟的矩阵的头指针

    • pitch:分配存储器的宽度,以字节为单位 (通常是 256B or 512B)

    • width:分配矩阵的列数

    • height:分配矩阵的行数

cudaMallocPitch 可以保证,虽然可能存在内存空间的浪费,但是访存的时候对于每一行数据的读取是对齐的。

这里为什么要单独返回 Pitch 呢?是因为后面的高阶拷贝要用到,用来计算2D数据中的某一个元素的地址: 对于一个给定类型的数组,Row 行 Column 列 的元素的地址为:

T* pElement = (T*)((char*)BaseAddress + Row * pitch) + Column;
  • host​ cudaError_t cudaMemcpy2D(void dst, size_t dpitch, const void src, size_t spitch, size_t width, size_t height, enum cudaMemcpyKind kind)

    • dst: 目的矩阵内存头指针

    • dpitch: dst指向的2D数组中的内存宽度,以字节为单位,是cuda为了读取方便,对齐过的内存宽度,可能大于一行元素占据的实际内存

    • src:源矩阵内存头指针

    • spitch: src指向的2D数组中的内存宽度,以字节为单位

    • width: src指向的2D数组中一行元素占据的实际宽度。以字节为单位,等于 width*sizeof(type)

    • height: src指向的2D数组的行数

Example

    size_t width = 6;
    size_t height = 5;
    float *h_data, *d_data;
    size_t pitch;

    // 这里只分配 host 端的内存
    h_data = (float *)malloc(sizeof(float)*width*height);
    for (int i = 0; i < width*height; i++)
        h_data[i] = (float)i;

    // 这里通过 cudaMallocPitch 来分配 device 端的内存,并获得 pitch 值
    printf("\n\tAlloc memory.");
    cudaMallocPitch((void **)&d_data, &pitch, sizeof(float)*width, height);
    printf("\n\tPitch = %d B\n", pitch);

    // 将数据从 host 端拷贝到 device 端
    // 由于 host 上是我们直接连续分配的,所以这个 spitch 就由我们直接写出来
    printf("\n\tCopy to Device.\n");
    cudaMemcpy2D(d_data, pitch, h_data, sizeof(float)*width, sizeof(float)*width, height, cudaMemcpyHostToDevice);

    myKernel << <1, 1 >> > (d_data, height, width, pitch);
    cudaDeviceSynchronize();

    // 将数据从 device 端拷贝回 host 端,此时 dpitch 是我们的 host
    printf("\n\tCopy back to Host.\n");
    cudaMemcpy2D(h_data, sizeof(float)*width, d_data, pitch, sizeof(float)*width, height, cudaMemcpyDeviceToHost);

Reference

在 host 上分配的数据内存可能是可以换页的(pageable),这就有可能会导致数据的缺页或者是脏页,所以当 cuda 需要进行 data transfer 的时候,会进行一个先在 host 上分配 pinned memory,将 pageable memory 上的数据拷贝到 pinned memory 上,再从 pinned memory 拷贝数据到 device 端的操作,如下图所示:

对于 高阶数组而言,为了达到更好的访存效果,我们有专门的内存分配和拷贝函数(这里只讨论到二维矩阵,因为三维及以上的矩阵拷贝时需要自定义额外的大量参数,而且使用也不是非常频繁,不在此讨论,有兴趣可以看 ):

https://devblogs.nvidia.com/how-optimize-data-transfers-cuda-cc/
Link
Link
程序园 | cudaMallocPitch 和 cudaMemcpy2D
博客园 | 二维数组 cudaMallocPitch() 和三维数组 cudaMalloc3D() 的使用
相关函数的cuda文档位置