cuda基础

最新推荐文章于 2024-07-05 20:32:07 发布

weixin_30795127

最新推荐文章于 2024-07-05 20:32:07 发布

阅读量339

点赞数

文章标签： python c/c++

原文链接：http://www.cnblogs.com/larry-xia/p/11574203.html

版权

CUDA项目配置

（1）打开vs，创建一个空win32程序,即cuda_test项目。

（2）选择cuda_test，点击右键–>生成依赖项–>生成自定义，选择CUDA10.0。

（3）右键源文件文件夹->添加->新建项->选择CUDA C/C++File，取名cuda_main。

（4）点击cuda_main.cu的属性,在配置属性–>常规–>项类型–>选择“CUDA C/C++”。

（5）包含目录配置：

右键点击项目属性–>属性–>配置属性–>VC++目录–>包含目录

添加包含目录：$(CUDA_PATH)\include

（6）库目录配置

添加库目录：$(CUDA_PATH)\lib\x64

（7）依赖项

配置属性–>链接器–>输入–>附加依赖项

添加库文件：cublas.lib;cuda.lib;cudadevrt.lib;cudart.lib;cudart_static.lib;OpenCL.lib

将CPU及其系统的内存称为主机host,将GPU及其内存称为设备device.

线程块Block由多个线程组成(可以组织为一维、二维和三维），各block是并行执行的，block间无法通信，也没有执行顺序。

线程格Grid由多个线程块组成

线程束Warp：指一个包含32个线程的集合，被“编织在一起”并且“步调一致”的形式执行。在程序中的每一行，线程束中的每个线程都将在不同数据上执行相同的命令。

核函数Kernel:在GPU上执行的函数通常称为核函数，一般通过标识符__global__修饰，调用通过<<<参数1,参数2>>>，用于说明内核函数中的线程数量，以及线程是如何组织的。

CUDA C需要使用某种语法将一个函数标记为“设备代码”，CADA C提供了与C在语言级别上的集成，使得设备调用看起来非常像主机函数调用。尖括号表示要将一些参数传递给运行时系统，告诉运行时如何启动设备代码。

CUDA编程模式

　　1. 定义需要在 device 端执行的核函数。( 函数声明前加 _golbal_ 关键字 )

2. 在显存中为待运算的数据以及需要存放结果的变量开辟显存空间。( cudaMalloc 函数实现 )

3. 将待运算的数据传输进显存。( cudaMemcpy，cublasSetVector 等函数实现 )

4. 调用 device 端函数，同时要将需要为 device 端函数创建的块数线程数等参数传递进 <<<>>>。( 注: <<<>>>下方编译器可能显示语法错误，不用管 )

5. 从显存中获取结果变量。( cudaMemcpy，cublasGetVector 等函数实现 )

6. 释放申请的显存空间。( cudaFree 实现 )

函数声明

　　　　1. __device__

　　　　表明此函数只能在 GPU 中被调用，在 GPU 中执行。这类函数只能被 __global__ 类型函数或 __device__ 类型函数调用。

　　　　2. __global__

　　　　表明此函数在 CPU 上调用，在 GPU 中执行。这也是以后会常提到的 "内核函数"，有时为了便于理解也称 "device" 端函数。

　　　　3. __host__

　　　　表明此函数在 CPU 上调用和执行，这也是默认情况。
　　内核函数配置运算符 <<<>>> - 这个运算符在调用内核函数的时候使用，一般情况下传递进三个参数：

　　　　1. 块数

　　　　2. 线程数

　　　　3. 共享内存大小 (此参数默认为0 )

几个内置变量

　　 threadIdx，顾名思义获取线程thread的ID索引；如果线程是一维的那么就取threadIdx.x，二维的还可以多取到一个值threadIdx.y，以此类推到三维 threadIdx.z。
blockIdx，线程块的ID索引；同样有blockIdx.x，blockIdx.y，blockIdx.z。
blockDim，线程块的维度，同样有blockDim.x，blockDim.y，blockDim.z。
gridDim，线程格的维度，同样有gridDim.x，gridDim.y，gridDim.z。

        5. 对于一维的block，线程的threadID=threadIdx.x。
        6. 对于大小为（blockDim.x, blockDim.y）的二维 block，线程的threadID=threadIdx.x+threadIdx.y*blockDim.x。
        7. 对于大小为（blockDim.x, blockDim.y, blockDim.z）的三维 block，线程的threadID=threadIdx.x+threadIdx.y*blockDim.x+threadIdx.z*blockDim.x*blockDim.y。
        8. 对于计算线程索引偏移增量为已启动线程的总数。如stride = blockDim.x * gridDim.x; threadId += stride。

GPU内存

全局内存

通俗意义上的设备内存

共享内存

使用__shared__关键字声明，例如__shared__ float cache[10]，对于GPU上启动的每个线程块，CUDA C编译器都将创建该共享变量的一个副本。线程块中的每个线程都共享这块内存，但线程却无法看到也不能修改其他线程块的变量副本。这样使得一个线程块中的多个线程能够在计算上通信和协作。

常量内存

使用关键字__constant__声明，为了提升性能。常量内存采取了不同于标准全局内存的处理方式。在某些情况下，用常量内存替换全局内存能有效地减少内存带宽。常量内存用于保存在核函数执行期间不会发生变化的数据。变量的访问限制为只读。NVIDIA硬件提供了64KB的常量内存。不再需要cudaMalloc()或者cudaFree(),而是在编译时，静态地分配空间。当我们需要拷贝数据到常量内存中应该使用cudaMemcpyToSymbol()，而cudaMemcpy()会复制到全局内存。

从常量内存中读取数据可以节约内存带宽，主要有两个原因：

对常量内存的单次读操作可以广播到邻近线程，这将节约约15次读取操作
常量内存的数据将缓存起来，因此对相同地址的连续操作将不会产生额外的内存通信量

纹理内存

固定内存

常用函数

cudaGetDeviceCount()　　获取显示设备数目

cudaGetDeviceProperties()　　获取设备属性

cudaChooseDevice()　　根据指定的属性条件选择设备

cudaSetDevice()　　指定使用的显示设备

cudaMalloc()　　在设备中分配空间

cudaMemcpy()　　host和device之间拷贝内存，

cudaFree()　　释放显存

__syncthreads()　　用于同一线程块内线程间的同步，__syncthreads() is you garden variety thread barrier. Any thread reaching the barrier waits until all of the other threads in that block also reach it.

矢量求和

 1 #include "cuda_runtime.h"
 2 #include "cuda.h"
 3 #include "device_launch_parameters.h"
 4 
 5 #include <iostream>
 6 #include <cstdlib>
 7 #define N 10
 8 
 9 __global__ void add(int* a, int* b, int*c)
10 {
11     int tid = blockIdx.x;
12     if (tid < N)
13     {
14         c[tid] = a[tid] + b[tid];
15     }
16 }
17 
18 int main(void)
19 {
20     int a[N], b[N], c[N];
21     int *dev_a, *dev_b, *dev_c;
22     cudaMalloc((void**)&dev_a, N * sizeof(int));
23     cudaMalloc((void**)&dev_b, N * sizeof(int));
24     cudaMalloc((void**)&dev_c, N * sizeof(int));
25     for (int i = 0; i < N; ++i)
26     {
27         a[i] = -i;
28         b[i] = i*i;
29     }
30     cudaMemcpy(dev_a, a, N * sizeof(int), cudaMemcpyHostToDevice);
31     cudaMemcpy(dev_b, b, N * sizeof(int), cudaMemcpyHostToDevice);
32     add << <N, 1 >> > (dev_a, dev_b, dev_c);
33     cudaMemcpy(c, dev_c, N * sizeof(int), cudaMemcpyDeviceToHost);
34     for (int i = 0; i < N; ++i)
35     {
36         printf("%d+%d=%d\n", a[i], b[i], c[i]);
37     }
38     cudaFree(dev_a);
39     cudaFree(dev_b);
40     cudaFree(dev_c);
41 
42     getchar();
43     return 0;
44 }

View Code

调用核函数<<<>>>中

第一个参数表示设备在执行核函数时使用的并行线程块数量，即创建多少个核函数的副本并以并行的方式执行它们。内置变量blockIdx包含的值就是当前执行设备代码的线程块的索引。硬件限制线程块数量不能超过65535，内置变量blockDim保存的是三维的线程块中线程的维度。即CUDA运行时允许启动一个二维线程格，且线程格中的每个线程块都是一个三维的线程数组。

第二个参数表示CUDA运行时在每个线程块中创建的线程数量，内置参数threadIdx为线程索引。硬件限制每个线程块中线程数量不能超过设备属性结构中maxThreadsPerBlock的值。

并行线程块集合也称为一个线程格Grid。

事件

cuda中的事件本质上是一个GPU时间戳

  1 /*
  2 * Copyright 1993-2010 NVIDIA Corporation.  All rights reserved.
  3 *
  4 * NVIDIA Corporation and its licensors retain all intellectual property and
  5 * proprietary rights in and to this software and related documentation.
  6 * Any use, reproduction, disclosure, or distribution of this software
  7 * and related documentation without an express license agreement from
  8 * NVIDIA Corporation is strictly prohibited.
  9 *
 10 * Please refer to the applicable NVIDIA end user license agreement (EULA)
 11 * associated with this source code for terms and conditions that govern
 12 * your use of this NVIDIA software.
 13 *
 14 */
 15 
 16 
 17 #include "cuda.h"
 18 #include "../common/book.h"
 19 #include "../common/cpu_bitmap.h"
 20 
 21 #define DIM 1024
 22 
 23 #define rnd( x ) (x * rand() / RAND_MAX)
 24 #define INF     2e10f
 25 
 26 struct Sphere {
 27     float   r, b, g;
 28     float   radius;
 29     float   x, y, z;
 30     __device__ float hit(float ox, float oy, float *n) {
 31         float dx = ox - x;
 32         float dy = oy - y;
 33         if (dx*dx + dy*dy < radius*radius) {
 34             float dz = sqrtf(radius*radius - dx*dx - dy*dy);
 35             *n = dz / sqrtf(radius * radius);
 36             return dz + z;
 37         }
 38         return -INF;
 39     }
 40 };
 41 #define SPHERES 200
 42 
 43 __constant__ Sphere s[SPHERES];
 44 
 45 __global__ void kernel(unsigned char *ptr) {
 46     // map from threadIdx/BlockIdx to pixel position
 47     int x = threadIdx.x + blockIdx.x * blockDim.x;
 48     int y = threadIdx.y + blockIdx.y * blockDim.y;
 49     int offset = x + y * blockDim.x * gridDim.x;
 50     float   ox = (x - DIM / 2);
 51     float   oy = (y - DIM / 2);
 52 
 53     float   r = 0, g = 0, b = 0;
 54     float   maxz = -INF;
 55     for (int i = 0; i<SPHERES; i++) {
 56         float   n;
 57         float   t = s[i].hit(ox, oy, &n);
 58         if (t > maxz) {
 59             float fscale = n;
 60             r = s[i].r * fscale;
 61             g = s[i].g * fscale;
 62             b = s[i].b * fscale;
 63             maxz = t;
 64         }
 65     }
 66 
 67     ptr[offset * 4 + 0] = (int)(r * 255);
 68     ptr[offset * 4 + 1] = (int)(g * 255);
 69     ptr[offset * 4 + 2] = (int)(b * 255);
 70     ptr[offset * 4 + 3] = 255;
 71 }
 72 
 73 // globals needed by the update routine
 74 struct DataBlock {
 75     unsigned char   *dev_bitmap;
 76 };
 77 
 78 int main(void) {
 79     DataBlock   data;
 80     // capture the start time
 81     cudaEvent_t     start, stop;
 82     HANDLE_ERROR(cudaEventCreate(&start));
 83     HANDLE_ERROR(cudaEventCreate(&stop));
 84     HANDLE_ERROR(cudaEventRecord(start, 0));
 85 
 86     CPUBitmap bitmap(DIM, DIM, &data);
 87     unsigned char   *dev_bitmap;
 88 
 89     // allocate memory on the GPU for the output bitmap
 90     HANDLE_ERROR(cudaMalloc((void**)&dev_bitmap,
 91         bitmap.image_size()));
 92 
 93     // allocate temp memory, initialize it, copy to constant
 94     // memory on the GPU, then free our temp memory
 95     Sphere *temp_s = (Sphere*)malloc(sizeof(Sphere) * SPHERES);
 96     for (int i = 0; i<SPHERES; i++) {
 97         temp_s[i].r = rnd(1.0f);
 98         temp_s[i].g = rnd(1.0f);
 99         temp_s[i].b = rnd(1.0f);
100         temp_s[i].x = rnd(1000.0f) - 500;
101         temp_s[i].y = rnd(1000.0f) - 500;
102         temp_s[i].z = rnd(1000.0f) - 500;
103         temp_s[i].radius = rnd(100.0f) + 20;
104     }
105     HANDLE_ERROR(cudaMemcpyToSymbol(s, temp_s,
106         sizeof(Sphere) * SPHERES));
107     free(temp_s);
108 
109     // generate a bitmap from our sphere data
110     dim3    grids(DIM / 16, DIM / 16);
111     dim3    threads(16, 16);
112     kernel << <grids, threads >> >(dev_bitmap);
113 
114     // copy our bitmap back from the GPU for display
115     HANDLE_ERROR(cudaMemcpy(bitmap.get_ptr(), dev_bitmap,
116         bitmap.image_size(),
117         cudaMemcpyDeviceToHost));
118 
119     // get stop time, and display the timing results
120     HANDLE_ERROR(cudaEventRecord(stop, 0));
121     HANDLE_ERROR(cudaEventSynchronize(stop));
122     float   elapsedTime;
123     HANDLE_ERROR(cudaEventElapsedTime(&elapsedTime,
124         start, stop));
125     printf("Time to generate:  %3.1f ms\n", elapsedTime);
126 
127     HANDLE_ERROR(cudaEventDestroy(start));
128     HANDLE_ERROR(cudaEventDestroy(stop));
129 
130     HANDLE_ERROR(cudaFree(dev_bitmap));
131 
132     // display
133     bitmap.display_and_exit();
134 }