首页 > 编程语言 >CUDA开始的GPU编程 - 第六章:thrust库

CUDA开始的GPU编程 - 第六章:thrust库

时间:2024-11-12 11:44:06浏览次数:3  
标签:int dev host vector CUDA thrust GPU include

第六章:thrust库

使用 CUDA 官方提供的 thrust::universal_vector

虽然自己实现 CudaAllocator 很有趣,也帮助我们理解了底层原理。但是既然 CUDA 官方已经提供了 thrust 库,那就用他们的好啦。

#include <cuda_runtime.h>
#include <thrust/universal_vector.h>  // trusth库

#include <cstdio>

#include "helper_cuda.h"

template <class Func>
__global__ void parallel_for(int n, Func func) {
  for (int i = blockDim.x * blockIdx.x + threadIdx.x; i < n;
       i += blockDim.x * gridDim.x) {
    func(i);
  }
}

int main() {
  int n = 65536;
  float a = 3.14f;
  thrust::universal_vector<float> x(n);  // 使用 thrust库的 universal_vector
  thrust::universal_vector<float> y(n);

  for (int i = 0; i < n; i++) {
    x[i] = std::rand() * (1.f / RAND_MAX);
    y[i] = std::rand() * (1.f / RAND_MAX);
  }

  parallel_for<<<n / 512, 128>>>(n, [a, x = x.data(), y = y.data()] __device__(
                                        int i) { x[i] = a * x[i] + y[i]; });
  checkCudaErrors(cudaDeviceSynchronize());

  for (int i = 0; i < n; i++) {
    printf("x[%d] = %f\n", i, x[i]);
  }

  return 0;
}

universal_vector 是基于统一内存分配的,因此无论在 GPU 还是 CPU 上,数据都可以直接访问。

使用分离的 device_vector 和 host_vector

device_vector 在 GPU 上分配内存,而 host_vector 则在 CPU 上分配内存。

可以通过 = 运算符在 device_vector 和 host_vector 之间拷贝数据,他会自动帮你调用 cudaMemcpy,非常智能。

比如这里的 x_dev = x_host 会将 x_host 中的数据复制到 GPU 上的 x_dev

#include <cuda_runtime.h>
#include <thrust/device_vector.h>
#include <thrust/host_vector.h>

#include <cstdio>

#include "helper_cuda.h"

template <class Func>
__global__ void parallel_for(int n, Func func) {
  for (int i = blockDim.x * blockIdx.x + threadIdx.x; i < n;
       i += blockDim.x * gridDim.x) {
    func(i);
  }
}

int main() {
  int n = 65536;
  float a = 3.14f;
  thrust::host_vector<float> x_host(n);
  thrust::host_vector<float> y_host(n);

  for (int i = 0; i < n; i++) {
    x_host[i] = std::rand() * (1.f / RAND_MAX);
    y_host[i] = std::rand() * (1.f / RAND_MAX);
  }

  thrust::device_vector<float> x_dev = x_host;
  thrust::device_vector<float> y_dev = x_host;

  parallel_for<<<n / 512, 128>>>(
      n, [a, x_dev = x_dev.data(), y_dev = y_dev.data()] __device__(int i) {
        x_dev[i] = a * x_dev[i] + y_dev[i];
      });

  x_host = x_dev;

  for (int i = 0; i < n; i++) {
    printf("x[%d] = %f\n", i, x_host[i]);
  }

  return 0;
}

使用 Thrust 模板函数:thrust::generate

Thrust 提供了与 C++ 标准库类似的模板函数,例如 thrust::generate(b, e, f),它的作用与 std::generate 相似,能够批量生成数据并填充到区间 [b, e) 中。第三个参数是一个函数,这里我们使用了一个 lambda 表达式。

前两个迭代器参数分别是 device_vectorhost_vector 的开始和结束迭代器,可以通过成员函数 begin()end() 获取。第三个参数可以是任意函数,这里用了 lambda 表达式。

#include <cuda_runtime.h>
#include <thrust/device_vector.h>
#include <thrust/generate.h>
#include <thrust/host_vector.h>

#include <cstdio>

#include "helper_cuda.h"

template <class Func>
__global__ void parallel_for(int n, Func func) {
  for (int i = blockDim.x * blockIdx.x + threadIdx.x; i < n;
       i += blockDim.x * gridDim.x) {
    func(i);
  }
}

int main() {
  int n = 65536;
  float a = 3.14f;
  thrust::host_vector<float> x_host(n);
  thrust::host_vector<float> y_host(n);

  auto float_rand = [] { return std::rand() * (1.f / RAND_MAX); };
  thrust::generate(x_host.begin(), x_host.end(), float_rand);
  thrust::generate(y_host.begin(), y_host.end(), float_rand);

  thrust::device_vector<float> x_dev = x_host;
  thrust::device_vector<float> y_dev = x_host;

  parallel_for<<<n / 512, 128>>>(
      n, [a, x_dev = x_dev.data(), y_dev = y_dev.data()] __device__(int i) {
        x_dev[i] = a * x_dev[i] + y_dev[i];
      });

  x_host = x_dev;

  for (int i = 0; i < n; i++) {
    printf("x[%d] = %f\n", i, x_host[i]);
  }

  return 0;
}

使用 Thrust 模板函数:thrust::for_each

同理,thrust::for_each(b, e, f) 对标 C++ 的 std::for_each,用于对区间 [b, e) 中的每个元素 x 调用函数 f(x)。其中,x 实际上是一个引用。如果使用常量迭代器,则是常引用,可以通过 cbegin()cend() 获取常值迭代器。

#include <cuda_runtime.h>
#include <thrust/device_vector.h>
#include <thrust/for_each.h>
#include <thrust/generate.h>
#include <thrust/host_vector.h>

#include <cstdio>

#include "helper_cuda.h"

template <class Func>
__global__ void parallel_for(int n, Func func) {
  for (int i = blockDim.x * blockIdx.x + threadIdx.x; i < n;
       i += blockDim.x * gridDim.x) {
    func(i);
  }
}

int main() {
  int n = 65536;
  float a = 3.14f;
  thrust::host_vector<float> x_host(n);
  thrust::host_vector<float> y_host(n);

  thrust::for_each(x_host.begin(), x_host.end(),
                   [](float &x) { x = std::rand() * (1.f / RAND_MAX); });
  thrust::for_each(y_host.begin(), y_host.end(),
                   [](float &y) { y = std::rand() * (1.f / RAND_MAX); });

  thrust::device_vector<float> x_dev = x_host;
  thrust::device_vector<float> y_dev = x_host;

  thrust::for_each(x_dev.begin(), x_dev.end(),
                   [] __device__(float &x) { x += 100.f; });

  thrust::for_each(x_dev.cbegin(), x_dev.cend(),
                   [] __device__(float const &x) { printf("%f\n", x); });

  parallel_for<<<n / 512, 128>>>(
      n, [a, x_dev = x_dev.data(), y_dev = y_dev.data()] __device__(int i) {
        x_dev[i] = a * x_dev[i] + y_dev[i];
      });

  x_host = x_dev;

  for (int i = 0; i < n; i++) {
    printf("x[%d] = %f\n", i, x_host[i]);
  }

  return 0;
}
  • 当然还有 thrust::reduce,thrust::sort,thrust::find_if,thrust::count_if,thrust::reverse,thrust::inclusive_scan 等。

Thrust 模板函数的特点:自动决定 CPU 或 GPU 执行

Thrust 的 for_each 可以作用于 device_vector 也可以作用于 host_vector。当作用于 host_vector 时,函数会在 CPU 上执行,而作用于 device_vector 时,则会在 GPU 上执行。

在这里插入图片描述

例如,针对 x_host 使用 for_each 时,lambda 表达式不需要修饰,而针对 x_dev 使用时,lambda 表达式需要加上 __device__ 修饰符。

使用 counting_iterator 实现整数区间循环

Thrust 中的迭代器区间操作(如 thrust::for_eachthrust::transformthrust::reduce 等)通常基于迭代器区间。counting_iterator 是一种特殊的迭代器,能够生成递增的整数序列,适合用于这种情况。

这是 Thrust 提供的一种特殊的迭代器,当需要在 Thrust 算法中使用一个递增整数序列时,counting_iterator 就可以作为迭代器传递

  • counting_iterator 实际上是一个生成器,它的作用是 生成一个递增的数值序列

thrust::make_counting_iterator(num) 构建一个计数迭代器,他作为区间表示的就是整数的区间。

#include <cuda_runtime.h>
#include <thrust/device_vector.h>
#include <thrust/for_each.h>
#include <thrust/generate.h>
#include <thrust/host_vector.h>

#include <cstdio>

int main() {
  thrust::for_each(thrust::make_counting_iterator(0),
                   thrust::make_counting_iterator(10),
                   [] __device__(int i) { printf("%d\n", i); });

  return 0;
}
//  0
//  1
//  2
//  3
//  4
//  5
//  6
//  7
//  8
//  9

使用 zip_iterator 合并多个迭代器

zip_iterator 可以看作是一个复合迭代器,它将多个容器(或迭代器)作为输入,生成一个新的迭代器,这个新的迭代器能够同时访问所有输入容器中对应位置的元素。在每次迭代时,zip_iterator 会返回多个容器中相应位置的元素。

可以通过 thrust::make_zip_iterator(a, b) 将多个迭代器合并,就像 Python 中的 zip 函数一样。

在使用时,可以通过 auto const &tup 来捕获每次迭代返回的元组,并使用 thrust::get<index>(tup) 获取其中第 index 个元素。之所以这么处理,是因为 Thrust 需要兼容一些使用较老标准(如 C++03)的程序,尽管现在可以更简洁地使用 C++11 的 std::tuple 和 C++17 的结构绑定语法。

#include <cuda_runtime.h>
#include <thrust/device_vector.h>
#include <thrust/for_each.h>
#include <thrust/generate.h>
#include <thrust/host_vector.h>

#include <cstdio>

#include "helper_cuda.h"

template <class Func>
__global__ void parallel_for(int n, Func func) {
  for (int i = blockDim.x * blockIdx.x + threadIdx.x; i < n;
       i += blockDim.x * gridDim.x) {
    func(i);
  }
}

int main() {
  int n = 65536;
  float a = 3.14f;
  thrust::host_vector<float> x_host(n);
  thrust::host_vector<float> y_host(n);

  auto float_rand = [] { return std::rand() * (1.f / RAND_MAX); };
  thrust::generate(x_host.begin(), x_host.end(), float_rand);
  thrust::generate(y_host.begin(), y_host.end(), float_rand);

  thrust::device_vector<float> x_dev = x_host;
  thrust::device_vector<float> y_dev = x_host;

  thrust::for_each(thrust::make_zip_iterator(x_dev.begin(), y_dev.cbegin()),
                   thrust::make_zip_iterator(x_dev.end(), y_dev.cend()),
                   [a] __device__(auto const &tup) {
                     auto &x = thrust::get<0>(tup);
                     auto const &y = thrust::get<1>(tup);
                     x = a * x + y;
                   });

  x_host = x_dev;

  for (int i = 0; i < n; i++) {
    printf("x[%d] = %f\n", i, x_host[i]);
  }

  return 0;
}

这样,我们能够通过 zip_iterator 同时操作多个容器。

标签:int,dev,host,vector,CUDA,thrust,GPU,include
From: https://blog.csdn.net/lurenze/article/details/143698735

相关文章

  • GPU云服务器的使用场景和功能有哪些?
    摘要:本文将全面介绍GPU云服务器的特点、优势及应用场景。并针对不同的使用需求,给出典型配置方案示例。包括:深度学习、高性能计算、3D渲染、区块链矿机、游戏直播等多种场景,旨在帮助用户深入理解GPU云服务器的功能,并快速上手应用。一、GPU云服务器简介1、GPU云服务器定义GPU......
  • miniconda Pytorch CUDA Cudnn onnxruntime
    FROMubuntu:22.04#docker启动方式#dockerrun-itd--gpusall--privileged=true--shm-size8G--nameonnx197271d29cb79/bin/bashMAINTAINERSuSu#切换阿里云源RUNapt-getupdate&&apt-getinstall-yvim&&apt-getinstall-ysudo&&......
  • GPU OpenGL 管线
    GPUOpenGL管线主要分为以下几个阶段:顶点数据输入:数据定义与准备:开发者定义要渲染的图形的顶点数据,这些数据包含了每个顶点的位置、颜色、纹理坐标、法线向量等信息。例如,对于一个简单的三角形,需要指定三个顶点的三维坐标以及相关属性。这些数据通常存储在内存中,可以通过数组......
  • ECE 4122/6122 CUDA program
    ECE4122/6122Lab4:CUDA-basedJohnConway’sGameofLife(100pts)Category:CUDADue:TuesdayNovember8th,2024by11:59PMObjective:ImplementaC++CUDAprogramtoruntheGameofLife.GameDescription:TheGameofLife(anexampleofacellulara......
  • OpenGL 纹理采样 在GPU中哪个部件完成
    OpenGL纹理采样主要在GPU的流式多处理器(StreamingMultiprocessor,SM)中完成。SM内部包含多个用于执行计算的核心(Core)以及纹理缓存(TextureCache)等部件,这些部件协同工作来实现纹理采样。具体过程如下:纹理数据获取:当需要进行纹理采样时,首先会从纹理内存(通常是显存中的一块区......
  • GPU 架构是图形处理器
    GPU架构是图形处理器(GPU)的内部设计和组织方式,它决定了GPU的性能、功能和效率。以下是GPU架构的一些主要组成部分和相关特点:流处理器(StreamingProcessor)或着色器核心(ShaderCore):这是GPU中最基本的计算单元,负责执行图形渲染和计算任务中的各种计算操作,例如顶点着色、像......
  • GPU不一定是最佳选择
    1GPU确实很快Bepuphysicsv2主要对两个架构细节非常敏感:内存带宽和浮点吞吐量。从配备双通道DDR3内存(如3770K)的四核4宽SIMDCPU到配备AVX2和更高频率DDR4的7700K,可以带来巨大的速度提升。尽管它仍然只是四核,而且从IvyBridge到KabyLake的通用IPC/时钟改进并......
  • 手把手教你搭建Windows+YOLO11+CUDA环境,以EMA注意演示如何改进YOLO11, 训练自定义数据
    YOLOv11目标检测创新改进与实战案例专栏文章目录:YOLOv11创新改进系列及项目实战目录包含卷积,主干注意力,检测头等创新机制以及各种目标检测分割项目实战案例专栏链接:YOLOv11目标检测创新改进与实战案例文章目录YOLOv11目标检测创新改进与实战案例专栏前言本......
  • Ubuntu安装Nvidia驱动与Cuda Toolkit详细教程 | 步骤解析与环境配置 - 幽络源
    步骤概述步骤1:检查是否安装Nvidia驱动步骤2:查询显卡匹配的驱动版本步骤3:安装Nvidia驱动步骤4:检查Nvidia是否安装并加载步骤5:禁用nouveau步骤6:重新加载nvidia驱动并检查步骤7:下载cudatoolkit步骤8:安装cudatoolkit到指定目录步骤9:配置cudatoolkit环境变量步骤10:测试cudatool......
  • GPU 服务器厂家:进博会上的机遇与挑战,AI 发展的强大助力
    2024年的进博会犹如一座璀璨的科技灯塔,照亮了人工智能发展的前行之路。在这场全球瞩目的科技盛宴上,GPU服务器厂家迎来了崭新的机遇与挑战,众多AI相关的前沿产品和技术精彩亮相,让人们深刻领略到GPU服务器在AI发展中举足轻重的重要作用。以AMD为例,其在进博会上展示的......