首页 > 其他分享 >cuda 流

cuda 流

时间:2024-03-31 11:23:02浏览次数:18  
标签:idx int void dev host cuda sizeof

如下图,将多个执行相同核函数的进程通过cuda流来使他们并发执行,提升效率

这很像cpu的流水线

想让下面这个核函数执行两次,每次都是不同的参数
我们需要用到cuda的流来并发的执行提升效率

__global__ void kernel( int *a, int *b, int *c ) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < N) {
        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;
    }
}

实现:

#include<iostream>
#include "cuda_runtime.h"
#include "device_launch_parameters.h"

# define N (1024*1024)
# define FULL_DATA_SIZE (N*20)

__global__ void kernel(int* a, int* b, int* c)
{
	int idx = threadIdx.x + blockIdx.x * blockDim.x;
	if (idx < N)
	{
		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(void)
{
    cudaStream_t    stream0, stream1;
    int* host_a, * host_b, * host_c;
    int* dev_a0, * dev_b0, * dev_c0;
    int* dev_a1, * dev_b1, * dev_c1;


    cudaStreamCreate(&stream0);    //初始化流
    cudaStreamCreate(&stream1);

    cudaMalloc((void**)&dev_a0, N * sizeof(int));
    cudaMalloc((void**)&dev_b0, N * sizeof(int));
    cudaMalloc((void**)&dev_c0, N * sizeof(int));
    cudaMalloc((void**)&dev_a1, N * sizeof(int));
    cudaMalloc((void**)&dev_b1, N * sizeof(int));
    cudaMalloc((void**)&dev_c1, N * sizeof(int));

    //分配一个大小为FULL_DATA_SIZE * sizeof(int)字节的主机内存空间,
    //cudaHostAllocDefault是分配标志,告诉CUDA运行时将主机内存分配为可被GPU直接访问的可锁定内存
    //这意味着数据可以直接从主机内存传输到GPU内存,而不需要中间的数据复制操作。
    cudaHostAlloc((void**)&host_a, FULL_DATA_SIZE * sizeof(int), cudaHostAllocDefault);
    cudaHostAlloc((void**)&host_b, FULL_DATA_SIZE * sizeof(int), cudaHostAllocDefault);
    cudaHostAlloc((void**)&host_c, FULL_DATA_SIZE * sizeof(int), cudaHostAllocDefault);

    //初始化
    for (int i = 0; i < FULL_DATA_SIZE; i++) { 
        host_a[i] = rand();
        host_b[i] = rand();
    }

    for (int i = 0; i < FULL_DATA_SIZE; i += N * 2) 
    {
        //将指定大小的数据从主机内存异步复制到设备内存。
        //由于是异步调用,函数调用将立即返回,而复制操作将在指定的CUDA流stream中异步执行。
        //这允许主机和设备之间的数据传输与其他CUDA操作重叠,从而提高程序性能。
        cudaMemcpyAsync(dev_a0, host_a + i, N * sizeof(int), cudaMemcpyHostToDevice, stream0);
        cudaMemcpyAsync(dev_a1, host_a + i + N, N * sizeof(int), cudaMemcpyHostToDevice, stream1);

        cudaMemcpyAsync(dev_b0, host_b + i, N * sizeof(int), cudaMemcpyHostToDevice, stream0);
        cudaMemcpyAsync(dev_b1, host_b + i + N, N * sizeof(int), cudaMemcpyHostToDevice, stream1);

        kernel << <N / 256, 256, 0, stream0 >> > (dev_a0, dev_b0, dev_c0);
        kernel << <N / 256, 256, 0, stream1 >> > (dev_a1, dev_b1, dev_c1);

        cudaMemcpyAsync(host_c + i, dev_c0, N * sizeof(int), cudaMemcpyDeviceToHost, stream0);
        cudaMemcpyAsync(host_c + i + N, dev_c1, N * sizeof(int), cudaMemcpyDeviceToHost, stream1);
    }

    //流同步 , 等待与指定流相关的所有CUDA操作完成
    cudaStreamSynchronize(stream0);
    cudaStreamSynchronize(stream1);


    //回收内存和流
    cudaFreeHost(host_a);
    cudaFreeHost(host_b);
    cudaFreeHost(host_c);
    cudaFree(dev_a0);
    cudaFree(dev_b0);
    cudaFree(dev_c0);
    cudaFree(dev_a1);
    cudaFree(dev_b1);
    cudaFree(dev_c1);
    cudaStreamDestroy(stream0); 
    cudaStreamDestroy(stream1);

}

标签:idx,int,void,dev,host,cuda,sizeof
From: https://www.cnblogs.com/algoshimo/p/18106513

相关文章

  • Windows安装CUDA 12.1及cudnn
    下载CUDA打开链接(https://developer.nvidia.com/cuda-toolkit-archive)选择 12.1.1 版本 选择Windows->x86_64->10->exe(local)->Download  下载完成后按提示安装到默认路径 下载cudnn点击进入nVidia下载cudnn(https://developer.download.nvidia.com/co......
  • 虚拟环境装torch与cuda
    遇到问题1在python环境中导入torchvision的时候,出现了以下错误ImportError:cannotimportname'PILLOW_VERSION'from'PIL'问题:Pillow包版本过高。解决方法:1.卸载新版本pipuninstallPillow2.安装新版本pipinstallPillow==6.2.2备注:通过conda进行uninstall好像会把......
  • Ubuntu22.04下Issac Gym/宇树机器人RL&gcc/g++,CUDA,CUDA ToolKit,Pytorch配置环境配
    前置条件本随笔写作Condition:在本人3050Ti笔记本上配好环境后,再在室友4060笔记本上边配边记录整理所得。室友的系统已经配好了相应驱动,因此,本随笔内容基于已经安装了NVIDIA显卡驱动的系统。下次搞到没装驱动的系统我再补一个随笔。宇树机器人宇树科技的文档中心有一个简单的安......
  • cuda原子操作
    如果不用原子操作,在进行计算直方图时会发生计算冲突d_b[i]为h_a中数字i有几个下面的代码将h_a全赋值为3,但d_b[3]却为1#include<iostream>#include"cuda_runtime.h"#include"device_launch_parameters.h"#defineN10__global__voidf(int*a,int*b){ intx=blo......
  • agx orin nx 安装 opencv4.9 cuda 版本
    1.卸载原来的CPU版opencvsudoaptpurgelibopencv*sudoaptupdate2.找一个路径,存放下载opencv安装包sudowget-Oopencv-4.9.0.zipwgethttps://github.com/opencv/opencv/archive/4.9.0.zipsudowget-Oopencv_contrib-4.9.0.zipwgethttps://github.com/openc......
  • CUTLASS: Fast Linear Algebra in CUDA C++
    https://developer.nvidia.com/blog/cutlass-linear-algebra-cuda/EfficientMatrixMultiplicationonGPUs计算密集度=(时间复杂度/空间复杂度)=O(N^3)/O(N^2)=O(N)//naivefor(inti=0;i<M;++i)for(intj=0;j<N;++j)for(intk=0;k<......
  • wsl2 ubuntu子系统安装显卡驱动与cuda
    wsl2安装参考文档:http://t.csdnimg.cn/ClwJ9演示安装ubuntu22列出可安装的子系统命令:wsl--list--onlinePSC:\Users\linyu>wsl--list--online以下是可安装的有效分发的列表。使用'wsl.exe--install<Distro>'安装。NAMEF......
  • CUDA编程入门历险记
    说来惭愧,做CUDA相关工作已经有两年多,但是对CUDA编程的研究并不深入,最近想夯实一下基础,于是找了一个教程“cudacrashcourse“。这个教程在B站和youtube上都有:B站:https://www.bilibili.com/video/BV127411G76m?p=1&vd_source=5d09aad9eacf6d90c0a17cf81ee41eefyoutube:https://w......
  • Gaussian Splatting CUDA结构
    给自己记录一下,不一定对。每次去重写的时候又要忘记。1.python部分在gaussian_renderer/__init__.py里面调用cuda写的rasterization,语句为:fromdiff_gaussian_rasterizationimportGaussianRasterizationSettings,GaussianRasterizer其中GaussianRasterizationSettings定义......
  • ubuntu安装cuda和cudnn,并测试tensorflow和pytorch库的与cuda的兼容性(2023年版)
    lspci|grep-invidia查看nvidia设备,看到GPUgcc--version检查是否安装上gcc软件包根据官方文档指示,pipinstalltorch==1.13.1+cu117-fhttps://download.pytorch.org/whl/torch_stable.html,pipinstalltorchaudio==0.13.1+cu117-fhttps://download.pytorch.org/whl/torch......