GPU cudaMallocManaged 统一内存的优劣点_蓝鲸123的博客

统一内存(UM)的迁移

分配 UM 时，内存尚未驻留在主机或设备上。主机或设备尝试访问内存时会发生页错误，此时主机或设备会批量迁移所需的数据。同理，当 CPU 或加速系统中的任何 GPU 尝试访问尚未驻留在其上的内存时，会发生页错误并触发迁移。

能够执行页错误并按需迁移内存对于在加速应用程序中简化开发流程大有助益。此外，在处理展示稀疏访问模式的数据时（例如，在应用程序实际运行之前无法得知需要处理的数据时），以及在具有多个 GPU 的加速系统中，数据可能由多个 GPU 设备访问时，按需迁移内存将会带来显著优势。

有些情况下（例如，在运行时之前需要得知数据，以及需要大量连续的内存块时），我们还能有效规避页错误和按需数据迁移所产生的开销。

本实验的后续内容将侧重于对按需迁移的理解，以及如何在分析器输出中识别按需迁移。这些知识可让您在享受按需迁移优势的同时，减少其产生的开销。

统一内存（UM）的页错误

可以使用nsys 工具分析。

enerating CUDA Memory Operation Statistics...
CUDA Memory Operation Statistics (nanoseconds)
Time(%)      Total Time  Operations         Average         Minimum         Maximum  Name                                                                            
-------  --------------  ----------  --------------  --------------  --------------  --------------------------------------------------------------------------------
   78.8        42212544        2304         18321.4            2751          109728  [CUDA Unified Memory memcpy HtoD]                                               
   21.2        11349888         768         14778.5            1791           95136  [CUDA Unified Memory memcpy DtoH]                                               
# 数据迁移
CUDA Memory Operation Statistics (KiB)
              Total      Operations              Average            Minimum              Maximum  Name                                                                            
-------------------  --------------  -------------------  -----------------  -------------------  --------------------------------------------------------------------------------
         393216.000            2304              170.667              4.000             1020.000  [CUDA Unified Memory memcpy HtoD]                                               
         131072.000             768              170.667              4.000             1020.000  [CUDA Unified Memory memcpy DtoH]                                               
 当仅通过CPU访问统一内存时，是否存在内存迁移和/或页面错误的证据？
 没有。
 
 当仅通过GPU访问统一内存时，是否有证据表明内存迁移和/或页面错误？
 没有。
 
 当先由CPU然后由GPU访问统一内存时，是否有证据表明存在内存迁移和/或页面错误？
 有。
 
 当先由GPU然后由CPU访问统一内存时，是否存在内存迁移和/或页面错误的证据？
 有。
  
统一内存管理存在内存的迁移，在kernel执行的时候会降低kernel的执行效率。 
优化：异步内存预取
 
在主机到设备和设备到主机的内存传输过程中，使用一种技术来减少页错误和按需内存迁移成本，此强大技术称为异步内存预取。通过此技术，程序员可以在应用程序代码使用统一内存 (UM) 之前，在后台将其异步迁移至系统中的任何 CPU 或 GPU 设备。此举可以减少页错误和按需数据迁移所带来的成本，并进而提高 GPU 核函数和 CPU 函数的性能。 
此外，预取往往会以更大的数据块来迁移数据，因此其迁移次数要低于按需迁移。此技术非常适用于以下情况：在运行时之前已知数据访问需求且数据访问并未采用稀疏模式。 
CUDA 可通过 cudaMemPrefetchAsync 函数，轻松将托管内存异步预取到 GPU 设备或 CPU。以下所示为如何使用该函数将数据预取到当前处于活动状态的 GPU 设备，然后再预取到 CPU： 
int deviceId;
cudaGetDevice(&deviceId);                                         // The ID of the currently active GPU device.
cudaMemPrefetchAsync(pointerToSomeUMData, size, deviceId);        // Prefetch to GPU device.
cudaMemPrefetchAsync(pointerToSomeUMData, size, cudaCpuDeviceId); // Prefetch to host. `cudaCpuDeviceId` is a
                                                                  // built-in CUDA variable.
#include <stdio.h>
void initWith(float num, float *a, int N)
  for(int i = 0; i < N; ++i)
    a[i] = num;
__global__
void addVectorsInto(float *result, float *a, float *b, int N)
  int index = threadIdx.x + blockIdx.x * blockDim.x;
  int stride = blockDim.x * gridDim.x;
  for(int i = index; i < N; i += stride)
    result[i] = a[i] + b[i];
void checkElementsAre(float target, float *vector, int N)
  for(int i = 0; i < N; i++)
    if(vector[i] != target)
      printf("FAIL: vector[%d] - %0.0f does not equal %0.0f\n", i, vector[i], target);
      exit(1);
  printf("Success! All values calculated correctly.\n");
int main()
  int deviceId;
  int numberOfSMs;
  cudaGetDevice(&deviceId);
  cudaDeviceGetAttribute(&numberOfSMs, cudaDevAttrMultiProcessorCount, deviceId);
  printf("Device ID: %d\tNumber of SMs: %d\n", deviceId, numberOfSMs);
  const int N = 2<<24;
  size_t size = N * sizeof(float);
  float *a;
  float *b;
  float *c;
  cudaMallocManaged(&a, size);
  cudaMallocManaged(&b, size);
  cudaMallocManaged(&c, size);
   * Prefetching can also be used to prevent CPU page faults.
//  将数据预取到CPU
  cudaMemPrefetchAsync(a, size, cudaCpuDeviceId);
  cudaMemPrefetchAsync(b, size, cudaCpuDeviceId);
  cudaMemPrefetchAsync(c, size, cudaCpuDeviceId);
  initWith(3, a, N);
  initWith(4, b, N);
  initWith(0, c, N);
//  将数据预取到GPU
  cudaMemPrefetchAsync(a, size, deviceId);
  cudaMemPrefetchAsync(b, size, deviceId);
  cudaMemPrefetchAsync(c, size, deviceId);
  size_t threadsPerBlock;
  size_t numberOfBlocks;
  threadsPerBlock = 256;
  numberOfBlocks = 32 * numberOfSMs;
  cudaError_t addVectorsErr;
  cudaError_t asyncErr;
  addVectorsInto<<<numberOfBlocks, threadsPerBlock>>>(c, a, b, N);
  addVectorsErr = cudaGetLastError();
  if(addVectorsErr != cudaSuccess) printf("Error: %s\n", cudaGetErrorString(addVectorsErr));
  asyncErr = cudaDeviceSynchronize();
  if(asyncErr != cudaSuccess) printf("Error: %s\n", cudaGetErrorString(asyncErr));
   * Prefetching can also be used to prevent CPU page faults.
//  将数据预取到CPU
  cudaMemPrefetchAsync(c, size, cudaCpuDeviceId);
  checkElementsAre(7, c, N);
  cudaFree(a);
  cudaFree(b);
  cudaFree(c);
在使用异步预取进行了一系列重构之后，您应该看到内存传输次数减少了，但是每次传输的量增加了，并且内核执行时间大大减少了。
                    cudaMallocManaged 分配旨在供主机或设备代码使用的内存，并且现在仍在享受这种方法的便利之处，即在实现自动内存迁移且简化编程的同时，而无需深入了解 cudaMallocManaged 所分配统一内存 (UM) 实际工作原理的详细信息。nsys profile 提供有关加速应用程序中 UM 管理的详细信息，并在利用这些信息的同时结合对 UM 工作原理的更深入理解，进而为优化加速应用程序创造更多机会。统一内存(UM)的迁移分配 UM 时，内存尚未驻留在主机或设备上。主机或设备尝试访问内存时会发
				到底CUDA里最传统的内存拷贝cudaMalloc，和显式地在主机创建内存cudaHostAlloc再开放给GPU访问，以及全局内存寻址cudaMallocManaged，这三种方式，哪种在GPU与CPU之间传输数据时，有更高的效率？
为了回答这个问题，于是有了以下对上述三种方式所创建的内存，在主机与设备之间传输速度的比较代码。......
CUDA8.0，SM6.x之后，使用cudaMalloc默认使用同一内存管理。
注：windows，mac Os 目前只支持基本的统一内存管理属性，不支持SM6.x之后的统一内存管理属性。
可以使用cuda...
				#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include "device_functions.h"
#include <stdio.h>
#include "opencv2/imgproc/imgproc.hpp"
#include "opencv2/imgproc/types_c.h"
				CUDA 6.0以后，出现了：统一内存，unified memory，其具有以下特性： 
1 其声明有以下两种方式： 
a 使用关键字managed，可选的和device,一起使用，比如：device managed int ret[1000]; 
b 使用函数cudaMallocManaged(): 
    int *ret; 
    cudaMallocManaged(&ret,1000*s
				Unified Memory是啥参见其他博主文章，主要就是简化代码，不需要各种cudamalloc和cudamemcpy。
cuda_runtime_api.h中的接口：extern host cudart_builtin cudaError_t CUDARTAPI cudaMallocManaged(void **devPtr, size_t size, unsigned int flags __dv(cudaMemAttachGlobal));
一般传前两个参数即可。
传指针时先在C代码中malloc，
				  CUDA编程的内存管理与C语言的类似，需要程序员显式管理主机和设备之间的数据移动。随着CUDA版本的升级，NVIDIA正系统地实现主机和设备内存空间的统一，但对于大多数应用程序来说，仍需要手动移动数据。对于CUDA内存管理来说，工作重点在于如何使用CUDA函数来显式地管理内存和数据移动，主要是两个方面：分配和释放设备内存；在主机和设备之间传输数据。为了达到最优性能，CUDA提供了在主机端准备设备内存地函数，并且显式地向设备传输数据和从设备中获取数据。
内 存 分 配 和 释 放
  CUDA编程模型假设
				内存管理
CUDA是C语言的扩展，内存方面基本集成了C语言的方式，由程序员控制CUDA内存，当然，这些内存的物理设备是在GPU上的，而且与CPU内存分配不同，CPU内存分配完就完事了，GPU还涉及到数据传输，主机和设备之间的传输。
为达到最优性能，CUDA提供了在主机端准备设备内存的函数，并且显式地向设备传递数据，显式的从设备取回数据。
统一内存Unified Memory
统一内存是可从系统中的任何处理器访问的单个内存地址空间（参见上图）。 这种硬件/软件技术允许应用程序分配可以从 CPU 或 GPU 上
                kubelet Container runtime network not ready“ networkReady=“NetworkReady=false reason:NetworkPluginNo
                    wuzongdefeng: 
                    不知道怎么回事，从k8s官方安装文档和其它地方的安装文档都看了。都没提到和apparmor的联系。没想到出现问题了安装你的操作就可以了，非常感谢
                Ubuntu20.04中安装MySQL 5.7.x
                    我是真的菜啊啊: 
                    而且他示例图片下载的是32位的，命令用的64位的，小白很容易踩坑
                C++ const_cast用法
                    如果是const char*呢