标签：HANDLE int 基础 prop 线程 cuda ERROR 内存
一：核函数调用与参数传递
1：设备指针
1）可以将cudaMalloc()分配的指针传递给在设备上执行的函数
2）可以用cudaMalloc()分配的指针在设备上进行内存读写操作
3）可以将设备指针传递给在主机上执行的函数
4）不能在主机代码中使用设备指针对内存进行读写操作
二：设备属性
1：使用
1）当编写支持双精度浮点数的应用程序时需要查询支持该功能的设备
三：共享内存与同步
1：对于共享变量，编译器为每个线程块生成一个副本，只需更具线程块中线程数量分配数量
2：归约算法
每个线程将cache[]中的两个数据想加，再放回cache[]。得到的结果为原始数据的一半，执行log2(threadsPerBlock)个步骤，得到所有cache[]的总和
3：线程块选择
min（32,(N+threadPerBlock)/thredPerBlock）
4:线程发散
某些线程执行一条指令，其他线程不执行。
当__synthreads()位于发散分支中，会使设备一直等待而出错。
四：常量内存
1:常量内存使用：
当所有线程束访问相同的只读数据时，常量内存可以大大提高性能。
常量内存可以广播半个线程束。
2：事件与性能分析
cudaEvent_t start,stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start,0);
//do somthing
cudaEventRecord(stop,0);
cudaEventSynchrnize(stop);//使CUP在某个事件上同步，因为GPU执行异步函数调用时，CPU会执行下一条语句。//运行时阻塞
int time;
cudaEventElapsedTime(&time,start,stop);//统计时间
cudaEventDestroy(start);
cudaEventDestory(stop);
五：纹理内存
1：使用--分配内存，绑定内存，核函数调用，解除绑定，释放内存。
1）将输入变量声明为texture类型引用
texture<DataType,type,readmode> texout;//type--参考系类型（1，2，3维）
2)将变量绑定到内存缓冲区
cudaBindTexture(NULL,texout,data.dev_outSrc,imageSize)
六：流
1：页锁定主机内存--固定内存
操作系统不会将固定内存映射到磁盘且不可分页，始终驻留在物理内存中（可以通过直接内存访问DMA进行主机与GPU之间的复制）
1）仅针对cudaMemcpy()调用中的目标内存或源内存，才使用页锁定内存。且在不使用时以及释放（cudaFreeHost）
2）分配流使用的固定内存
cudaHostAlloc(host_a,N*sizeof(int),cudaHostAllocDefault);
3）只能以异步方式对固定内存进行复制操作
cudaMemcpyAsync() ---- cudaHostAlloc()
异步方式不能保证好久执行操作，但可以保证在下一个操作之前执行完。
2:cuda流
1）设备支持重叠功能--prop.deviceOverlap
例如：支持核函数的同时还支持主机与设备之间的复制
int main()
{
cudaDeviceProp prop;
int whichDevice;
HANDLE_ERROR(cudaGetDevice(&whichDevice));
HANDLE_ERROR(cudaGetDeviceProperties(&prop,whichDevice));
if (!prop.deviceOverlap)
printf("Device not support overlap;\n");
return 0;
}
2)高效的使用多个cuda流
将操作放入队列时采用宽度优先方式：将流之间的操作交叉添加在队列中。
//由于GPU内存远小于主机内存，因此使用分块的方式执行计算。
for(int i = 0;i < FULL_DATA_SIZE;i += 2*N)
{
（1）深度优先方式
HANDLE_ERROR(cudaMemcpyAsync(dev_a0,host_a+i,N*sizeof(int),cudaMemcpyHostToDevice,stream0));
HANDLE_ERROR(cudaMemcpyAsync(dev_b0,host_b+i,N*sizeof(int),cudaMemcpyHostToDevice,stream0));
kernel<<<N/256,256,dev_a0,dev_b0,dev_c0>>>;
HANDLE_ERROR(cudaMemcpyAsync(host_c+i,dev_c0,N*sizeof(int),cudaMemcpyDeviceToHost,stream0));
（2）宽度优先方式
HANDLE_ERROR(cudaMemcpyAsync(dev_a0,host_a+i+N,N*sizeof(int),cudaMemcpyHostToDevice,stream1));
HANDLE_ERROR(cudaMemcpyAsync(dev_b0,host_b+i+N,N*sizeof(int),cudaMemcpyHostToDevice,stream1));
kernel<<<N/256,256,dev_a1,dev_b1,dev_c1>>>;
HANDLE_ERROR(cudaMemcpyAsync(host_c+i+N,dev_c0,N*sizeof(int),cudaMemcpyDeviceToHost,stream1));
}
计算a中三个值与b中三个值的平均值
#define N (1024*1024)
#define FULL_DATA_SIZE (n*1024)
__global__ void kernel(int *a,int *b,int *c)
{
int idx = threadIdx.x + blockDim.x*blockIdx.x;
if(idx < Ｎ)
｛
int idx1 = (idx + 1) % 256;
int idx2 = (idx + 2) % 256;
float as = (a[idx] + a[idx1] + a[idx2]) / 3.0f;
float bs = (b[idx] + b[idx1] + b[idx2]) / 3.0f;
c[idx] = (as + bs) / 2;
｝
}
int main()
{
cudaDeviceProp prop;
int whichDevice;
HANDLE_ERROR(cudaGetDevice(&whichDevice));
HANDLE_ERROR(cudaGetDeviceProperties(&prop,whichDevice));
if (!prop.deviceOverlap)
printf("Device not support overlap;\n");
//创建启动计时器
cudaEvent_t start,stop;
float elapsedtime;
HANDLE_ERROR(cudaEventCreate(&start));
HANDLE_ERROR(cudaEventCreate(&stop));
HANDLE_ERROR(cudaEventRecord(start,0));
//创建流
cudaStream_t stream0,stream1;
HANDLE_ERROR(cudaStreamCreate(&stream0));
HANDLE_ERROR(cudaStreamCreate(&stream1));
return 0;
}
七：零拷贝主机内存
首先应该查询设备：
cudaDeviceProp prop;
int whichDevice;
//查看设备支持零拷贝内存吗
HANDLE_ERROR(cudaGetDevice(&whichDevice));
HANDLE_ERROR(cudaGetDeviceProperties(&prop,whichDevice));
if(prop.cudaMapHostMemory != 1)
{
printf("cannot map memory\n");
return 0;
}
//标识希望设备映射主机内存
HANDLE_ERROR(cudaSetDeviceFlags( cudaDeviceMapHost));
1：获取零拷贝主机内存
cudaHostAlloc() ---- cudaHostAllocMapped
与固定内存有相同的属性，但可以在核函数中直接访问这种内存
cudaHostAlloc((void **)&a,size,cudaHostWriteCombined | cudaHostMapped) //合并式写入
2：将主机指针映射到设备指针
cudaHostGetDevicePointer(&dev_a,a,0);
3：性能
1）集成的GPU
与主机共享系统内存，使用零拷贝内存可提高性能，但是该内存与固定内存性能相同
2）单独的GPU
一次性读写操作可使用零拷贝内存，但主机不会缓存零拷贝内存的内容，多次读写时应该避开。
八：使用多GPU
在需要使用多GPU之前应查询设备数
应为每个GPU分配一个不同的线程来控制。
1：线程函数routine()
routine(&(data[1]));
可在应用程序的默认线程中调用。
2：线程辅助函数和start_routine()
CUTThread thread = start_routine(routine,&(data[0]));
3：线程阻塞函数end_thread(thread)
主应用线程将等待其他线程运行完。
九：可移动的固定内存
对于分配固定内存的线程来说是锁定内存，对于其他的来说是可分页的固定内存。要使所有线程都将这块内存视为固定内存，就需要使用可移动的固定内存机制
1）分配
cudaHostAlloc() --- cudaHostPortable
标签：HANDLE,int,基础,prop,线程,cuda,ERROR,内存 来源： https://www.cnblogs.com/pengtangtang/p/12762830.html
本站声明： 1. iCode9 技术分享网（下文简称本站）提供的所有内容，仅供技术学习、探讨和分享； 2. 关于本站的所有留言、评论、转载及引用，纯属内容发起人的个人观点，与本站观点和立场无关； 3. 关于本站的所有言论和文字，纯属内容发起人的个人观点，与本站观点和立场无关； 4. 本站文章均是网友提供，不完全保证技术分享内容的完整性、准确性、时效性、风险性和版权归属；如您发现该文章侵犯了您的权益，可联系我们第一时间进行删除； 5. 本站为非盈利性的个人网站，所有内容不会用来进行牟利，也不会利用任何形式的广告来间接获益，纯粹是为了广大技术爱好者提供技术内容和技术思想的分享性交流网站。