全局内存

Posted 2021-02-10 jetson-xie

1.一般情况下，应用程序不会在某一个时间点访问任意数据或运行任意代码。应用程序遵循局部性原则。这表明他们可以在任意时间点访问较小的局部内存空间。有两种不同类型的局部性：

• 时间局部性

• 空间局部性

2.时间局部性：认为一个数据在某一时间点被引用，那么这个数据在这个时间点所属的附近时间段内被引用的可能性高，被引用的可能性随着该时间离时间点的距离成反比。

空间局部性：认为一个数据在某一个时间点内被引用， 那么这个数据附近的数据也以有可能在这个时间点附近被引用。

3.随着处理器到内存访问延迟的增加，内存的容量也在相应的增加，常见的层次结构为：

寄存器（最快，但容量最小）→缓存 → 主存 → 磁盘存储器（容量最大，但是速度最慢）

4.磁盘存储器常有的特点：

• 更低的每比特位的平均成本

• 更高的容量

• 更高的延迟

• 更少的处理器访问频率

5.CPU和GPU的主存都采用的是DRAM(动态随机存取存储器)，而低延迟内存 (如CPU一级缓存)则采用SRAM(静态随机存取存储器)。内存层次中最大且最慢的级别通常采用磁盘或者闪存驱动来实现。当数据被处理器频繁使用的时候，该数据就保存在低延迟、低容量的存储器中，而当该数据被存储起来以备用时，数据就存储在高延迟，大容量的存储器中。

6.CPU与GPU在内存层次结构设计上使用相似的准则和模型，他们两者的差别就在于：CUDA编程模型能将内存层次结构更好地呈现给用户，能让我们显式地控制它的行为。

7.对于程序员来说，一般有两种类型的存储器：

• 可编程的：你可以显示的控制哪些数据放在可编程内存中

• 不可编程的：你无法决定要将程序存放在哪个位置中，程序将自动选择存放位置以获得良好的性能。

8.在CPU内存层次中，一级缓存和二级缓存都是不可编程的，但CUDA内存模型提出了多种可编程模型的类型：

• 寄存器

• 共享内存

• 常量内存

• 本地内存

• 纹理内存

• 全局内存

9.一个核函数中的一个线程拥有着自己的本地内存，一个线程块有着自己的共享内存，对同一个线程块中的所有线程都是可见的，其内容持续至整个线程块的运行周期。所有线程都可以访问全局内存，所有线程可以访问的只读内存为纹理内存和常量内存空间。常量内存空间，纹理内存空间和全局内存空间拥有着不同的用途。纹理内存空间为各种不同类型的内存空间提供不同的寻址模式和滤波模式，对于一个程序，常量内存，纹理内存和全局内存拥有着相同的生命周期。

10.寄存器：寄存器是GPU中访问速度最快的内存空间，一个核函数中没有其他修饰符的变量或者一个核函数中有固定常量的数组都会被存放在寄存器中，寄存器变量对于核函数来说是私有的，一个核函数通常使用寄存器来保存频繁访问的线程私有变量，寄存器变量与核函数的生命周期是一样的，当核函数执行完成，那么就再也不能对该寄存器进行访问了。

11.寄存器是一个在SM中由活跃线程束划分出的较少资源。在Fermi架构中限制一个线程中最多只能拥有63个寄存器，在KeplerSM中限制一个线程中最多只能拥有255个寄存器。在核函数中使用的寄存器越少，则每个SM对应的常驻线程块就更多，那么每个SM上的并发线程块越多，使用率和性能就越高。

12.在编译的时候可以通过nvcc编译器输出核函数使用的硬件资源情况

nvcc -arch=sm_61 -Xptxas -v,-abi=no add_big_data_Mat.cu -o add_big_data_Mat

可以输出寄存器的数量，共享内存的字节数和每个线程所使用的常量内存的字节数

如果一个核函数使用了超过硬件限制数量的寄存器，则会用本地内存来代替多占用的寄存器。这种寄存器溢出会给性能带来不利的影响，nvcc编译器会使用启发式策略来最小化寄存器的使用，以避免寄存器的溢出，同时我们也可以使用代码来帮助编译器进行优化:

__global__

void __launch_bounds__(maxThreadsPerBlocks,minBlocksPerMultiprocessor)

kernel(..){

}

其中maxThreadsPerBlocks表示一个线程块中包含的最大线程个数，minBlockPerMultiProcessor是可缺省参数，表示每个SM中最小的常驻线程块个数。

同时你还可以使用编译器选项来规定每个线程的最大寄存器个数

-maxrregcount=32

如果使用了制定的启动边界，则这里的指定的值(32)是无效的。

13.本地内存

核函数中符合进入寄存器中但是不能进入被该核函数分配的寄存器空间中的变量将溢出到本地内存中，编译器放置到本地内存的变量可能有：

• 在编译时使用未知索引引用的本地数组

• 可能会占用大量寄存器空间的较大本地结构体或数组

• 任何不满足核函数寄存器限定条件的变量

溢出到本地内存中的变量实质上与全局内存在同一存储空间中，因此本地内存访问的特点是高延迟和低带宽。

14.共享内存

在核函数中使用

_shared__ 修饰的变量存放在共享内存中。因为共享内存是片上内存，因此在访问时具有高宽带，低延迟的特点，类似于CPU的一级缓存，但是不同的是这块内存属于可编程内存。

15.每个线程块上都有一定数量的线程块分配的共享内存，因此必须非常小心的使用共享内存，否则可能会在不经意间限制活跃线程束的数量。

16.共享内存的生命周期就是一个线程块的生命周期，当一个线程块执行结束后，那么该线程块所对应的共享内存将会被释放并重新分配给新的线程块。

17.共享内存是线程间进行相互通信的基本方式，要想访问共享内存，就必须线程间进行同步，就得使用：void__syncthreads();

该函数可以避免访问数据之间的冲突。

18.SM中的共享内存和一级缓存因为都在同一个64KB的片上，可以动态分配他们两个之间的内存空间量：

cudaError_tcudaFuncSetCacheConfig(constvoid *func,enum cudaFuncCache cacheConfig);

这个函数在每个核函数的基础上进行划分，cacheConfig的类型有四种：

• cudaFuncCachePreferNone 没有参考值

• cudaFuncCachePreferShared 建议48KB的共享内存和16KB的一级缓存

• cudaFuncCachePreferL1 建议16KB的共享内存和32KB的一级缓存

• cudaFuncCachePreferEqual 建议32KB的共享内存和32KB的一级缓存

Fermi架构支持前三种配置，而Kepler架构支持所有配置。

19.常量内存驻留在设备内存中，并在每个SM专用的常量缓存中缓存，常量变量使用__constant__修饰符来修饰，对于所有计算能力的设备而言，常量内存都只能申请64KB，常量内存是静态申明，并且在所有核函数中均可见。

20.核函数只能从常量内存中读取数据，因此常量内存必须在主机端通过

cudaError_tcudaMemcpyToSymbol(constvoid *symbol,constvoid *src,size_t count);

这个函数将src所指的地址下的count个字节内容拷贝到symbol所指的全局内存或者常量内存中。

21.线程束中所有线程从相同的内存地址中读取数据，那么常量内存表现良好。如数学公式中的系数，就可以放在常量内存中，如果线程束中的线程需要对不同的地址进行读取，那么这时候常量内存就不是一个较好的选择。

22.全局内存：全局内存是GPU中容量最大，延迟最高，并且经常被使用的内存空间，global指的是其生命周期和作用域。这个声明表示可以在任何含有SM设备上被访问到，并且生命周期贯穿整个程序。

23.一个全局内存变量可以静态声明或者动态声明。你可以使用

__device__

静态的在设备中声明一个变量。在主机端使用cudaMalloc函数就可以动态声明一个全局内存空间，通过cudaFree来释放，

24.从多个线程访问全局内存时必须注意，由于线程的执行不能跨线程块同步，因此在多个线程对同一个全局地址进行处理可能会出现问题。

25.GPU缓存：与CPU缓存相同，GPU缓存也是不可编程的内存，在GPU上有四种缓存：

• 一级缓存

• 二级缓存

• 只读常量缓存

• 只读纹理缓存

26.每个SM上有一个一级缓存，所有SM共享一个二级缓存。一级缓存和二级缓存都是用来存放局部内存或者全局内存中的数据，也包括寄存器溢出的部分，CUDA允许我们选择读取一级缓存和二级缓存的数据还是单纯的使用读取二级缓存的数据。

27.通过主机端初始化常量内存值的源码

#include<iostream>

usingnamespace std;

#include<cuda_runtime.h>

#include<stdio.h>

__device__float d_var;

__global__voidCHECK_DATA(){

//cout << "data is :"<< d_var;

d_var += 2.0f;

printf("data is:%lf ",d_var);

}

intmain(){

floath_var = 3.14f;

cudaSetDevice(0);

cudaMemcpyToSymbol(d_var,&h_var,sizeof(float));

CHECK_DATA<<<1,1>>>();

cudaMemcpyFromSymbol(&h_var,d_var,sizeof(float));

cout << "d_var:"<<h_var<< endl;

cudaDeviceReset();

return 0;

}

在这个代码中，d_var只是一个符号，因此不能进行取址操作，它在主机端只是表示一个设备上的物理地址。但在核函数中，它表示一个变量。

那么如何在主机端获取得到这个变量对应的地址呢？

cudaError_tcudaGetSymbolAddress(void **devptr,constvoid* symbol);

函数来或的变量地址，

float *ptr;

cudaGetSymbolAddress((void **)&ptr,d_var);

cudaMemcpy(ptr,&h_var,sizeof(float),cudaMemcpyHostToDevice);

具体代码如上

28.在CUDA编程中，你需要对设备和主机端进行操作，一般情况下，核函数无法访问主机端变量，主机端无法访问设备端变量。

29.内存管理：CUDA编程模型内存管理和C语言内存管理一样，需要程序员显式的去管理主t机端和设备端的数据移动以及内存管理:

• 分配和释放设备内存

• 在主机和设备之间传输数据

内存分配和释放：cudaMalloc函数和cudaFree函数来对进行内存的申请和释放，但是初始化操作可以由

cudaError_tcudaMemset(void *devptr,int value,size_t count);

来对分配到的内存进行初始化，由value中所存储的值对devptr所指向的地址的count个字节的数据进行填充。

30.CUDA编程的一个基本原则是尽可能减少主机与设备之间的数据传输。

31.分配的主机内存默认情况下是分配到可分页内存上的，当从可分页主机内存传输到设备上时，CUDA模型首先会分配临时页面锁定的或主机固定的内存，将主机源数据复制到固定内存，再从固定内存复制到设备上。

32.CUDA运行时允许你使用如下指令直接分配固定主机内存：

cudaError_tcudaMallocHost(void **devptr,size_t count);

这个函数可以在主机端分配count字节的锁定页面内存，这些页面锁定内存对于GPU设备来说是可以访问的，因为这些固定内存可以直接被GPU设备访问，所以固定内存相对于GPU设备可以以更高的带宽来读取。但是过多的固定内存可能会降低主机系统的性能，因为它减少了用于可分页内存的数量，其中分页内存对主机是可以访问的。

固定主机内存必须通过

cudaError_tcudaFreeHost(void *ptr);

来释放固定内存。

33.相对于分页内存，固定内存在创建和释放时成本更高，但是它为主机和数据之间进行数据传输提供了高吞吐量。

34.相对于分页内存，使用固定内存获得的加速取决于设备计算能力，例如，当将一个10M的数据传输到Fermi架构上，那么此时将这些数据放置在固定内存上能够获得较好的性能。

35.零拷贝内存：通常来说，主机无法访问设备变量，同时设备也无法访问主机申请的变量，但有个例外：零拷贝内存，主机和设备都可以访问零拷贝内存。

36.GPU线程可以直接访问零拷贝内存，访问零拷贝内存有以下几种优势：

• 当设备内存不足时可以利用主机内存

• 避免了主机和设备之间 的传输造成时间上的浪费

• 提高了PCIe的传输速率

37.零拷贝内存可以通过下面这个函数进行申请，实际上零拷贝内存就是固定内存，通过下面这个函数在设备上进行一个地址映射：

cudaError_tcudaHostAlloc(void **ptr,size_t count,unsignedint flags);

这个函数分配了count 字节的主机内存，该页面是锁定的并且设备可以访问，用这个函数分配的内存也需要使用cudaFreeHost(void *ptr)来进行释放。

flags可以对已经分配的内存的特殊属性进行进一步配置

• cudaHostAllocDefalt

• cudaHostAllocPortable

• cudaHostAllocWriteCombined

• cudaHostAllocMapped

cudaHostAllocDefalt使函数cudaHostAlloc变成与函数cudaMallocHost(分配固定内存)一致，设置cudaHostAllocPortable函数能够返回所有能跟CUDA上下文执行的固定内存，而不仅仅只是CUDA执行的分配内存。cudaHostAllocWriteCombined可以在某些配置上通过PCIe总线更好的传输，但在大多数系统中不能够被读取。零拷贝内存最显著的标志就是cudaHostAllocMapped，该标志可以实现直接主机写入和设备读取被映射到设备地址上的主机内存。

cudaError_tcudaHostGetDevicePointer(void **pDevice,void *pHost,unsignedint flags);

pDevice是主机内存在设备上映射的地址，如果设备不支持映射得到的内存，那么该函数将会失效，flags将留作以后使用，目前需要置零。

38.通过零拷贝内存进行数组相加的源码

#include<iostream>

usingnamespace std;

#include<stdio.h>

#include<cuda_runtime.h>

#include<sys/time.h>

doubleget_sys_time(void){

structtimeval time_t;

gettimeofday(&time_t,NULL);

return ((double)time_t.tv_sec + (double)time_t.tv_usec * 1.e-6);

}

voidCHECK(cudaError_t status){

if(status != cudaSuccess)

{

cout << "CUDA ERROR:"<<cudaGetErrorString(status)<<endl;

exit(1);

}

}

voidinitial_data(float *A,size_t size){

for(size_t i = 0;i < size; i++)

{

A[i] = (float)i;

}

}

__global__voidKernel(float *A,float *B,float *C,int N){

int tid = threadIdx.x + blockIdx.x * blockDim.x;

if(tid < N )

C[tid] = A[tid] * B[tid];

}

voidshow_data(float *A,int size){

for(int i=0;i<size;i++)

{

cout << "data:" << A[i] <<endl;

}

}

intmain(){

float *h_A,*h_B,*h_C;

int size = 32;

int bytes = size * sizeof(float);

h_A = (float *)malloc(bytes);

h_B = (float *)malloc(bytes);

h_C = (float *)malloc(bytes);

unsignedint flags = cudaHostAllocMapped;

CHECK(cudaHostAlloc((void **)&h_A,bytes,flags));

CHECK(cudaHostAlloc((void **)&h_B,bytes,flags));

float *d_A,*d_B,*d_C;

CHECK(cudaHostGetDevicePointer((void **)&d_A,(void *)h_A,0));

CHECK(cudaHostGetDevicePointer((void **)&d_B,(void *)h_B,0));

CHECK(cudaMalloc((float **)&d_C,bytes));

initial_data(h_A,size);

initial_data(h_B,size);

dim3 block(4,1);

dim3 grid((block.x + size - 1)/block.x,1);

Kernel<<<grid,block>>>(d_A,d_B,d_C,size);

CHECK(cudaGetLastError());

cudaDeviceSynchronize();

CHECK(cudaMemcpy(h_C,d_C,bytes,cudaMemcpyDeviceToHost));

show_data(h_C,size);

cudaFreeHost((void *)d_A);

cudaFreeHost((void *)d_B);

cudaFree(d_C);

free(h_A);

free(h_B);

free(h_C);

}

39.如果针对小数据，则零拷贝内存会是个不错的选择，但是针对大数据的话，零拷贝内存不一定占有优势，数据越大，延迟越高。

40.有两种异构计算机系统：集成架构和离散架构

集成架构：CPU和GPU集成在一个芯片上，并且在物理地址上共享主存，因此在这种架构中，无需在PCIe总线上备份，所以使用零拷贝内存可能在性能上会更佳。

离散架构：使用PCIe总线将设备和主机联系在一起的，只有在特定时候才会在零拷贝内存上占有优势。

41.由于映射的固定内存在主机和设备之间是共享的，因此必须保证内存访问要同步，避免数据冲突，数据冲突一般是由于多线程访问同一内存地址造成的。

42.不要过度使用零拷贝内存，因为其延迟性高，所以在零拷贝内存中读取核函数可能会很慢。

43.计算能力2.0及以上版本的设备支持一种特殊的寻址方式，称为统一虚拟寻址（UVA），UVA在CUDA 4.0中被引用，支持64位Linux系统，通过UVA，主机和设备可以共享同一个虚拟地址空间。

44.通过UVA，由cudaAlloc分配的主机固定内存可以同时被设备访问到，因此在传入核函数的时候可以直接传入cudaAlloc返回的指针。在38中代码流程为：

• 分配映射的固定主机内存

• 使用CUDA函数获取映射到固定内存上的设备指针

• 将设备指针传递给核函数

但实际上，通过UVA可以节省使用CUDA函数获取映射到固定内存上的设备指针，直接将固定cudaAlloc函数返回的指针直接传入核函数中也可以得到结果。

45.统一内存寻址：在CUDA6.0中引入了统一内存寻址的概念，统一内存寻址用于简化编程模型中的内存管理，统一内存中创建了一个托管内存池，托管内存池已经分配的空间可以使用相同的内存地址（指针）在CPU和GPU中访问，底层系统自动在统一内存空间中完成CPU和GPU之间的传输。这种数据传输对于应用程序是透明的，这样简化了程序代码。

46.统一内存寻址与UVA（统一虚拟内存寻址）之间的差别是UVA虽然为设备和主机提供了相同的虚拟地址，但是不会实现物理层面的将一个数据从一个位置进行转移到另一个位置，这是统一内存寻址特有的功能。

*47.统一内存寻址提供了一个“单指针到数据”模型，有点类似于零拷贝内存，因为零拷贝内存是由主机端分配的固定内存，因此在访问的时候受到PCIe总线上访问零拷贝内存的影响，核函数的调用将具有较高的延迟。另一方面，统一内存寻址将数据和执行空间分离，因此可以根据需要将数据透明地传输到主机或者设备上，以提高局部性和性能。

48.托管内存：托管内存指的是由底层系统自动分配的统一内存，与特定于设备内存空间的数据可以相互操作，他们都是通过cudaMalloc进行创建，因此可以在核函数中使用以下两种内存：由系统控制的托管内存和由程序明确和调用的未托管内存。所有在设备中的操作也同样适用于托管内存，但与一般内存而言，托管内存同时还可以被主机端所访问和引用。

49.托管内存可以被静态声明也可以被动态声明，可以通过注释

__managed__

来静态声明托管内存，但是这个操作只能是在文件范围内或者全局范围内进行申请。

50.托管内存的动态分配：

cudaError_tcudaMallocManaged(void **devPtr,size_t count,unsignedint flags = 0);

可以通过上面的函数来动态的创建托管内存。这个函数分配了count字节的内存，并用devPtr指针指向该内存地址，该指针在所有设备上都是有效的，与未托管内存相比，使用托管内存的程序可以利用自动数据传输和指针重复消除的功能。

51.内存访问模式：为了在读写数据时达到最佳性能，内存访问操作必须满足一定条件。在
CUDA编程模型中，显著特征之一就是由指令必须以线程束为单位进行发布和执行。在执行内存指令时，线程束中的每个线程都提供了一个正在加载或者存储的内存地址，在线程束中的32个线程中，每个线程都提出了包含请求地址单一内存请求访问。根据线程束中内存地址的分布，内存可以被分成不同的模式。

52.全局内存通过缓存来加载和存储，全局内存是一个逻辑内存空间，所有应用程序最初都存在DRAM上，也就是设备内存中，核函数的请求通常都是在DRAM设备和片上内存间以128字节或者32字节内存事物来实现的。

53.所有对于全局内存的访问都会通过二级缓存，也有许多访问会通过一级缓存，取决于设备的访问类型和GPU架构，如果两级缓存都使用上，则内存访问是由一个128字节的内存事物实现的，如果单纯只使用二级缓存，那么内存访问是通过一个32字节的内存事物来实现的。对于全局缓存架构，可以在编译的时候选择是否使用一级缓存。

54.一行一级缓存是128字节，它映射到设备中是128个字节的对齐段，因此，若是一个线程束中的一个线程访问4个字节，那么每次请求就会获取128个字节的数据，这恰好与缓存行和设备内存段的大小相互契合。

54.在优化应用程序时，需要注意设备内存访问的两个特性：

• 对齐内存访问

• 合并内存访问

55.对齐内存访问：当设备内存事物访问的第一个地址是用于事务服务的缓存粒度的偶数倍时（32字节的二级缓存或者128字节的一级缓存），就会出现对齐内存访问，运行非对齐访问会造成带宽的严重浪费。

合并内存访问：当一个线程束中全部的32个线程访问一个连续内存地址时，那么这种内存访问就称为合并内存访问。

56.当出现非对齐内存访问的时候，这时候可能就需要多个内存事物来进行访问，那么这时候对于一些位置为非粒度偶数倍的位置，访问它的内存事物就要从偶数倍的地方进行读取，这样会对于内存事物中一部分无数据的造成带宽浪费。

57.一般来说，需要优化内存事物效率：用最少的内存事物满足最多的内存访问请求。事物数量和吞吐量的需求根据设备的计算能力变化而变化。

58.全局内存的读取：在SM中，数据通过以下三种路径对全局内存进行缓存/缓冲

• 一级/二级缓存

• 只读缓存

• 常量缓存

59.一/二级缓存是默认路径，要想通过其他缓存通道需要程序来显式说明，全局内存加载是否会通过一级缓存取决于以下两个因素：

• 设备的计算能力

• 编译器选项

60.在Fermi（计算能力2.1）架构和Kepler K40及以后的GPU中，可以通过编译器编译选项选择启用或者禁用全局内存负载的一级缓存，默认情况下在Fermi架构上可以使用一级缓存，在Kepler架构下禁用一级缓存。通过以下编译器命令来禁用一级缓存：

-Xptxas -dlcm=cg

可以通过下面这个命令来启用一级缓存：

-Xptxas -dlcm=ca

当一级缓存被禁用，则程序会将缓存放置在二级缓存中，若无二级缓存，则会选择在DRAM中完成操作。二级缓存每一次内存事物会通过一个、两个、或者四个部分来完成，每个部分含有32个字节。

61.当一级缓存开启时，全局内存加载首先会尝试通过一级缓存，如果一级缓存缺失，则通过二级缓存，如果二级缓存缺失，则会选择通过DRAM来完成，此时的加载求由128个字节的内存事物来完成。

62.在Kepler K10，K20和K20XGPU中，一级缓存不用来加载全局缓存，而是专门用来缓存寄存器溢出到本地内存中的缓存。

63.内存加载可以分为两类：

• 有缓存的加载（启用一级缓存）

• 没有缓存的加载（禁用一级缓存）

64.对齐与非对齐：如果内存访问的第一个地址是32的倍数，那么是对齐加载

合并与非合并：如果线程束访问一个连续的地址，则是加载合并

65.缓存加载经过一级缓存，在粒度为128字节的一级缓存行上由设备内存事物进行传输，缓存加载可以分为对齐/非对齐及合并/非合并加载。

66.CPU与GPU一级缓存加载的区别：CPU一级缓存优化了时间局部性和空间局部性，GPU一级缓存仅仅针对空间局部性，因此频繁的访问一个数据并不会使数据保存在缓存区中。

67.没有缓存的加载不经过一级缓存，因此它的内存段粒度是32字节，而非缓存池粒度128字节，这是更细粒度的加载，可以提高总线利用率。

68.全局加载效率=请求的全局内存加载吞吐量/所需的全局内存加载吞吐量

69.没有缓存的加载的整体性能要差于有缓存的加载的整体性能，如果有一级缓存，一个非对齐访问可能将数据放到一级缓存中，这个一级缓存用于后续的非对齐内存访问，如果缺失一级缓存，那么每一次加载都需要调用多个内存事物，且无法应对后续的请求。

70.只读缓存：只读缓存最初是预留给纹理内存使用的，但是在运算能力3.5及以上的设备中，只读缓存也可以作为全局内存的加载代替一级缓存。只读缓存的加载粒度是32字节，对于分散读取来说，这种缓存粒度要优于一级缓存。

71.有两种方式可以指导内存通过制度缓存加载：

• *使用函数_ldg(可能不行)

• 在间接引用的指针上使用修饰符

72.全局内存写入：一级缓存不能在Kepler和Fermi架构上进行存储操作，数据在到内存上存储之前只经过二级缓存，存储操作由粒度是32内存事物执行，组成可以是一段，两段或者是四段。

73.当存储操作为64个字节在一个128字节内部但不连续，那么这时候会使用一个四段事物来完成存储操作，也就是说，一个四段事物比两个一段事物执行效果好。

74.结构体数组（AOS）和数组结构体（SOA）：许多并行编程，更倾向于SOA（结构体数组），即结构体中含有数组的，在CUDA C编程中也倾向于SOA。

75.影响设备内存操作性能的主要因素有：

• 有效利用SM和DRAM之间的内存传送，为了提高利用率，因此要保证内存访问模式是对齐和合并的。

• 可通过以下两点实现最大化存储器操作数：

1. 展开

2. 修改核函数启动参数

76.核函数可达到的带宽：在分析核函数时，需要考虑内存延迟，即完成一次独立内存请求的时间，内存带宽，即SM访问设备内存的时间，它以每秒访问的字节数来进行衡量。

77.内存带宽：大多数的核函数对内存带宽较为敏感，也就是说他们收到内存带宽的限制。一般有两种类型的内存带宽：

• 理论带宽

• 有效带宽

理论带宽是当前硬件能实现的绝对最大带宽，有效带宽是核函数实际达到的带宽，它是测量带宽，可以用下列公式计算：

有效带宽（GB/s） = （读字节数 + 写字节数 * 10^9）/运行时间

78.通过设备进行矩阵转置操作源码：

#include <iostream>

usingnamespace std;

#include<cuda_runtime.h>

#include<sys/time.h>

#include<stdio.h>

#define IMG_HEIGHT 480

#define IMG_WIDTH 640

voidCHECK(cudaError_t status){

if(status != cudaSuccess)

{

printf("cudaError:%s ",cudaGetErrorString(status));

exit(1);

}

}

doubleget_sys_time(void){

structtimeval time_t;

gettimeofday(&time_t,NULL);

return ((double)time_t.tv_sec + (double)time_t.tv_usec *1.e-6);

}

__global__voidKernel(int *src,int *dst,int src_rows,int src_cols){

int idx = blockIdx.x * blockDim.x + threadIdx.x;

int idy = blockIdx.y * blockDim.y + threadIdx.y;

if(idx >= src_rows || idy >= src_cols) return ;

//printf("blockIdx : %d,threadIdx : %d,idx:%d ",blockIdx.x,threadIdx.x,idx);

//printf("blockIdy : %d,threadIdy : %d,idy:%d ",blockIdx.y,threadIdx.y,idy);

// printf("idx:%d,idy:%d，src:%d ",idx,idy,src[idx*src_cols+idy]);

dst[idy * src_rows + idx] = src[idx*src_cols + idy];

//dst[idy][idx] = src[idx][idy];

//printf("dst:%d,src:%d ",dst[idy][idx],src[idx][idy]);

//printf("debug.. ");

}

voidshow_data(int *data,int height,int width){

for(int rows =0;rows<height;rows++)

{

for(int cols = 0;cols < width; cols++)

{

printf("%9d",data[rows * width + cols]);

}

cout << endl;

}

}

voidinitial_data(int *src,int height,int width){

for(int rows = 0;rows< height ; rows++)

{

for(int cols = 0;cols < width;cols++)

{

src[rows * width + cols] = rows * cols;

}

}

}

intmain(){

cudaSetDevice(0);

int *src_array,*dst;

doubletime_t;

int size = IMG_HEIGHT * IMG_WIDTH;

int bytes = sizeof(int)*size;

cout << "bytes : "<<bytes << endl;

unsignedint flags = cudaHostAllocMapped;

CHECK(cudaHostAlloc((void **)&src_array,bytes,flags));

CHECK(cudaHostAlloc((void **)&dst,bytes,flags));

initial_data(src_array,IMG_HEIGHT,IMG_WIDTH);

dim3 block(1,32);

dim3 grid((block.x+IMG_HEIGHT-1)/block.x,(block.y + IMG_WIDTH-1)/block.y);

time_t = get_sys_time();

Kernel<<<grid,block>>>(src_array,dst,IMG_HEIGHT,IMG_WIDTH);

cudaDeviceSynchronize();

time_t = get_sys_time() - time_t;

cout << "cost_time:"<<time_t <<"s"<< endl;

//cout << "src:"<<endl;

//show_data(src_array,IMG_HEIGHT,IMG_WIDTH);

//cout << "dst:"<< endl;

//show_data(dst,IMG_WIDTH,IMG_HEIGHT);

cudaFreeHost(src_array);

cudaFreeHost(dst);

}