Featured image of post cuda_c_basic
cuda

cuda_c_basic

cuda基本概念和编程

引言

CUDA（Compute Unified Device Architecture）是 NVIDIA 推出的并行计算平台和编程模型，允许开发者利用 GPU（图形处理器）进行高性能通用计算（GPGPU）。它扩展了 C/C++ 等语言，使开发者能够编写直接在 GPU 上运行的代码，大幅加速计算密集型任务。

CUDA 核心概念

(1) 主机（Host）与设备（Device）

  • Host（主机）：CPU 及其内存（如 malloc 分配的内存）。
  • Device（设备）：GPU 及其显存（如 cudaMalloc 分配的内存）。
  • 数据传输：主机和设备之间需要通过 cudaMemcpy 显式传输数据。

(2) 核函数（Kernel）

  • GPU 上并行执行的函数，用

    1

    __global__ void(...) {}

    1
2
3
4

    __global__ void addKernel(int *a, int *b, int *c) {
    int i = threadIdx.x;
    c[i] = a[i] + b[i];
}

  • 调用方式：

    1

    <<<grid, block>>>

    指定并行执行配置：

    1

    addKernel<<<1, N>>>(d_a, d_b, d_c); // 1个block，N个threads

(3) 线程层次

  • 线程（Thread）：最基本的执行单元。
  • 线程块（Block）：一组线程，可共享 __shared__ 内存，可同步（__syncthreads()）。
  • 网格（Grid）：多个线程块的集合。
  • 索引计算：
    • threadIdx.x：线程在 block 内的索引。
    • blockIdx.x：block 在 grid 内的索引。
    • blockDim.x：block 的大小（线程数）。

CUDA 执行流程

  1. 分配设备内存

    1
2

    int *d_a;
cudaMalloc(&d_a, size);

  2. 主机→设备数据传输

    1

    cudaMemcpy(d_a, h_a, size, cudaMemcpyHostToDevice);

  3. 启动核函数

    1

    kernel<<<grid, block>>>(d_a, d_b, d_c);

  4. 设备→主机数据传输

    1

    cudaMemcpy(h_c, d_c, size, cudaMemcpyDeviceToHost);

  5. 释放设备内存

    1

    cudaFree(d_a);

向量加法(一维)执行流程 

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60

#include <stdio.h>

#include <vector>
#include <iostream>

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#define cudaCheckError(msg) \
	do { \
		cudaError_t __err = cudaGetLastError(); \
		if (__err  != cudaSuccess) { \
			fprintf(stderr, "Fatal error: %s (%s at %s:%d)\n", \
				msg, cudaGetErrorString(__err), \
				__FILE__, __LINE__); \
			fprintf(stderr, "*** FAILED - ABORTING\n"); \
			exit(1); \
		} \
	} while(0)

const int DSIZE = 4096;

__global__ void vector_add(const int* a, const int* b, int* c, int size) {
	int index = blockIdx.x * blockDim.x + threadIdx.x;	// 确定唯一id
	if (index < size) {
		c[index] = a[index] + b[index];
	}
}

int main() {
	std::vector<int> v1 = { 1,2,3,4,5 };
	std::vector<int> v2 = { 10,20,30,40,50 };
	std::vector<int> v3 = { 0,0,0,0,0 };

	int* d_a, * d_b, * d_c;	// device数据拷贝
	cudaMalloc((void**)&d_a, v1.size() * sizeof(int));
	cudaMalloc((void**)&d_b, v2.size() * sizeof(int));
	cudaMalloc((void**)&d_c, v3.size() * sizeof(int));
	cudaCheckError("cudaMalloc failuer");

	cudaMemcpy(d_a, v1.data(), v1.size() * sizeof(int), cudaMemcpyHostToDevice);
	cudaMemcpy(d_b, v2.data(), v2.size() * sizeof(int), cudaMemcpyHostToDevice);
	cudaCheckError("cudaMemcpy failuer");

	vector_add<<<(v1.size() + 512 - 1) / 512, 512>>>(d_a, d_b, d_c, v1.size());
	cudaCheckError("vector_add failuer");

	cudaMemcpy(v3.data(), d_c, v3.size() * sizeof(int), cudaMemcpyDeviceToHost);
	cudaCheckError("cudaMemcpy from device failuer");

	for (auto v : v3) {
		std::cout << v << " ";
	}
	std::cout << std::endl;
	
    // 手动释放
	cudaFree(d_a);
	cudaFree(d_b);
	cudaFree(d_c);
}
  • 通过RAII管理device内存，同时实现模板支撑更多的数据类型
 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76

#include <stdio.h>

#include <vector>
#include <iostream>

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#define cudaCheckError(msg) \
	do { \
		cudaError_t __err = cudaGetLastError(); \
		if (__err  != cudaSuccess) { \
			fprintf(stderr, "Fatal error: %s (%s at %s:%d)\n", \
				msg, cudaGetErrorString(__err), \
				__FILE__, __LINE__); \
			fprintf(stderr, "*** FAILED - ABORTING\n"); \
			exit(1); \
		} \
	} while(0)

const int DSIZE = 4096;

__global__ void vector_add(const int* a, const int* b, int* c, int size) {
	int index = blockIdx.x * blockDim.x + threadIdx.x;
	if (index < size) {
		c[index] = a[index] + b[index];
	}
}

// RAII进行管理,管理device的内存
// 通过模板支持更多数据类型
template<typename T>
class DeviceData {
public:
	DeviceData(size_t size) {
		cudaMalloc((void**)&data_, size);
	}

	~DeviceData() {
		cudaFree(data_);
	}

	int* getData() {
		return data_;
	}

private:
	T* data_;
};


int main() {
	std::vector<int> v1 = { 1,2,3,4,5 };
	std::vector<int> v2 = { 10,20,30,40,50 };
	std::vector<int> v3 = { 0,0,0,0,0 };

	DeviceData<int> d_a(v1.size() * sizeof(int));
	DeviceData<int> d_b(v2.size() * sizeof(int));
	DeviceData<int> d_c(v3.size() * sizeof(int));
	cudaCheckError("cudaMalloc failuer");

	cudaMemcpy(d_a.getData(), v1.data(), v1.size() * sizeof(int), cudaMemcpyHostToDevice);
	cudaMemcpy(d_b.getData(), v2.data(), v2.size() * sizeof(int), cudaMemcpyHostToDevice);
	cudaCheckError("cudaMemcpy failuer");

	vector_add<<<(v1.size() + 512 - 1) / 512, 512>>>(d_a.getData(), d_b.getData(), d_c.getData(), v1.size());
	cudaCheckError("vector_add failuer");

	cudaMemcpy(v3.data(), d_c.getData(), v3.size() * sizeof(int), cudaMemcpyDeviceToHost);
	cudaCheckError("cudaMemcpy from device failuer");

	for (auto v : v3) {
		std::cout << v << " ";
	}
	std::cout << std::endl;
}

矩阵乘法（多维）执行流程 

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <iostream>

#include <stdio.h>
#include <stdlib.h>

#define cudaCheckError(msg) \
	do { \
		cudaError_t __err = cudaGetLastError(); \
		if (__err  != cudaSuccess) { \
			fprintf(stderr, "Fatal error: %s (%s at %s:%d)\n", \
				msg, cudaGetErrorString(__err), \
				__FILE__, __LINE__); \
			fprintf(stderr, "*** FAILED - ABORTING\n"); \
			exit(1); \
		} \
	} while(0)

const int dy = 2;		
const int dx = 3;

// A 是dy * dx, B 是dx * dy, C 是dy * dy
// C = AB 实现
__global__ void mmul(const float* A, const float* B, float* C, int dy, int dx) {
	// m * k 和 k * n  矩阵乘法的通用写法
	int m = dy, k = dx, n = dy;
	// 计算当前线程对应的行和列

	// 可以认为，只用row和col能用来表示C中元素的线程才会进行计算
	// 不会发生线程的冗余计算
	int row = blockIdx.y * blockDim.y + threadIdx.y;
	int col = blockIdx.x * blockDim.x + threadIdx.x;

	// C[row][col] 为 A的第row行 * B的第col列
	if (row < m && col < n) {
		float sum = 0.0f;
		for (int i = 0; i < k; i++) {
			sum += A[row * k + i] * B[i * n + col]; // A 行主序，B 列主序
		}
		C[row * n + col] = sum; // C 行主序
	}
}

int main() {
	float* h_A, * h_B, * h_C, * d_A, * d_B, * d_C;
	h_A = new float[dy * dx];
	h_B = new float[dx * dy];
	h_C = new float[dy * dy];
	h_A[0] = 1.0f; h_A[1] = 2.0f; h_A[2] = 1.0f; 
	h_A[3] = 2.0f; h_A[4] = 2.0f; h_A[5] = 1.0f;

	h_B[0] = 1.0f; h_B[1] = 2.0f;
	h_B[2] = 2.0f; h_B[3] = 1.0f;
	h_B[4] = 2.0f; h_B[5] = 2.0f;

	cudaMalloc((void**)&d_A, dy * dx * sizeof(float));
	cudaMalloc((void**)&d_B, dx * dy * sizeof(float));
	cudaMalloc((void**)&d_C, dy * dy * sizeof(float));
	cudaCheckError("cudaMalloc failuer");

	cudaMemcpy(d_A, h_A, dy * dx * sizeof(float), cudaMemcpyHostToDevice);
	cudaMemcpy(d_B, h_B, dx * dy * sizeof(float), cudaMemcpyHostToDevice);
	cudaCheckError("cudaMemcpy error");

	// 设备大小
	dim3 block(2, 2);		// 每个block有2*2个线程

	// 通用公式
	// grid((结果矩阵列数 + block.x - 1) / block.x, 
	//	    (结果矩阵行数 + block.y - 1) / block.y);
	dim3 grid((dy + block.x - 1) / block.x, (dy + block.y - 1) / block.y);	// 用来定义grid中block的数量

	mmul << <grid, block >> > (d_A, d_B, d_C, dy, dx);
	cudaCheckError("kernel launch failure");

	cudaMemcpy(h_C, d_C, dy * dy * sizeof(float), cudaMemcpyDeviceToHost);

	for (int i = 0; i < 4; ++i) {
		std::cout << h_C[i] << " ";
		if ((i + 1) % 2 == 0) {
			std::cout << std::endl;
		}
	}

	cudaFree(d_A);
	cudaFree(d_B);
	cudaFree(d_C);
	delete[] h_A;
	delete[] h_B;
	delete[] h_C;
}

注意事项

  • 在对host主机的代码中，可以使用标准库，包括shared_ptr, unique_ptr等；但是在核函数当中，对device设备编程，则无法使用标准库里的东西。
  • cuda的device可以支持模板以实现更通用的功能
  • 可以自己实现RAII对device的资源进行管理