cuda
将一张输入图像缩小到指定大小的输出图像
#include "ImageCuda.h"
#include <cuda_runtime.h>
__global__ void downsample(uchar* srcImg,
uchar* desImg,
int s_width,
int s_height,
int d_width,
int d_height,
float scale_width,
float scale_height)
{
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
if (x < d_width && y < d_height) {
int s_x = int(scale_width * (x + 0.5));
int s_y = int(scale_height * (y + 0.5));
s_x = (s_x < s_width) ? s_x : s_width - 1;
s_y = (s_y < s_height) ? s_y : s_height - 1;
desImg[(y * d_width + x) * 3] = srcImg[(s_y * s_width + s_x) * 3];
desImg[(y * d_width + x) * 3 + 1] = srcImg[(s_y * s_width + s_x) * 3 + 1];
desImg[(y * d_width + x) * 3 + 2] = srcImg[(s_y * s_width + s_x) * 3 + 2];
}
}
实现了一个CUDA核函数 downsample
,用于将一张输入图像缩小到指定大小的输出图像。具体来说,输入图像以 srcImg
表示,输出图像以 desImg
表示,输入图像的大小为 s_width
× s_height
,输出图像的大小为 d_width
× d_height
,缩放比例分别为 scale_width
和 scale_height
。
该核函数中,每个线程负责输出图像中的一个像素值。首先,线程计算输出图像中对应的像素位置 (x, y)
。然后,通过缩放比例计算对应的输入图像中的像素位置 (s_x, s_y)
。由于输入图像和输出图像的像素数不同,因此需要进行插值处理。这里采用最近邻插值的方式,即将输入图像中距离输出像素最近的像素值赋给输出像素。
网格概念
当我们在CUDA程序中启动一个核函数时,需要指定一个网格(grid)和块(block)的大小。网格是由多个块组成的,每个块中又有多个线程。
下面是一个简单的示意图,展示了一个包含4个块和每个块包含4个线程的网格:
|<-------blockDim.x--------->|
|<-threadIdx.x->
+----------------------------+-------------------------------+
| Block (0,0) | Block (1,0) |
| | |
| Thread (0,0) Thread (0,1)| Thread (0,0) Thread (1,0) |
| Thread (1,0) Thread (1,1)| Thread (0,1) Thread (1,1) |
+----------------------------+-------------------------------+
| Block (0,1) | Block (1,1) |
| | |
| Thread (0,0) Thread (0,1)| Thread (0,0) Thread (1,0) |
| Thread (1,0) Thread (1,1)| Thread (0,1) Thread (1,1) |
+----------------------------+-------------------------------+
每个线程都有一个唯一的ID,它的计算方式是blockIdx.x * blockDim.x + threadIdx.x
,这个ID就对应着一个像素的位置。在这个示意图中,假设每个块中有4个线程,那么每个块就可以处理4个像素。整个网格中有16个线程,因此可以处理16个像素。可以看到,通过网格和块的划分,我们可以有效地利用CUDA的并行计算能力,同时也方便地控制并行计算的粒度和规模。
blockDim
是一个编译时常量,它是一个dim3
类型的变量,用于表示每个块中的线程数。在示意图中,假设每个块中有2个线程,则blockDim
应该是(2, 2, 1)
。blockDim
通常在程序中被定义为一个常量,例如:
#define BLOCK_SIZE 16
dim3 blockDim(BLOCK_SIZE, BLOCK_SIZE, 1);
在这个示例中,每个块中有BLOCK_SIZE * BLOCK_SIZE
个线程。在实际应用中,blockDim
的值通常需要根据具体的算法和硬件设备进行调整。
cuda架构
在CUDA中,程序的执行是由多个线程块(blocks)组成的,每个线程块中又包含多个线程(threads)。线程块和线程是CUDA执行模型的两个重要概念,它们是CUDA中并行计算的基本单位。
在CUDA程序执行时,线程块会被分配到CUDA设备上的不同计算单元上并发执行,而线程则被映射到各个计算单元上执行计算任务。每个线程块中的线程是可以互相通信的,同一个线程块中的线程可以共享共享内存(Shared Memory)。
因此,可以将整个CUDA执行模型视为由多个线程块和线程组成的。线程块是CUDA程序的最小执行单位,而线程则是执行具体计算任务的基本单位。在CUDA程序中,通常会将一个问题分解成多个线程块,每个线程块中再由多个线程协同完成计算任务。
代码分解
核函数很像一个for_each_func,告诉每个像素(线程)要做什么
//定义了一个CUDA核函数 downsample,包括8个参数:输入图像 srcImg,输出图像 desImg,输入图像的宽度 s_width 和高度 s_height,
//输出图像的宽度 d_width 和高度 d_height,以及缩放比例 scale_width 和 scale_height。
__global__ void downsample(uchar* srcImg,
uchar* desImg,
int s_width,
int s_height,
int d_width,
int d_height,
float scale_width,
float scale_height)
{
//计算当前线程处理的像素在输出图像中的坐标 (x, y)。blockIdx.x 和 blockIdx.y 表示当前线程所在的线程块在x和y维度上的索引,
//blockDim.x 和 blockDim.y 则表示线程块的大小。threadIdx.x 和 threadIdx.y 则是当前线程在线程块内的索引。
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
//判断当前线程是否越界,如果越界则不进行任何操作。如果不越界,则计算当前像素在输入图像中的坐标 (s_x, s_y)。
//scale_width 和 scale_height 表示输入图像和输出图像之间的缩放比例,加上0.5是为了进行四舍五入处理。
//然后,如果计算得到的输入图像中的像素位置 (s_x, s_y) 超出了输入图像的范围,则将它设置为输入图像中最后一个像素的位置。
if (x < d_width && y < d_height) {
int s_x = int(scale_width * (x + 0.5));
int s_y = int(scale_height * (y + 0.5));
s_x = (s_x < s_width) ? s_x : s_width - 1;
s_y = (s_y < s_height) ? s_y : s_height - 1;
//根据当前像素在输入图像中的坐标 (s_x, s_y),将对应的像素值复制到输出图像的相应位置。由于每个像素包含三个颜色通道,所以需要乘以3。
desImg[(y * d_width + x) * 3] = srcImg[(s_y * s_width + s_x) * 3];
desImg[(y * d_width + x) * 3 + 1] = srcImg[(s_y * s_width + s_x) * 3 + 1];
desImg[(y * d_width + x) * 3 + 2] = srcImg[(s_y * s_width + s_x) * 3 + 2];
}
}
异构编程
异构编程是指在一个计算系统中,利用多种不同类型的处理器或协处理器来共同完成一些任务,这些处理器或协处理器在架构、功能和特性上可能存在很大的差异。CUDA的异构编程就是利用CPU和GPU协同完成计算任务,其中CPU和GPU的架构、功能和特性都有所不同。
CPU是一种通用处理器,适用于各种类型的应用程序。CPU的主要优势是其强大的控制单元和灵活的数据处理能力,可以执行各种类型的指令,同时提供高度灵活的编程模型和多种编程语言。
GPU是一种专用处理器,适用于并行计算和大规模数据处理。GPU的主要优势是其高度并行的架构和大规模的数据处理能力,可以同时处理数百个线程,从而实现高效的并行计算。但是,由于GPU的架构和编程模型不同于CPU,因此需要专门的编程技术和工具来进行GPU编程。
在CUDA的异构编程中,程序员可以将任务分配到CPU和GPU之间,以充分利用两者的优势。一般来说,CPU用于串行代码的执行和控制流程,而GPU则用于并行计算和数据处理。CUDA提供了一种编程模型,使程序员可以轻松地将CPU和GPU之间的数据传输和任务调度进行管理,从而实现更高效的计算和更快的运算速度。
一个简单的异构编程例子是在CPU和GPU之间实现向量加法。具体的实现步骤如下:
- 在CPU上生成两个向量,并将它们存储在主机内存中。
- 分配两个向量的空间,并将它们复制到GPU设备内存中。
- 在GPU上启动核函数,对两个向量进行逐元素相加。
- 将结果从GPU设备内存复制回主机内存。
- 释放GPU设备内存并删除所有分配的对象。
以下是一个简单的CUDA C++代码示例,演示了如何在CPU和GPU之间执行向量加法:
#include <stdio.h>
#include <stdlib.h>
#include <cuda_runtime.h>
void vector_add_cpu(float* a, float* b, float* c, int n) {
for (int i = 0; i < n; i++) {
c[i] = a[i] + b[i];
}
}
__global__ void vector_add_gpu(float* a, float* b, float* c, int n) {
int idx = blockDim.x * blockIdx.x + threadIdx.x;
if (idx < n) {
c[idx] = a[idx] + b[idx];
}
}
int main() {
int n = 1000000;
float *a, *b, *c;
float *d_a, *d_b, *d_c;
// Allocate host memory
a = (float*) malloc(n * sizeof(float));
b = (float*) malloc(n * sizeof(float));
c = (float*) malloc(n * sizeof(float));
// Initialize host arrays
for (int i = 0; i < n; i++) {
a[i] = i;
b[i] = i + 1;
}
// Allocate device memory
cudaMalloc(&d_a, n * sizeof(float));
cudaMalloc(&d_b, n * sizeof(float));
cudaMalloc(&d_c, n * sizeof(float));
// Copy host data to device
cudaMemcpy(d_a, a, n * sizeof(float), cudaMemcpyHostToDevice);
cudaMemcpy(d_b, b, n * sizeof(float), cudaMemcpyHostToDevice);
// Launch kernel on device
int threads_per_block = 256;
int num_blocks = (n + threads_per_block - 1) / threads_per_block;
vector_add_gpu<<<num_blocks, threads_per_block>>>(d_a, d_b, d_c, n);
// Copy result from device to host
cudaMemcpy(c, d_c, n * sizeof(float), cudaMemcpyDeviceToHost);
// Check results
vector_add_cpu(a, b, c, n);
for (int i = 0; i < n; i++) {
if (c[i] != a[i] + b[i]) {
printf("Error at index %d: %f + %f != %f\n", i, a[i], b[i], c[i]);
break;
}
}
// Free memory
free(a);
free(b);
free(c);
cudaFree(d_a);
cudaFree(d_b);
cudaFree(d_c);
return 0;
}
CUDA的核函数启动语法。vector_add_gpu
是要在GPU上执行的核函数,<<<num_blocks, threads_per_block>>>
用于指定执行核函数的线程块和线程数。其中,num_blocks
是指总的线程块数,threads_per_block
是指每个线程块中包含的线程数。这种语法可以在主机端启动核函数,在设备端执行核函数,并返回结果。
g++ -o my_program my_program.cpp -L/usr/local/cuda/lib64 -lcudart
常用函数
- CUDA内存管理函数:包括
cudaMalloc()
,cudaMemcpy()
,cudaFree()
等函数,用于在GPU上分配、拷贝和释放内存。 - CUDA流(Stream):一个抽象的执行上下文,用于管理CUDA的异步执行。可以使用
cudaStreamCreate()
创建一个流,使用cudaStreamDestroy()
销毁一个流,使用cudaStreamSynchronize()
等函数控制流的同步和异步执行。 - CUDA事件(Event):用于测量CUDA程序的执行时间和控制程序的执行顺序。可以使用
cudaEventCreate()
创建一个事件,使用cudaEventRecord()
记录一个事件的发生,使用cudaEventSynchronize()
等函数控制事件的同步和异步执行。 - CUDA线程块同步函数:包括
__syncthreads()
函数,用于同步一个线程块中的所有线程的执行。 - CUDA的并行计算工具库:包括cuBLAS(用于矩阵乘法等运算)、cuFFT(用于快速傅里叶变换)、cuRAND(用于随机数生成)等库。
学习资料
- CUDA官方文档:https://docs.nvidia.com/cuda/index.html
- 《CUDA C Programming Guide》(CUDA C编程指南),CUDA官方出版物,介绍了CUDA编程的基础知识和高级特性。可以在官网下载或者在Amazon购买。
- Udacity在线课程:https://www.udacity.com/course/intro-to-parallel-programming–cs344
- Coursera在线课程:https://www.coursera.org/learn/heterogeneous-parallel-programming
- 《CUDA by Example: An Introduction to General-Purpose GPU Programming》(通过实例学习CUDA:一种通用GPU编程的介绍),Jason Sanders和Edward Kandrot著,是一本以实例为主的CUDA编程教材,适合初学者阅读。
文章评论