C++ GPU编程(英伟达CUDA)

安装编译环境
https://developer.download.nvidia.com/compute/cuda/12.5.0/local_installers/cuda_12.5.0_555.85_windows.exe

CMakeLists.txt

cmake_minimum_required(VERSION 3.10)set(CMAKE_CXX_STANDARD 17)
set(CMAKE_BUILD_TYPE Release)
#set(CMAKE_CUDA_ARCHITECTURES 52;70;75;86)
set(CMAKE_CUDA_SEPARABLE_COMPILATION ON)project(hellocuda LANGUAGES CXX CUDA)add_executable(main main.cu hello.cu)target_include_directories(main PUBLIC ../../include)
target_compile_options(main PUBLIC $<$<COMPILE_LANGUAGE:CUDA>:--extended-lambda>)
target_compile_options(main PUBLIC $<$<COMPILE_LANGUAGE:CUDA>:--expt-relaxed-constexpr>)
# --use_fast_math   sinf 替换成 __sinf
# --ftz=true 会把极小数(denormal)退化为0
# --prec-div=false 降低除法的精度换取速度。
# --prec-sqrt=false 降低开方的精度换取速度。
# --fmad 因为非常重要，所以默认就是开启的，会自动把 a * b + c 优化成乘加(FMA)指令。
# 开启 --use_fast_math 后会自动开启上述所有

 CudaAllocator.h

template <class T>
struct CudaAllocator {using value_type = T;T* allocate(size_t size) {T* ptr = nullptr;checkCudaErrors(cudaMallocManaged(&ptr, size * sizeof(T)));return ptr;}void deallocate(T* ptr, size_t size = 0) {checkCudaErrors(cudaFree(ptr));}template <class ...Args>void construct(T* p, Args &&...args) {if constexpr (!(sizeof...(Args) == 0 && std::is_pod_v<T>))::new((void*)p) T(std::forward<Args>(args)...);}
};

 hello.cu

#include <cstdio>
#include <cuda_runtime.h>__device__ void say_hello3() {  // 定义printf("Hello, world!\n");
}

 main.cu

#include <cstdio>
#include <cuda_runtime.h>
#include "helper_cuda.h"
#include <vector>
#include "CudaAllocator.h"
#include <thrust/universal_vector.h>
#include <thrust/device_vector.h>
#include <thrust/host_vector.h>
#include <thrust/generate.h>
#include <thrust/for_each.h>__device__ void say_hello3();  // 声明__device__ __inline__ void say_hello() {printf("Hello, world, world from GPU!\n");
}__host__ void say_hello_host() {printf("Hello, world from CPU!\n");
}__host__ __device__ void say_hello2() {
#ifdef __CUDA_ARCH__printf("Hello, world from GPU architecture %d!\n", __CUDA_ARCH__);
#elseprintf("Hello, world from CPU!\n");
#endif
}__global__ void another() {int localVal = 1;printf("another: Thread %d of %d\n", threadIdx.x, blockDim.x);
}__global__ void kernel() {unsigned int tid = blockDim.x * blockIdx.x + threadIdx.x;unsigned int tnum = blockDim.x * gridDim.x;printf("Flattened Thread %d of %d\n", tid, tnum);say_hello2();say_hello3();another << <1, tnum >> > ();  // 核函数调用核函数printf("kernel: called another with %d threads\n", tnum);
}__global__ void kernel2() {printf("Block (%d,%d,%d) of (%d,%d,%d), Thread (%d,%d,%d) of (%d,%d,%d)\n",blockIdx.x, blockIdx.y, blockIdx.z,gridDim.x, gridDim.y, gridDim.z,threadIdx.x, threadIdx.y, threadIdx.z,blockDim.x, blockDim.y, blockDim.z);
}// 默认host
constexpr const char* cuthead(const char* p) {return p + 1;
}// 内存访问
__global__ void kernel_memory(int* pret) {*pret = 42;
}// 数组并行处理
__global__ void kernel_array(int* arr, int n) {for (int i = blockDim.x * blockIdx.x + threadIdx.x;i < n; i += blockDim.x * gridDim.x) {arr[i] = i;}
}// 并行处理模板
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);}
}// 自定义加减乘除
__device__ __inline__ int float_atomic_add(float* dst, float src) {int old = __float_as_int(*dst), expect;do {expect = old;old = atomicCAS((int*)dst, expect,__float_as_int(__int_as_float(expect) + src));} while (expect != old);return old;
}// map
template <int blockSize, class T>
__global__ void parallel_sum_kernel(T* sum, T const* arr, int n) {__shared__ volatile int local_sum[blockSize];int j = threadIdx.x;int i = blockIdx.x;T temp_sum = 0;for (int t = i * blockSize + j; t < n; t += blockSize * gridDim.x) {temp_sum += arr[t];}local_sum[j] = temp_sum;__syncthreads();if constexpr (blockSize >= 1024) {if (j < 512)local_sum[j] += local_sum[j + 512];__syncthreads();}if constexpr (blockSize >= 512) {if (j < 256)local_sum[j] += local_sum[j + 256];__syncthreads();}if constexpr (blockSize >= 256) {if (j < 128)local_sum[j] += local_sum[j + 128];__syncthreads();}if constexpr (blockSize >= 128) {if (j < 64)local_sum[j] += local_sum[j + 64];__syncthreads();}if (j < 32) {if constexpr (blockSize >= 64)local_sum[j] += local_sum[j + 32];if constexpr (blockSize >= 32)local_sum[j] += local_sum[j + 16];if constexpr (blockSize >= 16)local_sum[j] += local_sum[j + 8];if constexpr (blockSize >= 8)local_sum[j] += local_sum[j + 4];if constexpr (blockSize >= 4)local_sum[j] += local_sum[j + 2];if (j == 0) {sum[i] = local_sum[0] + local_sum[1];}}
}// reduce
template <int reduceScale = 4096, int blockSize = 256, int cutoffSize = reduceScale * 2, class T>
int parallel_sum(T const* arr, int n) {if (n > cutoffSize) {std::vector<int, CudaAllocator<int>> sum(n / reduceScale);parallel_sum_kernel<blockSize> << <n / reduceScale, blockSize >> > (sum.data(), arr, n);return parallel_sum(sum.data(), n / reduceScale);}else {checkCudaErrors(cudaDeviceSynchronize());T final_sum = 0;for (int i = 0; i < n; i++) {final_sum += arr[i];}return final_sum;}
}// 共享内存
template <int blockSize, class T>
__global__ void parallel_transpose(T* out, T const* in, int nx, int ny) {int x = blockIdx.x * blockSize + threadIdx.x;int y = blockIdx.y * blockSize + threadIdx.y;if (x >= nx || y >= ny) return;__shared__ T tmp[blockSize * blockSize];int rx = blockIdx.y * blockSize + threadIdx.x;int ry = blockIdx.x * blockSize + threadIdx.y;tmp[threadIdx.y * blockSize + threadIdx.x] = in[ry * nx + rx];__syncthreads();out[y * nx + x] = tmp[threadIdx.x * blockSize + threadIdx.y];
}void testArray() {int n = 65535;int* arr;checkCudaErrors(cudaMallocManaged(&arr, n * sizeof(int)));int nthreads = 128;int nblocks = (n + nthreads + 1) / nthreads;kernel_array << <nblocks, nthreads >> > (arr, n);checkCudaErrors(cudaDeviceSynchronize());for (int i = 0; i < n; i++) {printf("arr[%d]: %d\n", i, arr[i]);}cudaFree(arr);
}void testVector() {int n = 65536;std::vector<int, CudaAllocator<int>> arr(n);parallel_for << <32, 128 >> > (n, [arr = arr.data()] __device__(int i) {arr[i] = i;});checkCudaErrors(cudaDeviceSynchronize());for (int i = 0; i < n; i++) {printf("arr[%d] = %d\n", i, arr[i]);}}void testMath() {int n = 1 << 25;std::vector<float, CudaAllocator<float>> gpu(n);std::vector<float> cpu(n);for (int i = 0; i < n; i++) {cpu[i] = sinf(i);}parallel_for << <n / 512, 128 >> > (n, [gpu = gpu.data()] __device__(int i) {gpu[i] = __sinf(i);});checkCudaErrors(cudaDeviceSynchronize());//for (int i = 0; i < n; i++) {//printf("diff %d = %f\n", i, gpu[i] - cpu[i]);//}
}void testTrust() {int n = 65536;thrust::universal_vector<float> x(n);   // CPU或GPUthrust::host_vector<float> x_host(n);   // CPUthrust::device_vector<float> x_dev = x_host;    // GPUthrust::for_each(x_dev.begin(), x_dev.end(), [] __device__(float& x) {x += 100.f;});
}void testAtomic() {// atomicAdd(dst, src) // atomicSub// atomicOr// atomicAnd// atomicXor// atomicMax// atomicMin// atomicExch   old = *dst; *dst = src// atomicCAS    automic::exchange
}void test_share() {int nx = 1 << 14, ny = 1 << 14;std::vector<int, CudaAllocator<int>> in(nx * ny);std::vector<int, CudaAllocator<int>> out(nx * ny);parallel_transpose<32> << <dim3(nx / 32, ny / 32, 1), dim3(32, 32, 1) >> >(out.data(), in.data(), nx, ny);
}// host 可以调用 global；global 可以调用 device；device 可以调用 device
// 实际上 GPU 的板块相当于 CPU 的线程，GPU 的线程相当于 CPU 的SIMD
// 数学函数sqrtf，rsqrtf，cbrtf，rcbrtf，powf，sinf，cosf，sinpif，cospif，sincosf，sincospif，logf，log2f，log10f，expf，exp2f，exp10f，tanf，atanf，asinf，acosf，fmodf，fabsf，fminf，fmaxf
// 低精度，高性能 __sinf,__expf、__logf、__cosf、__powf, __fdividef(x, y) 
// old = atomicAdd(dst, src) 相当于 old = *dst; *dst += src 
int main() {// 三重尖括号里的，第一个参数表示板块数量，第二个参数决定着启动 kernel 时所用 GPU 的线程数量kernel<<<1, 1>>>(); // CPU调用GPU函数异步执行kernel2<<<dim3(2, 1, 1), dim3(2, 2, 2)>>>();cudaDeviceSynchronize(); // CPU等待GPU完成say_hello_host();int* pret;checkCudaErrors(cudaMallocManaged(&pret, sizeof(int)));kernel_memory << <1, 1 >> > (pret);checkCudaErrors(cudaDeviceSynchronize());printf("result: %d\n", *pret);cudaFree(pret);testArray();float sin10 = sinf(10.f);/// map-reduceint n = 1 << 24;std::vector<int, CudaAllocator<int>> arr(n);int final_sum = parallel_sum(arr.data(), n);return 0;
}

参考

cuda-samples/Common/helper_string.h at 5f97d7d0dff880bc6567faa4c5e62e389a6d6999 · NVIDIA/cuda-samples · GitHub

 https://github.com/NVIDIA/cuda-samples/blob/5f97d7d0dff880bc6567faa4c5e62e389a6d6999/Common/helper_cuda.h#L31

https://www.nvidia.cn/docs/IO/51635/NVIDIA_CUDA_Programming_Guide_1.1_chs.pdf

https://developer.download.nvidia.cn/CUDA/training/StreamsAndConcurrencyWebinar.pdf

CUDA Pro Tip: Write Flexible Kernels with Grid-Stride Loops | NVIDIA Technical Blog

创作不易，小小的支持一下吧！

本文来自互联网用户投稿，该文观点仅代表作者本人，不代表本站立场。本站仅提供信息存储空间服务，不拥有所有权，不承担相关法律责任。如若转载，请注明出处：http://www.mzph.cn/pingmian/31752.shtml

如若内容造成侵权/违法违规/事实不符，请联系多彩编程网进行投诉反馈email:809451989@qq.com，一经查实，立即删除！

相关文章

域名反查IP多种方式

域名反查IP多种方式

域名反查IP是指通过已知的域名来查询其对应的IP地址。以下是几种常用的域名反查IP的方法&#xff1a; 使用ping命令&#xff1a; 在命令行中输入“ping 域名”&#xff0c;例如“ping www.example.com”。 系统会返回该域名对应的IP地址。 这种方法在Windows和Linux系统上都适…
阅读更多...
微服务中不同服务使用openfeign 相互调用

微服务中不同服务使用openfeign 相互调用

首先 我们上文 已经知道了 nacos 的注册服务&#xff0c;现在 我们 在不同服务中相互调用就可以使用openfeign 直接调用&#xff0c;而不是 再写冗余的调用代码啦 首先 我们的微服务组件如下 因为我这个微服务是我在 员工登录demo 中 拆出来的&#xff0c;在userlogin模块中…
阅读更多...
【计算机毕业设计】​206校园顺路代送微信小程序

【计算机毕业设计】​206校园顺路代送微信小程序

&#x1f64a;作者简介&#xff1a;拥有多年开发工作经验&#xff0c;分享技术代码帮助学生学习&#xff0c;独立完成自己的项目或者毕业设计。 代码可以私聊博主获取。&#x1f339;赠送计算机毕业设计600个选题excel文件&#xff0c;帮助大学选题。赠送开题报告模板&#xff…
阅读更多...
qmt量化交易策略小白学习笔记第43期【qmt编程之期货数据--如何获取历史主力合约--原生python】

qmt量化交易策略小白学习笔记第43期【qmt编程之期货数据--如何获取历史主力合约--原生python】

qmt编程之获取期货数据 qmt更加详细的教程方法&#xff0c;会持续慢慢梳理。 也可找寻博主的历史文章&#xff0c;搜索关键词查看解决方案 &#xff01; 感谢关注&#xff0c;咨询免费开通量化回测与获取实盘权限&#xff0c;欢迎和博主联系&#xff01; 获取历史主力合约 …
阅读更多...
Hi3861 OpenHarmony嵌入式应用入门--LiteOS Timer

Hi3861 OpenHarmony嵌入式应用入门--LiteOS Timer

iteOS Timer&#xff08;定时器&#xff09;是LiteOS操作系统中的一个重要组件&#xff0c;它提供了一种基于软件模拟的定时器功能&#xff0c;用于满足在硬件定时器数量不足时的定时需求。 软件定时器&#xff1a;基于系统Tick时钟中断&#xff0c;由软件来模拟的定时器。当经…
阅读更多...
【计算机网络体系结构】计算机网络体系结构实验-FTP实验

【计算机网络体系结构】计算机网络体系结构实验-FTP实验

1. 2. 3. wireshark 第一行&#xff1a;帧Frame 545&#xff1a;要发送的数据块&#xff0c;所抓帧的序号为545&#xff0c;捕获字节数等于传送字节数&#xff1a;451字节第二行&#xff1a;源Mac地址为a4:bb:6d:6e:28:9a&#xff1b;目标Mac地址为24:00:fa:e4:df:d8第三行&…
阅读更多...
mysql学习——多表查询

mysql学习——多表查询

多表查询 内连接外连接自连接自连接查询联合查询 子查询 学习黑马MySQL课程&#xff0c;记录笔记&#xff0c;用于复习。 添加外键 alter table emp add constraint fk_emp_dept_id foreign key (dept_id) references dept(id);多表查询 select * from emp , dept where emp…
阅读更多...
深度学习Day-21：ResNet与DenseNet结合

深度学习Day-21：ResNet与DenseNet结合

&#x1f368; 本文为&#xff1a;[&#x1f517;365天深度学习训练营] 中的学习记录博客 &#x1f356; 原作者&#xff1a;[K同学啊 | 接辅导、项目定制] 要求&#xff1a; 探索ResNet与DenseNet结合的可能性根据模型特性构建新的模型框架验证改进后模型的效果 一、 基础配…
阅读更多...
【linux】dup文件描述符复制函数和管道详解

【linux】dup文件描述符复制函数和管道详解

目录 一、文件描述符复制 1、dup函数&#xff08;复制文件描述符&#xff09; ​编辑 2、dup2函数&#xff08;复制文件描述符&#xff09; ​编辑 二、无名管道pipe 1、概述 2、无名管道的创建 3、无名管道读写的特点 4、无名管道ps -A | grep bash实现 三、有名管道FI…
阅读更多...
Java8使用Stream流实现List列表查询、统计、排序、分组、合并

Java8使用Stream流实现List列表查询、统计、排序、分组、合并

Java8使用Stream流实现List列表查询、统计、排序以及分组 目录 一、查询方法1.1 forEach1.2 filter(T -> boolean)1.3 filterAny() 和 filterFirst()1.4 map(T -> R) 和 flatMap(T -> Stream)1.5 distinct()1.6 limit(long n) 和 skip(long n) 二、判断方法2.1 anyMa…
阅读更多...
threejs-- add()和attach（）的区别(不受父对象影响)

threejs-- add()和attach（）的区别(不受父对象影响)

add和attach的区别 add()方法:attach()方法:总结区别: 在Three.js中&#xff0c;add()和attach()方法都涉及将一个物体&#xff08;object&#xff09;添加到另一个物体&#xff08;Object3D&#xff09;上&#xff0c;但它们有不同的作用和用法&#xff1a; add()方法: add(…
阅读更多...
容器之按钮盒构件演示

容器之按钮盒构件演示

代码; #include <gtk-2.0/gtk/gtk.h> #include <glib-2.0/glib.h> #include <gtk-2.0/gdk/gdkkeysyms.h> #include <stdio.h>int main(int argc, char *argv[]) {gtk_init(&argc, &argv);GtkWidget *window;window gtk_window_new(GTK_WINDO…
阅读更多...
Eclipse使用SpringXml to Java没有反应或者报错

Eclipse使用SpringXml to Java没有反应或者报错

Eclipse使用SpringXml to Java没有反应或者报错 定位错误方法&#xff1a; 通过Window -> Show View -> Error Log打开错误日志视图。 错误日志会记录Eclipse运行时发生的各种错误和警告&#xff0c;包括插件和工具的问题。 在错误日志中查找与你执行的Spring XML to J…
阅读更多...
Python学习系列之三目运算

Python学习系列之三目运算

Python学习系列之三目运算 前言C#的三目运算Python的三目运算总结 前言 在项目常有一些运算比较&#xff0c;之前使用的C#常用三目运算&#xff0c;减少使用switch或者if else来减少语句。 当C#转化为python时&#xff0c;三目运算使用不同了。 C#的三目运算 这里举个例子&am…
阅读更多...
xargs 传参

xargs 传参

xargs的默认命令是 echo&#xff0c;空格是默认定界符。这意味着通过管道传递给 xargs的输入将会包含换行和空白&#xff0c;不过通过 xargs 的处理&#xff0c;换行和空白将被空格取代。xargs是构建单行命令的重要组件之一。 xargs -n1 // 一次输出一个参数到一行&#xf…
阅读更多...
使用NestJS构建安全密码重置功能的完整指南：实现短信链接跳转验证功能

使用NestJS构建安全密码重置功能的完整指南：实现短信链接跳转验证功能

引言 实现忘记密码的短信链接验证功能&#xff0c;可以按照以下步骤进行&#xff1a; 用户请求重置密码&#xff1a;用户提供注册手机号码&#xff0c;系统生成一个唯一的重置令牌&#xff08;token&#xff09;&#xff0c;将令牌和用户信息存储在数据库中&#xff0c;并将包…
阅读更多...
Linux系统异常进程管理

Linux系统异常进程管理

Linux系统异常进程管理 1、异常关闭服务和进程 1&#xff09;【杀】进程 kill 进程【号】 ##温和、优雅 pkill 进程【名】 ##一下爆头 killall 进程【名】 ##优雅&#xff0c;可能需要多次反复 2&#xff09;杀不掉处理&#xff08;慎用&#xff09; 强制&#xff0c;一…
阅读更多...
Python学习笔记16：进阶篇(五)异常处理

Python学习笔记16：进阶篇(五)异常处理

异常 在编程中&#xff0c;异常是指程序运行过程中发生的意外事件&#xff0c;这些事件通常中断了正常的指令流程。它们可能是由于错误的输入数据、资源不足、非法操作或其他未预料到的情况引起的。Python中&#xff0c;当遇到这类情况时&#xff0c;会抛出一个异常对象&#…
阅读更多...
最详细的Selenium+Pytest自动化测试框架实战

最详细的Selenium+Pytest自动化测试框架实战

前言 selenium自动化 pytest测试框架 本章你需要 一定的python基础——至少明白类与对象&#xff0c;封装继承 一定的selenium基础——本篇不讲selenium&#xff0c; 测试框架简介 测试框架有什么优点呢&#xff1a; 代码复用率高&#xff0c;如果不使用框架的话&#xff…
阅读更多...
一个虚拟空间可以放多个不同类型的网站吗

一个虚拟空间可以放多个不同类型的网站吗

通常一些个人站长或者公司可能同时拥有几个网站&#xff0c;由于其他几个网站流量不高&#xff0c;而每个网站都租用一个虚拟主机空间的话&#xff0c;感觉有点浪费。大家可能会想虚拟主机能不能也像独立服务器那样放置多个网站呢&#xff1f;答案是肯定的&#xff0c;确定主机…
阅读更多...
最新文章

Copyright @ 2022~2023