![](https://img-blog.csdnimg.cn/20201014180756757.png?x-oss-process=image/resize,m_fixed,h_64,w_64)
CUDA
文章平均质量分 76
USST_Coder
这个作者很懒,什么都没留下…
展开
-
CUDA性能优化系列——Kmeans算法调优(三)
本篇对调度方式进行优化,实现内存拷贝和计算overlap。单流同步调用/*单流同步*/void CallKmeansSync(){ //TODO:init host memory float* h_Src/*[Coords][SrcCount]*/, * h_Clusters/*[ClusterCount][Coords]*/; int* h_MemberShip/*[kSrcCount]*/, * h_MemberCount/*[kClusterCount]*/; h_Src原创 2021-04-07 15:47:42 · 1827 阅读 · 0 评论 -
CUDA性能优化系列——Kmeans算法调优(二)
本篇介绍Kmeans算法中计算新的聚类中心部分。这部分主要逻辑:根据计算出的新的分类信息,对全部数据点依次对每个类别求出所属当前类别的数据点个数与坐标和。本质上就是进行规约运算。V1 Atomic实现全局规约由于最终生成16个聚类中心,因此这里的规约操作需要针对算法进行一定的调整。V1在实现逻辑为:先在共享内存上分别通过原子操作,对16个类进行规约,再通过原子操作进行设备内存的全局规约操作。/*V1 atomic规约*/__global__ void kKmeansSumAtomic(int Src原创 2021-03-24 11:53:37 · 1700 阅读 · 0 评论 -
CUDA性能优化系列——Kmeans算法调优(一)
__global__ void k_kmeans(float* d_src, int srcsize, int dim, float* d_cluster, int clustersize, float* d_dst){ extern __shared__ float sm_cluster[]; float regData[4] = { 0.0f,0.0f, 0.0f, 0.0f };//维度不确定的情况下how to use registers???????? int tid = threadId原创 2021-03-09 00:26:21 · 3340 阅读 · 2 评论 -
三种基于CUDA的归约计算
归约在并行计算中很常见,并且在实现上具有一定的套路。本文分别基于三种机制(Intrinsic,共享内存,atomic),实现三个版本的归约操作,完成一个warp大小的数据的归约求和计算。Intrinsic版本基于Intrinsic函数 __shfl_down_sync 实现,使一个warp内的线程通过读取相邻线程寄存器中数据,完成归约操作。实现如下:__global__ void kIntrinsicWarpReduce(int* d_src, int* d_dst){ int val = d_原创 2021-03-02 14:46:20 · 1855 阅读 · 0 评论 -
CUDA 统计数组直方图
#include "cuda_runtime.h"#include "device_launch_parameters.h"#include <stdio.h>#include <stdlib.h>#include <string>#include <cooperative_groups.h>#include <iostream>#define HISTOGRAMGRID 36#define HISTOGRAMBLOCK 1.原创 2020-12-25 14:47:41 · 872 阅读 · 0 评论 -
Pinned Memory 多设备异步拷贝
void TestHostPinnedMem(){ //compare pinned and pagable memcpy const int size = 1024 * 1024 * 100; int* h1 = (int*)malloc(size * sizeof(int)); int* d1; cudaMalloc(&d1, size * sizeof(int)); cudaMemcpy(d1, h1, size * sizeof(int), cudaMemcpyHostToDe原创 2020-11-23 15:59:32 · 464 阅读 · 0 评论 -
CUDA处理归约问题
归约问题由于计算操作的对称性,非常适合并行处理。本文以数组求和为例,通过CUDA先实基础版本,并基于基础版本尝试通过不同的优化手段实现几个方案,最后将所有优化手段集成到最终的实现。其中核函数执行时间是通过Nsight System工具抓取,命令参数如下nsys profile --stats true a.out 5基础版本实现核函数实现__global__ void kReduceWholeWarpSafe(int64_t* d_src, int64_t* d_sum){ int64_t*原创 2020-11-20 14:24:51 · 470 阅读 · 0 评论 -
cuda 配置共享内存Bank模式API
记录GPU共享内存4Byte or 8Byte模式的cuda API。Access Mode Configuration对Kepler来说,默认情况是4-byte模式,可以用下面的API来查看:cudaError_t cudaDeviceGetSharedMemConfig(cudaSharedMemConfig *pConfig);返回结果放在pConfig中,其结果可以是下面两种:cudaSharedMemBankSizeFourBytecudaSharedMemBankSizeEigh原创 2020-11-02 11:15:06 · 861 阅读 · 0 评论 -
Nsight Visual Studio 调试
官方手册:https://docs.nvidia.com/nsight-visual-studio-edition/5.4/Nsight_Visual_Stud...转载 2020-08-27 22:58:23 · 5190 阅读 · 1 评论 -
CUDA C编程权威指南笔记 第二章 编程模型
原创 2020-08-08 00:13:29 · 268 阅读 · 0 评论 -
CUDA C编程权威指南笔记 第三章 执行模型
原创 2020-08-10 00:06:48 · 177 阅读 · 0 评论 -
CUDA C编程权威指南笔记 第四章 全局内存
原创 2020-08-20 23:53:31 · 194 阅读 · 0 评论 -
CUDA从入门到精通:性能剖析和Visual Profiler
入门后的进一步学习的内容,就是如何优化自己的代码。我们前面的例子没有考虑任何性能方面优化,是为了更好地学习基本知识点,而不是其他细节问题。从本节开始,我们要从性能出发考虑问题,不断优化代码,使执行速度提高是并行处理的唯一目的。测试代码运行速度有很多方法,C语言里提供了类似于SystemTime()这样的API获得系统时间,然后计算两个事件之间的时长从而完成计时功能。在CUDA中,我们有专门测量...转载 2020-08-12 19:10:52 · 435 阅读 · 0 评论 -
CUDA C 编程指南
《CUDA C Programming Guide》(《CUDA C 编程指南》)导读田子宸浙大水硕在读184 人赞同了该文章说明转自知乎《CUDA C Programming Guide》(《CUDA C 编程指南》)导读 - 田子宸的文章 - 知乎https://zhuanlan.zhihu.com/p/53773183最近在学习CUDA,感觉看完就忘,于是这里写一个导读,整理一下重点主要内...转载 2020-02-22 19:43:12 · 8062 阅读 · 1 评论