zoukankan      html  css  js  c++  java
  • cuda编程学习3——VectorSum

    这个程序是把两个向量相加

    add<<<N,1>>>(dev_a,dev_b,dev_c);//<N,1>,第一个参数N代表block的数量,第二个参数1代表每个block中thread的数量

    tid=blockIdx.x;//blockIdx是一个内置变量,blockIdx.x代表这是一个2维索引

    下面对这个程序做几个变化,并指出相应的程序应该改变的地方:

    1.若启动1个block,每个block中有N个线程。改变:

    add<<<1,N>>>(dev_a,dev_b,dev_c);

    tid=threadIdx.x

    2.此程序N的大小为10,N是向量的大小,若N更大呢?硬件对线程块的数量是有限制的,对每个线程块中线程的数量也是有限制的,可以通过运行deviceQuery进行查看,我是NVIDIA Geforce 750Ti,cuda7.5,线程块数量不能超过65535,每个线程块中线程的数量不能超过1024。若所求的向量中元素大于最多能启动的线程块数量和每个线程块中最大线程数量呢?改变:

    add<<<m,n>>>(dev_a,dev_b,dev_c)//m为启动的block数量,n为每个block中thread数量

    tid=threadIdx.x+blockIdx.x*blockDim.x;//blockDim是一个内置变量,保存的是线程块中每一维的线程的数量

    注意:m应该=N/n,但是若n不能整除N呢?答案:m=(N+n-1)/n,即为大于或等于N的n的最小倍数

    3.若向量大小N大于总的可启动线程数量(最大block数量×block中最大thread数量)呢?改变:

    add<<<128,1024>>>(dev_a,dev_b,dev_c);

    tid=threadIdx.x+blockIdx.x*blockDim.x;

    while(tid<N)

    {
    c[tid]=a[tid]+b[tid];

    tid+=blockDim.x*gridDim.x;
    }

    注意:此处启动了128个block,每个block有1024个thread,用户可以更改这个值,此处只是以<128,1024>为例。

    tid+=blockDim.x*gridDim.x;一次循环,执行的线程数量为blockDim.x*gridDim.x,即线程块大小×线程块数量。

    代码:

    /*
    ============================================================================
    Name : VectorSum-CUDA.cu
    Author : can
    Version :
    Copyright : Your copyright notice
    Description : CUDA compute reciprocals
    ============================================================================
    */
    #include<iostream>
    using namespace std;
    #define N 10
    __global__ void add(int *a,int *b,int *c);
    static void checkCudaErrorAux(const char *,unsigned, const char *,cudaError_t);
    #define CUDA_CHECK_RETURN(value) checkCudaErrorAux(__FILE__,__LINE__,#value,value)
    int main()
    {
    int a[N],b[N],c[N];
    int *dev_a,*dev_b,*dev_c;
    for(int i=0;i<N;i++)
    {
    a[i]=i;
    b[i]=i*i;
    }
    CUDA_CHECK_RETURN(cudaMalloc((void **)&dev_a,N*sizeof(int)));
    CUDA_CHECK_RETURN(cudaMalloc((void **)&dev_b,N*sizeof(int)));
    CUDA_CHECK_RETURN(cudaMalloc((void **)&dev_c,N*sizeof(int)));
    CUDA_CHECK_RETURN(cudaMemcpy(dev_a,a,N*sizeof(int),cudaMemcpyHostToDevice));
    CUDA_CHECK_RETURN(cudaMemcpy(dev_b,b,N*sizeof(int),cudaMemcpyHostToDevice));
    add<<<N,1>>>(dev_a,dev_b,dev_c);//<N,1>,第一个参数N代表block的数量,第二个参数1代表每个block中thread的数量
    CUDA_CHECK_RETURN(cudaMemcpy(c,dev_c,N*sizeof(int),cudaMemcpyDeviceToHost));
    for(int j=0;j<N;j++)
    {
    cout<<a[j]<<" + "<<b[j]<<" = "<<c[j]<<endl;
    }
    return 0;
    }

    __global__ void add(int* a,int* b,int* c)
    {
    int tid=blockIdx.x;//blockIdx是一个内置变量,blockIdx.x代表这是一个2维索引
    if(tid<N)//避免出错造成非法内存访问
    {
    c[tid]=a[tid]+b[tid];
    }
    }

    static void checkCudaErrorAux(const char *file,unsigned line, const char *statement,cudaError_t error)
    {
    if(error==cudaSuccess)
    {
    return;
    }
    cout<<statement<<"returned:"<<cudaGetErrorString(error)<<" at file:"<<file<<" line:"<<line<<endl;
    exit(1);
    }

  • 相关阅读:
    jMeter 里 CSV Data Set Config Sharing Mode 的含义详解
    如何使用 jMeter Parallel Controller
    使用 Chrome 开发者工具 coverage 功能分析 web 应用的渲染阻止资源的执行分布情况
    使用 Chrome 开发者工具的 lighthouse 功能分析 web 应用的性能问题
    关于 SAP 电商云首页加载时触发的 OCC API 请求
    SAP UI5 确保控件 id 全局唯一的实现方法
    SAP 电商云 Accelerator 和 Spartacus UI 的工作机制差异
    介绍一个好用的能让网页变成黑色背景的护眼 Chrome 扩展应用
    Chrome 开发者工具 performance 标签页的用法
    Client Side Cache 和 Server Side Cache 的区别
  • 原文地址:https://www.cnblogs.com/shrimp-can/p/5034284.html
Copyright © 2011-2022 走看看