CUDA 学习1

最新推荐文章于 2024-07-04 10:21:11 发布

蓝天巨人

最新推荐文章于 2024-07-04 10:21:11 发布

阅读量124

点赞数

分类专栏： CUDA 学习使用文章标签：学习 c++ c语言

本文链接：https://blog.csdn.net/qq_44632658/article/details/129346107

版权

CUDA 学习使用专栏收录该内容

8 篇文章 3 订阅

订阅专栏

简单的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一次，只执行一次回调回调函数执行时, 阻塞队列中排在其后面的工作，直到回调函数完成