并行计算01-基于cuda编程01

摘要：第一个cuda编程，主要是想看一下编程思维的与一般的编程的区别；实践一个小例子进行感性入门测试。先用起来再说，后面再去理解原理，学与做相结合的学习方法。

1. 安装cuda

这个PPT讲了cuda的安装
https://wenku.baidu.com/view/a2fab32f5fbfc77da269b1f8.html

2. 创建项目

3. Host与device编程的区别

C编译，在Host上运行：

#include <stdio.h>int main(){    // 这个完全在主机上运行    printf("hello world\n");    system("pause");    return 0;}

经典的hello word, 这个只在host上运行，并没有考虑主机之外的设备【定义CUP及系统的内存称为主机；GPU及其内存称为设备；在GPU运行的函数称为核函数（Kernel）】。
增一个核函数，在设备上运行：

#include "cuda_runtime.h"#include "device_launch_parameters.h"#include <stdio.h>/* 核函数:__global__修饰符告诉编译器，函数应该编译为设备运行的。这个函数将交给设备的编译器编译。*/__global__ void aKernel(){}// main函数仍然是主机编译int main(){    aKernel<<<1,1>>>();    printf("hello world\n");    system("pause");    return 0;}

4. 矢量分布式求和并行模型

这个模型来自JS 与 EK《GPU高性能编程 Cuda实践》

5. 相关demo代码

#include "cuda_runtime.h"#include "device_launch_parameters.h"#include <stdio.h>#include<iostream>const int N =10; // 核函数__global__ void aKernel(int a, int b, int *c){    // 设备上的指针与C指针的使用是一样的    *c = a + b;}// 在CPU进行计算，这是一个串行的计算void add( int *a ,int *b , int *c)  {      int tid = 0;  // 这是第0个cup,因此索引从0开始；    while(tid < N)           {          c[tid] = a[tid] + b[tid];          tid += 1;  // 由于只有一个CPU，所以每次增1              }  }  // 核加__global__ void kernelAdd( int *a ,int *b , int *c)  {      int tid = blockIdx.x;   // 计算这个索引处理的数据    if(tid < N)      {          c[tid] = a[tid]+b[tid];      }  }  int main(){       printf("############demo01 hello word###################");    int c;    int *dev_c;    // 分配内存，第一个参数是一个指针，指向用于保存新分配内存地址的变量，第二个参数是分配内存的大小    // 主机可以作用dev_c作为参数的引用，不用可读取与写数据，也不要对它进行free。因为这个是在设备上的。    cudaMalloc( (void**)&dev_c,sizeof(int));    aKernel<<<1,1>>>(2,7,dev_c);    // 把结果复制回主机：内存指针，设备指针，大小，内存复制方向；另外还有两个cudaMemcpyHostToDevice&&cudaMemcpyDeviceToDevice ; 如果两个是主内存就是C了。    cudaMemcpy(&c, dev_c, sizeof(int), cudaMemcpyDeviceToHost);    printf("2+7=%d\n",c);    cudaFree(dev_c);    /*       主机上的指针只能访问主机的上的内存，设备上的指针只能访问设备上的内存。两个不可以相互操作。主机可以调用cudaMemcpy来访问设备的内存。    */    printf("############demo02 device info###################\n");    cudaDeviceProp prop;    int count;    cudaGetDeviceCount(&count);    printf("count: %d\n",count);    for(int i = 0 ; i  < count; i++){        cudaGetDeviceProperties(&prop, i);        printf("devide No:%d\n",i);        printf("标识设备的ASCII字符串: %s\n",prop.name);        printf("设备上全局内存的总量，单位为字节: %u\n",prop.totalGlobalMem);        printf("在一个线程块（Block）中可使用的共享内存总量，单位为字节: %u\n",prop.sharedMemPerBlock);        printf("每个线程块中可用的32位寄存器数量: %d\n",prop.regsPerBlock);        printf("在一个线程束（warp）中包含的线程数量: %d\n",prop.warpSize);        printf("在内存复制中最大的修正量（Pitch），单位为字节: %u\n",prop.memPitch);        printf("在一个线程块中包含的最大线程数目: %d\n",prop.maxThreadsPerBlock);        printf("在多维线程块数组中，每一维包含的最大线程数量:[%d,%d,%d]\n",prop.maxThreadsDim[0],prop.maxThreadsDim[1],prop.maxThreadsDim[2]);        printf("在一个线程格（Grid）中，每一维可以包含的线程块的数量: [%d,%d,%d]\n",prop.maxGridSize[0],prop.maxGridSize[1],prop.maxGridSize[2]);        printf("常量内存的总量: %u\n",prop.totalConstMem);        printf("设备计算功能集的主版本号: %d\n",prop.major);        printf("设备计算功能集的次版本号: %d\n",prop.minor);        printf("clock Rate: %d\n",prop.clockRate);        printf("设备的纹理对齐要求: %d\n",prop.textureAlignment);        printf("一个布尔类型值，表示设备是否可以同时执行一个cudaMemory()调用和一个核函数调用: %d\n",prop.deviceOverlap);        printf("设备上多处理器的数量: %d\n",prop.multiProcessorCount);        printf("一个布尔值，表示该设备上执行的核函数是否存在运行时限制: %d\n",prop.kernelExecTimeoutEnabled);        printf("一个布尔值，表示设备是否是一个集成的GPU: %d\n",prop.integrated);        printf("一个布尔值，表示设备是否将主机内存映射到cuda设备地址空间: %d\n",prop.canMapHostMemory);        printf("表示设备的计算模式：默认，独占或禁止: %d\n",prop.computeMode);        printf("一维纹理的最大大小: %d\n",prop.maxTexture1D);        printf("二维纹理的最大维数: [%d,%d]\n",prop.maxTexture2D[0],prop.maxTexture2D[1]);        printf("三维纹理的最大维数: [%d,%d,%d]\n",prop.maxTexture3D[0],prop.maxTexture3D[1],prop.maxTexture3D[2]);        //printf("二维纹理数组的最大维数: [%d,%d,%d]\n",prop.maxTexture2DArray[0],prop.maxTexture2DArray[1],prop.maxTexture2DArray[2]);        printf("一个布尔值，表示设备是否支持在同一个上下文中同时执行多个核函数: %d\n",prop.concurrentKernels);    }    printf("############demo03 矢量求和【CPU】###################\n");    int a[N],b[N],csum[N];      //在CPU上对数组'a'和'b'赋值      for(int i=0;i<N;i++)         {          a[i] = -1;          b[i] = i * i;              }      add(a,b,csum);    //打印结果      for(int i=0;i<N;i++)       {          printf("%d+%d=%d\n",a[i],b[i],csum[i]);              }      printf("############demo03 矢量求和【GPU】###################\n");    int d_a[N],d_b[N],d_c[N];      int *dev_a0, *dev_b0, *dev_c0;      //allocate memory on GPU      cudaMalloc( (void**)&dev_a0, N*sizeof(int) ) ;      cudaMalloc( (void**)&dev_b0, N*sizeof(int) ) ;      cudaMalloc( (void**)&dev_c0, N*sizeof(int) ) ;      // 初始化数    for(int i=0;i<N;i++)      {          d_a[i] = -1;          d_b[i] = i * i ;      }      // 从内存复制到设备    cudaMemcpy( dev_a0 , d_a, N*sizeof(int), cudaMemcpyHostToDevice ) ;      cudaMemcpy( dev_b0 , d_b, N*sizeof(int), cudaMemcpyHostToDevice ) ;      // 调用核函数：N表示设备在执行函数时使用的并行线程块的数量。这里也就意味着会有N个副本，通过blockIdx.x可获取线程块来计算，blockIdx为内置变量。    // cuda支持二维的线程块数组，例如矩阵运算或图像处理。直接使用二维索引就可以了。。    // 尖括号的1表示，每个线程块中线程的个数为1    kernelAdd<<<N,1>>>(dev_a0,dev_b0,dev_c0);      //把处理的结果从设备中复制数据回内存    cudaMemcpy(d_c, dev_c0, N*sizeof(int), cudaMemcpyDeviceToHost) ;      for(int i=0;i<N;i++)      {          printf("%d+%d=%d\n",d_a[i],d_b[i],d_c[i]);       }      //release the memory on GPU      cudaFree(dev_a0);      cudaFree(dev_b0);      cudaFree(dev_c0);      system("pause");    return 0;}

运行结果

6. 遇到问题与解决

生成项目出错：

1>CUDACOMPILE : nvcc warning : The 'compute_20', 'sm_20', and 'sm_21' architectures are deprecated, and may be removed in a future release (Use -Wno-deprecated-gpu-targets to suppress warning).1>CUDACOMPILE : nvcc warning : nvcc support for Microsoft Visual Studio 2010 and earlier has been deprecated and is no longer being maintained1>  kernel.cu1>  support for Microsoft Visual Studio 2010 has been deprecated!1>ManifestResourceCompile:1>  所有输出均为最新。1>LINK : fatal error LNK1123: 转换到 COFF 期间失败: 文件无效或损坏1>1>生成失败。1>1>已用时间 00:00:01.57========== 生成: 成功 0 个，失败 1 个，最新 0 个，跳过 0 个 ==========

解决：
参考
http://yacare.iteye.com/blog/2010049
具体操作，在代码空白处点击，选择项目-属性xxxx-如下图选择，最后点应用，就OK了。

运行结果【加入system(“pause”);】让cmd窗口停下来。

【作者：happyprince, http://blog.csdn.net/ld326/article/details/78572168】