并行计算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 maintained
1>  kernel.cu
1>  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】