CUDA教程,深入浅出谈cuda
库达加速的非常核心的一个应用是对矩阵运算进行加速，所以这篇博客，我会记录利用库达写矩阵乘法加速的过程。
矩阵乘法所谓矩阵乘法就是了AB=C，其中A为n*m的矩阵，B为m*q的矩阵，C为n*q的矩阵，矩阵乘法的代码如下：
const LL mod=1e 9 7；
结构矩阵{
LL a[2][2]；
void set1(){
memset(a，0，sizeof(a))；
}
void set2(){
set1()；
for(int I=0；I I)a[I][I]=1；
}
};
矩阵运算符*(常数矩阵一，常数矩阵b){
矩阵分辨率；
RES . set1()；
for(int I=0；i i ){
for(int j=0；j j ){
for(int k=0；k k ){
RES . a[I][j]=(RES . a[I][j]a . a[I][k]* b . a[k][j]% mod)% mod；
}
}
}
返回表示留数
}这里为了方便，接下来涉及的矩阵均为n*n的方形矩阵。
并行矩阵乘法核函数
//n * n矩阵乘A * B=C
_ _ global _ _ static void MatrixMulCUDA(const float * a，const float *b，float *c，int n)
{
const int tid=threadIdx.x//目前的线是第几个线
const int bid=blockIdx.x//目前的线是属于哪个街区
const int idx=bid * THREAD _ NUM tid//从出价和每日三次推出目前的线应该计算的A矩阵的行数和B矩阵列数
const int row=idx/n；
const int col=idx % n；
如果(第n行第n列){
float sum=0；
for(int I=0；i i ){
sum=a[row * n I]* b[I * n col]；
}
c[行*列]=总和；
}
}生成随机矩阵
//生成随机矩阵
void RandomMatrix(float *A，int n){
for(int I=0；i i ){
for(int j=0；j j ){
a[I * n j]=(float)RAND()/MY _ RAND _ MAX(float)RAND()/(MY _ RAND _ MAX * MY _ RAND _ MAX)；
}
}
}矩阵运算
void MatrixMul(){
//定义矩阵
float *a，*b，*c，* d；
int n=矩阵大小
a=(float *)malloc(sizeof(float)* n * n)；
b=(float *)malloc(sizeof(float)* n * n)；
c=(float *)malloc(sizeof(float)* n * n)；
d=(float *)malloc(sizeof(float)* n * n)；
srand(time(NULL))；
随机矩阵(a，n)；
随机矩阵(b，n)；
//把数据复制到显卡内存中
float *cuda_a，*cuda_b，* cuda _ c；
//cudaMalloc取得一块显存内存
cudaMalloc((void**) cuda_a，sizeof(float)* n * n)；
cudaMalloc((void**) cuda_b，sizeof(float)* n * n)；
cudaMalloc((void**) cuda_c，sizeof(float)* n * n)；
//cudamemacpy将产生的矩阵复制到显卡内存中
//cudamemacphyhosttodevice-从内存复制到显卡内存
//cudamemacpydevicetohost-从显卡内存复制到内存
cudamemacpy(cuda _ a，a，sizeof(float)*n*n，cudamemacphyhosttodevice)；
cudamemacpy(cuda _ b，b，sizeof(float)*n*n，cudamemacphyhosttodevice)；
float time _ elapsed=0；
cudaEvent_t启动，停止；
cudaEventCreate(开始);//创建事件
cudaEventCreate(停止);
cudaEventRecord( start，0)；//记录当前时间
MatrixMulCUDA blocks_num，THREAD_NUM，0 (cuda_a，cuda_b，cuda_c，n)；
cudaEventRecord( stop，0)；//记录当前时间
cudaEventSynchronize(启动);//等待事件完成。
cudaEventSynchronize(停止);//等待事件完成。记录之前的任务
//把结果从显示芯片复制回主内存
//cudamemacpy将结果从显存中复制回内存
cudamemacpy(c，cuda_c，sizeof(float) * n * n，cudamemacpydevicetohost)；
cudaEventElapsedTime(time _ elapsed，start，stop)；//计算时间差
cudaEventDestroy(开始);//销毁事件
cudaEventDestroy(停止);
//免费
cuda免费(cuda _ a)；
cuda免费(cuda _ b)；
cuda free(cuda _ c)；
printf(矩阵乘法国家政治保卫局。参见OGPU时间：%.10f\n ，time _ elapsed)；
时钟_t起始时间=clock()；
//CPU计算矩阵乘法
for(int I=0；i i ){
for(int j=0；j j ){
双温度=0;
for(int k=0；k k ){
temp=a[I * n k]* b[k * n j]；
}
d[I * n j]=temp；
}
}
时钟_t结束_时间=时钟();
//验证正确性和准确性
float max_error=0.0，average _ error=0；
for(int I=0；i i ){
for(int j=0；j j ){
if(d[i * n j]！=0){
float err=fabs((c[I * n j]-d[I * n j])/d[I * n j])；
if(max _ error err)max _ error=err；
平均值_误差=错误
}
}
}
double CPU _ time=(double)(end _ time-start _ time)/CLOCKS _ PER _ SEC * 1000.0；
printf(矩阵乘中央处理器时间：%.10f\n ，CPU _ time)；
printf(最大误差：%.10f平均误差：%.10f\n ，max_error，Average _ error/(n * n))；
printf(%.10f\n ，CPU _ time/time _ elapsed)；
}矩阵运算存在精度问题。因为我们在CPU上使用double(即64位浮点数)使计算过程递进，所以在GPU上只能使用float(32位浮点数)。在累加大量数字时，由于累加结果很快会变大，后面的数字很容易因位数过多而被截断。例如，下面的代码：
浮点数a=100998
浮动b=2.338
a=a b；
printf(总和为%.10f\n ，a)；输出为：101000.3359375000。很容易看到一个0.008被截断，1000次运算后，一个整数8就没了，这就造成了很大的精度问题。
卡汉求和公式为了在一定程度上解决这个精度问题，我们需要记住这个损失的误差值。假设误差为temp=(a b)-a-b，在上述问题中，temp=-0.008，在下一次计算中加入下一个加数，可以在一定程度上减小误差。
有一篇关于这个算法的好文章。该算法的伪代码描述如下：
def KahanSum(输入):
var sum=0.0
var c=0.0
对于i=1的输入.长度do
var y=input[i] - c //最初，c为零；那么它会补偿之前的精度。
var t=sum y//y的低位数字丢失
c=(t - sum) - y //恢复y的低位数字，带负号
总和=t
接下来我
Return sum在上面的伪代码中，变量C代表小数补偿的补码部分，更严格的说应该是负补码部分。随着这个完成部分的不断积累，当这些截断误差积累到一定量级时，在求和时就不会被截断，这样整个求和过程的精度就可以控制得比较好。然后将该算法应用于我们的矩阵乘法核函数。
//n * n矩阵乘A * B=C
_ _ global _ _ static void MatrixMulCUDA(const float * a，const float *b，float *c，int n)
{
const int tid=threadIdx.x//当前线程是什么线程？
const int bid=blockIdx.x//当前线程属于哪个块？
const int idx=bid * THREAD _ NUM tid//从bid和tid中推导出当前线程应该计算的矩阵A和矩阵B的行数和列数
const int row=idx/n；
const int col=idx % n；
if(第n行第n列){
//float sum=0；
//for(int I=0；i i ){
//sum=a[row * n I]* b[I * n col]；
//}
浮点t=0；
浮点y=0；
for(int I=0；i i ){
浮动r；
y-=a[row * n I]* b[I * n col]；
r=t-y；
y=(r-t)y；
t=r；
}
c[row * n col]=t；
}
}然后不知道为什么，我的GPU上精度没有提高。
转载请联系作者授权，否则将追究法律责任。
郑重声明：本文由网友发布，不代表盛行IT的观点，版权归原作者所有，仅为传播更多信息之目的，如有侵权请联系，我们将第一时间修改或删除，多谢。