CUDA Reduction
cuda Reduction
有一段时间没更了,之前一直在想记录博客的意义何在,也没想出个所以然,但是不能三分钟热度,所以还是把之前想发的存货找一找
这段时间在准备工作的事情,复习了下cuda,cuda的reduction是非常常见的,而众多的方法中尤以warp shfl的方法最为简便和高效,这里整理了下代码,简述思路,供自己以后查阅,也欢迎大家交流
顺便一提,其实最近已经有了采用Tensor Core来进行reduction和scan操作的工作出现了,可以参考论文《Accelerating Reduction and Scan Using Tensor Core Units》,取得了比cub库更好的效果。期待随着WMMA API的进化,能让Tensor Core为除了MatMul以外的算法提供更好的计算效果。 *
言归正传,基于warp shuffle 的reduction分为三部分:warp粒度的reduction,block粒度的reduction和整个device的reduction
warp Reduction
其中warp粒度的reduction比较简洁,但也是整个re ...
cuda异步内存复制
cuda异步内存复制
在cuda handbook中发现了一个利用cuda event来进行的异步的内存的复制,觉得比较巧妙,但是美中不足的是书中代码里存在一个bug,不过思想还是很不错的
利用两个buffer来充当异步复制的缓冲区
1void *g_hostBuffers[2];
定义需要的两个event,这里定义了之后就立马record,以便将来函数内第一次Sync时不会出错。
12345cudaEventCreate( &g_events[0] );cudaEventCreate( &g_events[1] );// record events so they are signaled on first synchronizecudaEventRecord( g_events[0], 0 );cudaEventRecord( g_events[1], 0 );
之后的内存分配和数据初始化部分我们就省略掉。接下来是最重要的函数部分
1234567891011121314151617181920voidchMemcpyHtoD( void *device, const ...
一些项目和大作业汇总(持续更新)
一些项目和大作业汇总
最近整理电脑,发现了以前写的一些小项目啊,大作业什么的,虽然大多数随着换电脑丢掉了,但是也保存下来了一部分,最近整理了一下传到了github上,这里也做一个汇总,如果大家有需要可以作参考。
都是很久以前的代码了,写的比较烂,希望对大家能有一丁点的帮助
这些代码的编写也或多或少的参考过网络的博客以及师兄师姐同学们的代码,感谢大家的帮助
计算机网络作业
研究生课程的计算机网络的大作业
https://github.com/WhatGhost/nku_network_hw
其中每个里面都有当时提交作业时候的文档
HW3是一个FTP的客户端,通过发送请求判断返回码来进行一些操作。GUI用java写的。
HW4是一个UDP聊天室,有服务器端和客户端两部分。
简单C编译器
本科时候编译原理课程的大作业,简单的只支持浮点和整数的简单赋值和计算,以及if,while,for等简单分支语句的编译器。
https://github.com/WhatGhost/simple-c-compiler
二次剩余判定与求解
二次剩余
最近在做一些关于椭圆曲线的优化,其中message到点的映射需要用到二次剩余相关的知识,也是第一次接触相关算法,查了一些资料之后自己在GPU端实现了一下,记录一下相关的心得。
一个整数X的平方 X2X^{2}X2 除以另一个整数p的余数,就叫做X对于整数p的二次剩余。也就是当存在某个X使得
X2≡d(modp)X^{2} \equiv d (mod p)
X2≡d(modp)
则d就是模p的二次剩余。一般来说,p都是大素数。
当需要判断一个x是否在椭圆曲线上时,需要计算 y2=x3+ax+by^{2} = x^{3} +ax +by2=x3+ax+b
是否成立,也就是算x3+ax+bx^{3} +ax +bx3+ax+b是不是模p上的二次剩余
判断二次剩余
判断二次剩余可以用欧拉准则,但是欧拉准则也只能判断是不是,但是不能求解出具体的结果,用来验证足够了。
欧拉准则提出,d是模p的二次剩余的充要条件就是
d(p−1)/2≡1(modp)d^{(p-1)/2} \equiv 1 (modp)
d(p−1)/2≡1(modp)
d是模p的非二次剩余的充要条件就是
d(p−1) ...
leetcode-97
Leetcode 97题 交错字符串
给定三个字符串 s1, s2, s3, 验证 s3 是否是由 s1 和 s2 交错组成的。
1234示例 1:输入: s1 = "aabcc", s2 = "dbbca", s3 = "aadbbcbcac"输出: true
1234示例 2:输入: s1 = "aabcc", s2 = "dbbca", s3 = "aadbbbaccc"输出: false
题解
这个问题我开始想用双指针,但是出来的报错,后来发现有些例子的结果是错误的,因为双指针说白了就是个贪心的策略。
后来参考了leetcode的题解,发现用动态规划能获得正确的解,不会陷入贪心的错误解中。
我们设f(i,j)是第一个字符串的前i位和第二个字符串的前j位能不能交错组合成第三个字符串的前i+j位。
若s3[i+j]==s1[i],则f(i,j)=f(i-1,j)。
若s3[i+j]==s2[j],则f(i,j)=f(i,j-1)。
初始状态f(0,0)=true; ...
cuda(Unresolved extern function错误) 编译调用其他文件中的device函数
Unresolved extern function错误
编译过程中遇到Unresolved extern function xxxxxx 错误,一般导致这个错误的原因是在__global__和一些__device__函数放在了分开编译的不同文件中,并且在__global__中调用了这些__device__函数。
解决的方法也很简单,有两种解决方法
将二者放到同一文件中进行编译
通过添加-dc标志来显示启用relocatable device code模式,-dc选项等价于 -rdc=true -c
补充问题
在实际操作中按照文档中如下的的加上-dc进行编译
12nvcc -arch=sm_35 -dc xx1.cu -o xx1.onvcc -arch=sm_35 main.cu xx1.o -o main
还是会报错Unresolved extern function,没能解决问题
后来在编译main.cu时添加 -rdc=true,不再报错,问题解决
但是文档中都只是说添加-dc就可以,为什么还需要再添加-rdc暂时还没搞明白。
12nvcc -arch=sm_35 -dc ...
leetcode 365 水壶问题
leetcode 365题 水壶问题
题目
有两个容量分别为 x升 和 y升 的水壶以及无限多的水。请判断能否通过使用这两个水壶,从而可以得到恰好 z升 的水?
如果可以,最后请用以上水壶中的一或两个来盛放取得的 z升 水。
你允许:
· 装满任意一个水壶
· 清空任意一个水壶
· 从一个水壶向另外一个水壶倒水,直到装满或者倒空
1234示例 1:输入: x = 3, y = 5, z = 4输出: True
1234示例 2:输入: x = 2, y = 6, z = 5输出: False
解题思路
这个问题最初始的解法就是深度优先搜索,这里不再详细介绍,网上有很多代码和方法
在查看题解的时候发现有一个很巧妙的方法,在这里分享
就是用数学方法来做,因为题目要求,不存在两个桶都不满的情况,所以我们每次的操作都只会对整体水量添加或减少x或者是y的水量
把不满的桶倒掉,如果另一桶是满的,相当于从空直接向另一桶灌满的操作,若果另一桶是空的,则相当于回归到最初的空状态
把不满的桶灌满,如果另一桶是满的,相当于从空直接把两个桶灌满的操作,若果另一桶是空的,则相当于把这个桶灌满的操作
所以我们可 ...
cuda 函数前缀 __host__ __device__ __global__ ____noinline__ 和 __forceinline__ 简介
cuda 函数前缀 host device global __noinline 和 forceinline 简介
这些函数前缀在官方的文档里被称为函数执行环境标识符Function execution space specifiers,也就是他指明了这段函数是在哪里被调用的。
global
***在GPU上执行*** ,但是需要 ***在CPU端调用***123456789101112131415161718192021PS.计算能力3.5以上的引入了 CUDA Dynamic Parallelism机制可以才devices端调用核函数```__global__```修饰的函数必须采用```void```返回值,并且需要在调用时制定 ***运行的参数*** (也就是<<<>>>里的block数和线程数)同时```__global___```函数是异步的,这也代表着函数没被执行完就返回了控制权,所以测量核函数的时间需要同步操作才能获得准确的结果。## __device__```__device__```函数是在GPU端调用且在GPU端执行的的函数``` ...
leetcode 322 零钱兑换问题
Leetcode 322题 零钱兑换问题
题目及实例
给定不同面额的硬币 coins 和一个总金额 amount。编写一个函数来计算可以凑成总金额所需的最少的硬币个数。如果没有任何一种硬币组合能组成总金额,返回 -1。
12345示例 1:输入: coins = [1, 2, 5], amount = 11输出: 3 解释: 11 = 5 + 5 + 1
1234示例 2:输入: coins = [2], amount = 3输出: -1
这个题目乍一看像是贪心算法,但是贪心算法容易陷入局部最优,如coins=[1,2,7,10],全局最优解应为[7,7],但是谈心会陷入局部最优解[2,2,10];
零钱问题曾经是我大二算法课考试的题目,一般会采用动态规划的解题思路,但是在leetcode的题解中看到了曾经我觉得性能可能会很差的回溯(+剪枝)的方法竟有很不错的性能,给了我很大的启发,所以下面我回把两个方法都介绍一边
动态规划
解析
动态规划其实是一个填表的过程,数组dp是一个长度为amount+1的数组(为什么要加1,因为钱数从0到amount有amount+1个数),其中第i个元素代 ...
leetcode 1071 字符串最大公因子问题
leetcode 1071题 字符串最大公因子问题
题目
对于字符串 S 和 T,只有在 S = T + … + T(T 与自身连接 1 次或多次)时,我们才认定 “T 能除尽 S”。
返回最长字符串 X,要求满足 X 能除尽 str1 且 X 能除尽 str2。
1234示例 1:输入:str1 = "ABCABC", str2 = "ABC"输出:"ABC"
1234示例 2:输入:str1 = "ABABAB", str2 = "ABAB"输出:"AB"
1234示例 3:输入:str1 = "LEET", str2 = "CODE"输出:""
解题思路
这个问题由于找的是最大的可以整除的字符串,其实就是求最大公因子的问题
首先由于这个串必须能除尽str1和str2,我们假设str1=n*X ,str2=m*X,所以str1+str2==str2+str1=(n*m)*X
所以我们只需要求str1.s ...