cuda初学(1):稀疏矩阵向量乘法(单精度)_xlsp的专栏-程序员宅基地_稀疏矩阵向量乘法

技术标签: cuda  character  CUDA学习  C/C++学习  fortran  thread  float  

初步学习CUDA编程,实现简单稀疏矩阵向量乘法运算,由于硬件限制,目前只测试了单精度程序

GPU计算子程序gpu_fmmv.cu:

#include <stdio.h>
#include <stdlib.h>
// CUDA-C includes
#include <cuda_runtime.h>

#ifdef __cplusplus
extern "C" {
#endif
//   For Fortran interface  //
#define GPU_fmmv gpu_fmmv_
extern void GPU_fmmv(int *, int *, int *, float *, float *, float *);

#ifdef __cplusplus
  }
#endif

#define THREAD_NUM 512

__global__ static void fmmv(int *neq, int *numcol, int *ia, float *a, float *v, float *w)
{
   const int tId = threadIdx.x;
   int row, col;

   for(row = tId; row < *neq; row += THREAD_NUM){
      w[row] = 0.0;
      for(int num = numcol[row]; num < numcol[row+1]; num ++){
         col = ia[num]-1;
         w[row] += a[num]*v[col];

      }
   }

}

void GPU_fmmv(int *neqi, int *numcol, int *ia, float *a, float *v, float *w)
{
   int *gpu_neq;
   int *gpu_ia, *gpu_numcol;
   float *gpu_a, *gpu_v, *gpu_w;

   int neq = *neqi;
   int nnz = numcol[neq];

/*   Malloc space on GPU device   */
   cudaMalloc((void **) &gpu_neq, sizeof(int));
   cudaMalloc((void **) &gpu_numcol, sizeof(float)*(neq+1));
   cudaMalloc((void **) &gpu_ia, sizeof(int)*nnz);
   cudaMalloc((void **) &gpu_a, sizeof(float)*nnz);
   cudaMalloc((void **) &gpu_v, sizeof(float)*neq);
   cudaMalloc((void **) &gpu_w, sizeof(float)*neq);

/*  Copy data to GPU device  */
   cudaMemcpy(gpu_neq, &neq, sizeof(int), cudaMemcpyHostToDevice);
   cudaMemcpy(gpu_numcol, numcol, sizeof(int) * (neq+1), cudaMemcpyHostToDevice);
   cudaMemcpy(gpu_ia, ia, sizeof(int) * nnz, cudaMemcpyHostToDevice);
   cudaMemcpy(gpu_a, a, sizeof(float) * nnz, cudaMemcpyHostToDevice);
   cudaMemcpy(gpu_v, v, sizeof(float) * neq, cudaMemcpyHostToDevice);

   fmmv<<<1,THREAD_NUM,0>>>(gpu_neq, gpu_numcol, gpu_ia, gpu_a, gpu_v, gpu_w);

   cudaMemcpy(w, gpu_w, sizeof(float) * neq, cudaMemcpyDeviceToHost);

   cudaFree(gpu_neq);
   cudaFree(gpu_numcol);
   cudaFree(gpu_ia);
   cudaFree(gpu_a);
   cudaFree(gpu_v);
   cudaFree(gpu_w);

}

主调用程序testgpu.f:

      implicit real*4(a-h, o-z)

      dimension numcol(1000), ia(10000)
      dimension a(10000), v(1000), w(1000)
      character*12 fname

      numarg = 0
      numarg = numarg + 1
      call getarg(numarg, fname)
      print *,'fname : ', fname
      open(21,file=fname,form='formatted',status='old')
      read(21,*) neq
      read(21,*) (numcol(i),i=1,neq+1)
      do i=1, neq
         read(21,*) (ia(j),j=numcol(i)+1, numcol(i+1))
      enddo
      do i=1,neq
         read(21,*) (a(j),j=numcol(i)+1, numcol(i+1))
      enddo
      read(21,*) (v(i),i=1,neq)
      close(21)

      call gpu_fmmv(neq, numcol, ia, a, v, w)

      print *,'w ======'
      print *,(w(i),i=1,neq)

      stop 0000
      end

编译命令:

nvcc -c gpu_fmmv.cu

ifort -o fmvv -L /usr/local/cuda/lib64/ -lcudart gpu_fmmv.o testgpu.f

版权声明:本文为博主原创文章,遵循 CC 4.0 BY-SA 版权协议,转载请附上原文出处链接和本声明。
本文链接:https://blog.csdn.net/xlsp/article/details/7393671

智能推荐

[原]一种使用android设备为ipad提供外置无线扩展存储的方法_iteye_14668的博客-程序员宅基地

由于带4G模块的ipad要贵1k 大洋, 64G的版本要比16G版本又要贵1k 大洋, 手头紧的A君想了一个办法来解决这个问题, 使乞丐版的ipad可以具有高大上皇帝版ipad类似的功能.  首先你要有一款安卓设备来做无线AP和移动无线外置扩展存储. 最好是可以插sd卡的 最大支持128G那种. 如果没有就没办法了. 需要root 其次, 需要一款ipad上面的文件浏览软件...

php json修改,PHP读写json文件及修改json的方法-php文件_科技体验者的博客-程序员宅基地

今天扣丁学堂小编主要是给大家分享一下PHP读写json文件及修改json的方法,喜欢PHP开发的小伙伴们可以随着小编一起来了解一下。PHP在线教程实例如下所示:// 追加写入用户名下文件$code="001";//动态数据$json_string = file_get_contents("text.json");// 从文件中读取数据到PHP变量$data = json_decode($json_...

能控能观判断_NIckWJJ的博客-程序员宅基地_能控能观怎么判断

能控能观判断1.能控判断1).约旦标准型下判别①T^{-1}AT为对角线,且特征值互异,则T^{-1}B不存在全零行时系统能控;若有相同特征值,则在满足上述要求的前提下,各特征值对应的T^{-1}B的行元素线性无关时系统能控。②T^{-1}AT中有约旦块,且每个约旦块特征值互异,则T^{-1}B中对应各约旦块的最后一行为非全零行时系统能控;若存在特征值相同的约旦块,则在满足上述要求的前提下,特征值相同的各约旦块对应的T^{-1}B的最后一行线性无关时系统能控。2).能控阵判别

C++ STL 常用(持续更新)__Epsilon_的博客-程序员宅基地

C++ STL 常用(持续更新)vectormap写LeetCode总是忘记一些常用的模板类用法,在此总结,望牢记。vector头文件#include &lt;vector&gt;using namespace std;初始化vector&lt;int&gt; nums = { 1,2,2,3,1};排序// 注意:1. sort用法要包含&lt;algorithm&gt;头文件。// 2. 如果不想改变原数组,可以提前copy一个#include &lt;a

用户数突破5亿,钉钉首次阐述“两个数字化”战略_趣味科技v的博客-程序员宅基地

10月13日,钉钉召开2021未来组织大会。会上,钉钉发布了未来组织年度趋势关键词“数字生产力”,首次对外公布并阐述“两个数字化”战略,正式发布钉钉6.3版本,推出全新协同产品“钉闪会”。...

深入理解ReentrantLock的实现原理_weixin_33836874的博客-程序员宅基地

ReentrantLock简介ReentrantLock是Java在JDK1.5引入的显式锁,在实现原理和功能上都和内置锁(synchronized)上都有区别,在文章最后我们再比较这两个锁。首先我们要知道ReentrantLock是基于AQS实现的,所以我们得对AQS有所了解才能更好的去学习掌握ReentrantLock,关于AQS的介绍可以参考我之前写的一篇文章《一文带你快速掌握AQS》...

随便推点

python3实现接口_python3.6+requests实现接口自动化1_weixin_39608394的博客-程序员宅基地

逐步完善中……以一个登录接口为例,展示一下目录结构目录1、项目目录2、登录接口和登录用例3、配置文件4、run_all.py1、项目目录以之前搭建的aiopms为平台写接口自动化,其中case中放模块和用例,common中放数据库连接信息等,config放邮箱登录信息,logs存放日志文件,report放报告文件,run_all.py执行脚本2、登录接口和登录用例login_api.py# cod...

ubuntu 18.04 安装 cuda_mixboot的博客-程序员宅基地

ubuntu 安装1, cuda安装1, cuda安装# nvidia-smiSat Jan 11 16:00:28 2020+-----------------------------------------------------------------------------+| NVIDIA-SMI 435.21 Driver Version: 435.21 ...

无法连接 MKS:套接字连接尝试次数太多正在放弃_斗转星移3的博客-程序员宅基地

VMware Workstation Pro打开虚拟机是提示:无法连接 MKS:套接字连接尝试次数太多正在放弃。这个是由于VMware Workstation部分服务未打开造成的,按window键+r,输入services.msc,回车,弹出服务设置对话框,把VMware 相关的右键属性都设置为手动,然后点击右上角点启动。就可以打开虚拟机了。...

Instant Run is disabled for non-debug variants_猎羽的博客-程序员宅基地

报错: Instant Run is disabled for non-debug variants解决办法:我是在用JRebel for Android插件时出现了上面的提示,原因是编译版本是release版本,因此不能使用。需要在Build Variants中换成debug版本。网上提供的解决办法如下(添加内容到build.gradle文件中):buildType

java 公钥 验签 xml_c# 对XML进行数字签名并且让java验签成功_美自的博客-程序员宅基地

实现:1.c#将xml报文做数字签名发送到java服务,java服务成功验签。2.c#服务对收到java服务推送的xml报文成功验签。前提:1.java服务要求遇到问题:1.Java和.net 之前 RSA 私钥秘钥格式问题。2.framework4.0 SignatureDescription could not be created for the signature algorithm su...

推荐文章

热门文章

相关标签