基础

一般内存的设计:寄存器->缓存->主存->磁盘存储器
GPU内存设计

修饰符	变量名	存储器	作用域	生命周期
float var	寄存器	线程	线程
float var[100]	本地	线程	线程
`__shared__`	float var +	共享	块	块
`__device__`	float var +	全局	全局	应用程序
`__constant__`	float var +	常量	全局	应用程序

全局变量声明`device`

#include <cuda_runtime.h>
#include <stdio.h>__device__ float devData;__global__ void checkGlobalVariable(){printf("Device: the value of the global variable is %f\n", devData);devData += 2.0f;
}int main(void)
{float value = 3.14f;cudaMemcpyToSymbol(devData, &value, sizeof(float));printf("Host:   copied %f to the global variable\n", value);checkGlobalVariable<<<1, 1>>>();cudaMemcpyFromSymbol(&value, devData, sizeof(float));printf("Host:   the value changed by the kernel to %f\n", value);cudaDeviceReset();return EXIT_SUCCESS;
}

cudaMemcpyToSymbol(devData, &value, sizeof(float)); 将host内存拷贝到device
cudaMemcpyFromSymbol(&value, devData, sizeof(float));将device内存拷贝到device

不能使用cudaMemcpy(&devData,&value,siezeof(float),cudaMemcopyHostToDevice);,因为不能使用"&",&devData在device端,而不在host端

通过cudaError_t cudaGetSymbolAddress(void**,devPtr, const void* symbol);获得全局变量的地址:

float *dftr = NULL;
cudaGetSymbolAddress((void**)&dptr, devData);
cudaMemcpy(dptr,&value,sizeof(float), cudaMemcpyHostToDevice);

内存管理

内存分配与释放

分配和释放内存耗时较高
cudaError_t cudaMalloc(void** devPtr, size_t count)来分配全局内存
cudaError_t cudaMemset(void *devPtr, int value, size_t count);将host内存拷贝到device
cuda Error_t cudaFree(void *devPtr);释放内存

内存传输

cudaError_t cudaMemcpy(void *dst, const void *src,size_t count, enum cudaMemcpyKind kind);从host的src赋值count字节到device的dst

#include <cuda_runtime.h>
#include <stdio.h>int main(int argc, char **argv)
{unsigned int isize = 1 << 22;unsigned int nbytes = isize * sizeof(float);cudaDeviceProp deviceProp;cudaGetDeviceProperties(&deviceProp, dev);printf("%s starting at ", argv[0]);printf("device %d: %s memory size %d nbyte %5.2fMB\n", dev,deviceProp.name, isize, nbytes / (1024.0f * 1024.0f));float *h_a = (float *)malloc(nbytes);float *d_a;cudaMalloc((float **)&d_a, nbytes);for(unsigned int i = 0; i < isize; i++) h_a[i] = 0.5f;cudaMemcpy(d_a, h_a, nbytes, cudaMemcpyHostToDevice);cudaMemcpy(h_a, d_a, nbytes, cudaMemcpyDeviceToHost);cudaFree(d_a);free(h_a);cudaDeviceReset();return EXIT_SUCCESS;
}

固定内存

cudaError_t cudaMallocHost(void **devPtr, size_t count);在host分配count字节内存,页面锁定,且device可以访问,比可分页内存有更高的带宽

固定内存的分配和释放成本更高,但是它为大规模数据传输提供了更高的传输吞吐量。

#include <cuda_runtime.h>
#include <stdio.h>int main(int argc, char **argv){// memory sizeunsigned int isize = 1 << 22;unsigned int nbytes = isize * sizeof(float);// host 固定内存(pin memory)float *h_a;cudaMallocHost ((float **)&h_a, nbytes);// 分配设备内存float *d_a;cudaMalloc((float **)&d_a, nbytes);memset(h_a, 0, nbytes);for (int i = 0; i < isize; i++) h_a[i] = 100.10f;// transfer data from the host to the devicecudaMemcpy(d_a, h_a, nbytes, cudaMemcpyHostToDevice);// transfer data from the device to the hostcudaMemcpy(h_a, d_a, nbytes, cudaMemcpyDeviceToHost);// free memorycudaFree(d_a);cudaFreeHost(h_a);//固定内存必须通过cudaFreeHost释放cudaDeviceReset();return EXIT_SUCCESS;
}

零拷贝内存

一般来说host不能访问device内存,device不能访问host内存
host和device都可以访问零拷贝内存
零拷贝内存是固定(不可分页)内存
零拷贝内存作用:
当device内存不足时可以利用host内存
避免主机和设备的显示数据传输
提高PCIe传输率

cuda Error_t cudaHostAlloc(void **pHost, size_t count, unsigned int flags);分配count字节host内存,必须使用cudaFreeHost函数释放

flags	描述
cudaHostAllocDefault	使cudaHostAlloc函数的行为与cudaMallocHost函数一致
cudaHostAllocPortable	返回能被所有CUDA上下文使用的固定内存，而不仅是执行内存分配的那一个
cudaHostAllocWriteCombined	返回写结合内存，该内存可以在某些系统配置上通过PCIe总线上更快地传输
cudaHostAllocMapped	该标志返回，可以实现主机写入和设备读取被映射到设备地址空间中的主机内存

通过cudaError_t cudaHostGetDevicePointer(void **pDevice , void *pHost, unsigned int flags);获得映射到固定内存的设备指针
该指针可以在device上被引用得到固定主机内存.
进行频繁读写的时候,零拷贝性能将显著降低,因为要经过PCIe总线来传递数据.
在集成架构(集显),CPU和GPU集成在一个芯片上,并且在物理地址上共享主存,零拷贝内存性能可能更好
在离散架构(独显),需要经过PCIe总线,零拷贝只有在特殊情况下才有优势

代码

统一寻址

统一虚拟寻址(UVA)

#include <cuda_runtime.h>
#include <stdio.h>void checkResult(float *hostRef, float *gpuRef, const int N){double epsilon = 1.0E-8;for (int i = 0; i < N; i++){if (abs(hostRef[i] - gpuRef[i]) > epsilon){printf("Arrays do not match!\n");printf("host %5.2f gpu %5.2f at current %d\n", hostRef[i],gpuRef[i], i);break;}}return;
}void initialData(float *ip, int size){int i;for (i = 0; i < size; i++){ip[i] = (float)( rand() & 0xFF ) / 10.0f;}return;
}void sumArraysOnHost(float *A, float *B, float *C, const int N){for (int idx = 0; idx < N; idx++){C[idx] = A[idx] + B[idx];}
}__global__ void sumArrays(float *A, float *B, float *C, const int N){int i = blockIdx.x * blockDim.x + threadIdx.x;if (i < N) C[i] = A[i] + B[i];
}__global__ void sumArraysZeroCopy(float *A, float *B, float *C, const int N){int i = blockIdx.x * blockDim.x + threadIdx.x;if (i < N) C[i] = A[i] + B[i];
}int main(int argc, char **argv){// set up data size of vectorsint ipower = 10;int nElem = 1 << ipower;size_t nBytes = nElem * sizeof(float);if (ipower < 18){printf("Vector size %d power %d  nbytes  %3.0f KB\n", nElem, ipower,(float)nBytes / (1024.0f));}else{printf("Vector size %d power %d  nbytes  %3.0f MB\n", nElem, ipower,(float)nBytes / (1024.0f * 1024.0f));}// part 1: using device memory// malloc host memoryfloat *h_A, *h_B, *hostRef, *gpuRef;h_A     = (float *)malloc(nBytes);h_B     = (float *)malloc(nBytes);hostRef = (float *)malloc(nBytes);gpuRef  = (float *)malloc(nBytes);// initialize data at host sideinitialData(h_A, nElem);initialData(h_B, nElem);memset(hostRef, 0, nBytes);memset(gpuRef,  0, nBytes);// add vector at host side for result checkssumArraysOnHost(h_A, h_B, hostRef, nElem);// malloc device global memoryfloat *d_A, *d_B, *d_C;cudaMalloc((float**)&d_A, nBytes);cudaMalloc((float**)&d_B, nBytes);cudaMalloc((float**)&d_C, nBytes);// transfer data from host to devicecudaMemcpy(d_A, h_A, nBytes, cudaMemcpyHostToDevice);cudaMemcpy(d_B, h_B, nBytes, cudaMemcpyHostToDevice);// set up execution configurationint iLen = 512;dim3 block (iLen);dim3 grid  ((nElem + block.x - 1) / block.x);sumArrays<<<grid, block>>>(d_A, d_B, d_C, nElem);// copy kernel result back to host sidecudaMemcpy(gpuRef, d_C, nBytes, cudaMemcpyDeviceToHost);// check device resultscheckResult(hostRef, gpuRef, nElem);// free device global memorycudaFree(d_A);cudaFree(d_B);// free host memoryfree(h_A);free(h_B);// part 2: using zerocopy memory for array A and B// allocate zerocpy memorycudaHostAlloc((void **)&h_A, nBytes, cudaHostAllocMapped);cudaHostAlloc((void **)&h_B, nBytes, cudaHostAllocMapped);// initialize data at host sideinitialData(h_A, nElem);initialData(h_B, nElem);memset(hostRef, 0, nBytes);memset(gpuRef,  0, nBytes);// pass the pointer to devicecudaHostGetDevicePointer((void **)&d_A, (void *)h_A, 0);cudaHostGetDevicePointer((void **)&d_B, (void *)h_B, 0);// add at host side for result checkssumArraysOnHost(h_A, h_B, hostRef, nElem);// execute kernel with zero copy memory//零拷贝的意义在哪里?sumArraysZeroCopy<<<grid, block>>>(d_A, d_B, d_C, nElem);//(h+A,H_B,d_C,nElem)才对吧// copy kernel result back to host sidecudaMemcpy(gpuRef, d_C, nBytes, cudaMemcpyDeviceToHost);// check device resultscheckResult(hostRef, gpuRef, nElem);// free  memorycudaFree(d_C);cudaFreeHost(h_A);cudaFreeHost(h_B);free(hostRef);free(gpuRef);// reset devicecudaDeviceReset();return EXIT_SUCCESS;
}

内存访问模式

核函数的内存请求通常是在DRAM设备和片上内存间以128字节或32字节内存事务来实现的。
如果这两级缓存都被用到，那么内存访问是由一个128字节的内存事务实现的。如果只使用了二级缓存，那么这个内存访问是由一个32字节的内存事务实现
一行一级缓存是128个字节，它映射到设备内存中一个128字节的对齐段。如果线程束中的每个线程请求一个4字节的值，那么每次请求就会获取128字节的数据，这恰好与缓存行和设备内存段的大小相契合。
优化应用程序时，你需要注意设备内存访问的两个特性：对齐内存访问-·合并内存访问
1. 对齐与合并内存访问,warp中所有的thread请求地址都在128B的缓存行范围内,完成内存加载只需要一个128B的事务,总线的使用率为100%

2. 访问时对齐的,引用的地址不是连续的thread ID,而是128B范围内的随机值.warp 中的Thread请求地址,仍然在一个缓存行范围内,所以只需要一个128B的事务来完成这一内存加载操作.总线利用率仍是100%.

3. warp请求32个连续4个字节的非对齐数据元素.在全局内存中warp的thread请求落在2个128字节的范围内.因此有两个128字节的事务来加载内存操作,总线利利用率为50%,并且在两个事务中加载的字节有一半是未使用的.

4. warp所有thread请求相同的地址,因为被引用的字节落在一个缓存范围内,所以只需要请求一个内存事务,但是总线的利用率非常低,如果加载的是4字节的,那么总线利用率就是4/128=3.125%

5. 最坏的情况:warp的thread分散于全局内存中的32个4字节地址.,地址要占用N个缓存(0<N<32),完成一次内存加载需要申请N次内存事务

结构体数组和数组结构体

SoA模式存储数据充分利用了GPU的内存带宽,GPU上SoA布局提供给了合并内存访问,对全局内存实现更高效的利用.

核函数可达到的带宽

CUDA C编程权威指南第四章全局内存相关推荐

CUDA C编程权威指南第五章共享内存和常量内存
共享内存是较小的片上内存,具有较低的延迟(相比全局,低20~30倍),提供更高的带宽(相比全局,10倍) block通信用于全局内存数据的缓存 __shared__来申请共享变量,如果共享内存大小在 ...
CUDA C编程权威指南第六章流和并发
流 cuda流流分为两种类型:1)隐式声明流(空流),2)显式声明流(非空流) cudaError_t cudaMemcpyAsync(void* dst,const void* src, size ...
CUDA C编程权威指南第七章调整指令级原语
将程序分为两类:IO密集型和计算密集型 double value = in1 * in2 + in3 乘法后紧跟加法的模式被称为乘加法,或者MAD 简单的编译器会将一个MAD指令转换为:一个乘法指令和 ...
CUDA C编程权威指南第三章 CUDA执行模型
基础每个GPU有多个SM(streaming multiprocessor) 当启动一个grid时,它的block会被分配给多个SM上执行,一个block一旦被调度到一个SM上,则这个block只会 ...
《CUDA C编程权威指南》——3.4　避免分支分化
本节书摘来自华章计算机<CUDA C编程权威指南>一书中的第3章,第3.4节,作者 [美] 马克斯·格罗斯曼(Max Grossman),译颜成钢殷建李亮,更多章节内容可以访问云栖社 ...
《CUDA C编程权威指南》——1.5节总结
本节书摘来自华章社区<CUDA C编程权威指南>一书中的第1章,第1.5节总结,作者[美] 马克斯·格罗斯曼(Max Grossman) ,更多章节内容可以访问云栖社区"华章社区 ...
c cuda 指定gpu_《CUDA C编程权威指南》——1.3　用GPU输出Hello World-阿里云开发者社区...
本节书摘来自华章计算机<CUDA C编程权威指南>一书中的第1章,第1.3节,作者 [美] 马克斯·格罗斯曼(Max Grossman),译颜成钢殷建李亮,更多章节内容可以访问云栖社 ...
《CUDA C编程权威指南》——2.2　给核函数计时
本节书摘来自华章计算机<CUDA C编程权威指南>一书中的第2章,第2.2节,作者 [美] 马克斯·格罗斯曼(Max Grossman),译颜成钢殷建李亮,更多章节内容可以访问云栖社 ...
《CUDA C编程权威指南》——2.4节设备管理
本节书摘来自华章社区<CUDA C编程权威指南>一书中的第2章,第2.4节设备管理,作者[美] 马克斯·格罗斯曼(Max Grossman) ,更多章节内容可以访问云栖社区"华章 ...

CUDA C编程权威指南第四章全局内存

基础

全局变量声明`device`

内存管理

内存分配与释放

内存传输

固定内存

零拷贝内存

统一寻址

内存访问模式

结构体数组和数组结构体

核函数可达到的带宽

CUDA C编程权威指南第四章全局内存相关推荐

最新文章

热门文章

CUDA C编程权威指南 第四章 全局内存

基础

全局变量声明__device__

内存管理

内存分配与释放

内存传输

固定内存

零拷贝内存

统一寻址

内存访问模式

结构体数组和数组结构体

核函数可达到的带宽

CUDA C编程权威指南 第四章 全局内存相关推荐

最新文章

热门文章

CUDA C编程权威指南第四章全局内存

全局变量声明`device`

CUDA C编程权威指南第四章全局内存相关推荐