简单的CUDA程序
#include <stdio.h>
__global__ void helloFromGPU( void ){
printf("hello world int gpus\n");
}
int main(){
printf("hello world for gpus\n");
// 10表示应用了10个线程
helloFromGPU <<<1,10>>>();
// 显示的清除内存中内容
cudaDeviceReset();
return 0;
}
helloFromGPU <<<1,10>>>();进行系统的调用,在的该程序中启动了10个线程。
cudaDeviceReset() 显示的释放和清空当前进行中与当前设备的有关的所有资源。
CUDa程序的编程的主要过程
- 分配对应的GPU内存空间
- 从CPU内存中拷贝数据到GPU内存中
- 调用CUDA内核函数的来完成指定程序的执行
- 将数据的从GPU拷贝到对应的CPU内存中
- 释放GPU内存的空间
内存管理
#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#include <time.h>
#include </usr/local/cuda/include/cuda_runtime.h>
#include </usr/local/cuda/include/cuda_runtime_api.h>
#define CHECK(call){
const cudaError err = call;
if(err != cudaSuccess){
printf("Error: %s:%d\n",__FILE__,__LINE__);
printf("Code: %d,reason:%s\n",err,cudaErrorToString(err));
exit(-1);
}
}
void sumArrayOnHost(float *a, float *b,float *c,const int N){
for (int idx = 0; idx < N; idx++)
{
c[idx] = a[idx] + b[idx];
}
}
void initialData(float *a, int size){
time_t t;
srand((unsigned int)time(&t));
for (int i = 0; i < size; i++)
{
a[i] = (float)(rand()&0XFF)/10.0f;
}
}
__global__ void sumArrayOnDevice(float *a, float *b, float *c){
const int tid = threadIdx.x;
c[tid] =a[tid] + b[tid];
}
int main(int argc, char const *argv[])
{
// 设置对应参数的size
int n_num = 32;
size_t nByte = n_num * sizeof(float);
float *h_a;
float *h_b;
float *host_ref;
float *dev_ref;
// 在主机上分配对应的空间
h_a = (float*)malloc(nByte);
h_b = (float*)malloc(nByte);
host_ref = (float*)malloc(nByte);
dev_ref = (float*)malloc(nByte);
initialData(h_a,n_num);
initialData(h_b,n_num);
memset(host_ref,0,nByte);
memset(dev_ref,0,nByte);
float *d_a;
float *d_b;
float *d_c;
// 分配cuda的内存空间
cudaMalloc((float **)&d_a,nByte );
cudaMalloc((float **)&d_b,nByte );
cudaMalloc((float **)&d_c,nByte );
// 将数据拷贝到对应的cuda空间中
cudaMemcpy(d_a,h_a,nByte,cudaMemcpyHostToDevice);
cudaMemcpy(d_b,h_b,nByte,cudaMemcpyHostToDevice);
dim3 block(n_num);
dim3 grid(n_num/block.x);
sumArrayOnHost(h_a,h_b,h_c,n_num);
sumArrayOnDevice<<<grid,block >>>(d_a,d_b,d_c);
// 会阻塞主机端的线程的运行直达设备所有的任务请求结束
cudaDeviceSynchronize();
cudaMemcpy(dev_ref,d_c,nByte,cudaMemcpyDeviceToHost);
cudaFree(d_a);
cudaFree(d_c);
cudaFree(d_b);
free(h_a);
free(h_b);
free(h_c);
return 0;
}
代码展示了一个完整的cuda程序的运行开始到结束的整个流程。
__global__ void sumArrayOnDevice(float *a, float *b, float *c){
const int tid = threadIdx.x;
c[tid] =a[tid] + b[tid];
}
__global__ void sumArrayOnDevice(float *a, float *b, float *c)是一个核函数,cuda对核函数的具有诸多的限制,其只能访问设备的内存,必须返回的类型为void类型,不支持可变参数
__host__ __device__ cudaError_t cudaMalloc ( void** devPtr, size_t size )
分配指定空间的内存
cudaError_t cudaMemcpy(void *dst, const void *src, size_t count, cudaMemcpyKind kind, cudaStream_t stream = (cudaStream_t)0)
进行将设备和主机之间内容进行赋值
cudaMemcpy(dev_ref,d_c,nByte,cudaMemcpyDeviceToHost);
cudaMemcpy(dev_ref,d_c,nByte,cudaMemcpyHostToDevice);
cudaMemcpy(dev_ref,d_c,nByte,cudaMemcpyHostToHost);
cudaMemcpy(dev_ref,d_c,nByte,cudaMemcpyDeviceToDevice);
// 会阻塞主机端的线程的运行直达设备所有的任务请求结束
cudaDeviceSynchronize();
CUDA流和并发
在网格并发中,多个内核在一个设备上同时执行,这本身具有较高的效率。所有的CUDA操作(包括内核和数据传输)都在一个流中显式或隐式地运行,并且这些流包含对应隐式声明以及显示声明。
如果的没有定义一个对应的流,那么内核函数的启动和数据传输将默认使用空流,如下所示使用为空流。
cudaMemcpy(...,cudaMemcpyHostToDevice);
kernel<<<grid,block>>>(void);
cudaMemcpy(...,cudaMemcpyDeviceToHost);
同时非空流也可以被显式的创建和管理,如果想要重叠不同的CUDA操作,必须使用非空流。
流可以将Memcpy操作也变成异步的:
__host__ cudaError_t cudaStreamCreate(cudaStream_t * pStream)
基于流的异步的内核启动和数据传输支持以下类型的粗粒度并发:
__host__ cudaError_t cudaMemcpyAsync(void * dst, const void * src, size_t count, cudaMemcpyKind kind, cudaStream_t stream __dv (0))
-
重叠主机计算和设备计算
-
重叠主机计算和主机与设备间的数据传输
-
重叠主机与设备间的数据传输和设备计算
-
并发设备计算
-
参数解释: 前头的四个参数与普通的Memcpy相同 第五个参数则是指定使用的流 为空的情况则使用默认流 (标识符为0)
这里需要额外注意的是, 使用异步数据传输时, 必须使用固定(非分页)的主机内存
cudaMallocHost(void **ptr,size_t size);
cudaHostAlloc(void **pHost,size_t size,unsigned int flags);
流的使用
在核函数的启动上,也与使用默认流不同,需要使用额外一个的标识符作为第四个参数
kernel_name <<<grid,block,sharedMemSize,stream>>>(argment list)
流使用结束之后,可以进行手动的销毁
cudaError_t cudaStreamDestory(cudaStream_t stream);
当流中的工作尚未完成时, 销毁失败, cudaStreamDestroy函数将立刻返回.只有当流内的操作全部完成时, 才会正常销毁.
流的查询
因为的所有的CUDA流操作都是异步的,所以CUDA的API提供了两个函数来检查六中所有的操作是否已经完成。
cudaError_t cudaStreamSynchronize(cudaStream_t stream);
cudaError_t cudaStreamQuery(cudaStream_t stream);
-
cudaStreamSynchronize强制阻塞主机,直到在给定流中所有的操作都完成了
-
cuda-StreamQuery会检查流中所有操作是否都已经完成,但在它们完成前不会阻塞主机:
流的优先级:
对于计算能力>=3.5的设备, 支持给流分配优先级:
可以使用此函数查询优先级的允许范围.这个函数的返回值存放在leastPriority和greatestPriority中,分别对应于当前设备的最低和最高优先级
__host__ __cudart_builtin__ cudaDeviceGetStreamPriorityRange(int * leastPriority, int * greatestPriority)
通过该函数设置的对应的优先级别
__host__ cudaError_t cudaStreamCreateWithPriority(cudaStream_t * pStream, unsigned int flags, int priority)
CUDA事件
CUDA中事件本质上是CUDA流中标记,它与该流内操作六中特定点的相关连,使用事件来执行以下的两个基本任务。
同步流的执行,监控设备的进展
只有当一个给定CUDA流中先前的所有操作都执行结束后,记录在该流内的事件才会起作用(即完成)可以理解为:事件就是在流中插入的一个特殊的操作
CUDA事件的声明和创建以及销毁
cudaEvent_t event;
cudaEventCreate( &event);
cudaEventDestroy(event);
当cudaEventDestroy函数被调用时,如果事件尚未起作用,则调用立即返回,当事件被标记完成时自动释放与该事件相关的资源
记录事件 & 计算运行时间
事件在流执行中标记了一个点
它们可以用来检查正在执行的流操作是否已经到达了给定点
cudaEventRecord(cudaEvent_t event, cudaStream_t stream __dv (0))
已经排队进入到CUDA流中的事件可以用于等待或者测试在指定流中先前操作的完成情况。等待的一个事件的会阻塞主机等待的流执行的中间点。
cudaEventSynchronize(cudaEvent_t event)
此操作会阻塞主机线程
其作用类似于上头的cudaStreamSynchronize, 但其可以等待一个流的中间点(事件插入点)
不阻塞的查询版本:
cudaEventQuery(cudaEvent_t event)
类似于上头的cudaStreamQuery
而后还有个函数用来计算两个事件完成时间隔的时间:
cudaEventElapsedTime(float * ms, cudaEvent_t start, cudaEvent_t end)
此函数返回事件启动和停止之间的运行时间,以毫秒为单位, 此时事件通常已经完成
startEvent & endEvent不必在同一个CUDA流中
显示同步:
之前介绍了几种 主机-设备的显示同步:
cudaDeviceSynchronize,cudaStreamSynchronize & cudaEventSynchronize
阻塞主机线程直到设备完成所有先前的任务:
cudaDeviceSynchronize(void)
阻塞主机线程直到指定的流中所有任务完成
cudaStreamSynchronize(cudaStream_t stream);
cudaStreamQuery(cudaStream_t stream);
阻塞主机线程直到指定的事件完成
还有个查询函数
cudaEventSynchronize(cudaEvent_t event);
cudaEventQuery(cudaEvent_t event);
此外, 还有个比较灵活的方法:
cudaStreamWaitEvent(cudaStream_t stream, cudaEvent_t event, unsigned int flags __dv (0))
该函数会指定该stream等待特定的event, 仅有该event触发之后才会启动流该event可以关联到相同或者不同的stream
流回调
回调这个玩意应该不陌生, 所以能很快的猜到,相当于往CUDA流中塞入了自建函数而非CUDA API
即自建函数与CUDA API一样能在CUDA流中排队执行
回调功能十分强大,因为它们是第一个GPU操作的例子,此操作与之前所学的都相反, 是GPU在主机上创建任务
流回调函数是由应用程序提供的一个主机函数,其有特殊的格式:
(由于C++的函数指针都有带参数的, 所以必须格式相同)
void my_callback(cudaStream_t stream, cudaError_t status, void *data) {
//Do something
并在流中使用以下的API函数注册:
__host__ cudaError_t CUDARTAPI cudaStreamAddCallback(cudaStream_t stream, cudaStreamCallback_t callback, void * userData, unsigned int flags)
callback:
主机端指定的回调函数
userData
就是传递给回调函数的参数
flags:
回调函数有俩限制:
从回调函数中不可以调用CUDA的API函数
在回调函数中不可以执行同步
回调函数的俩特点:
每使用cudaStreamAddCallback一次,只执行一次回调回调函数执行时, 阻塞队列中排在其后面的工作,直到回调函数完成