cuda demo的执行

转载

冷月星 2025-01-06 08:15:05

文章标签 cuda demo的执行 c++ cuda gpu 核函数 文章分类 机器学习人工智能

欢迎大家一起来观摩我学习CUDA编程，为我鼓掌

一.基本使用

示例代码解析
存储器

总结陈词

一.基本使用

这里我礼貌性的介绍一下最基本的操作：

cuda编程最基础的模式就是创建一个用__global__修饰的核函数，之后在主程序中创建设备数据指针和主机数据变量，设备指针先通过cudaMalloc开辟空间，再通过cudaMemcpy进行赋值，调用核函数后再通过cudaMemcpy拷贝结果。

示例代码解析

#include <stdio.h>
#include <cuda_runtime.h>
#include <opencv2/opencv.hpp>

#define NUM_THREADS 2000000000
#define SIZE  10
#define BLOCK_WIDTH 1000

__global__ void gpu_increment_atomic(int *d_a)
{
	int tid = blockIdx.x * blockDim.x + threadIdx.x;
	__shared__ int sh_arr[10];
	if (threadIdx.x < 10) {
		sh_arr[threadIdx.x] = 0;
	}
	__syncthreads();

	tid = tid % SIZE;
	atomicAdd(&sh_arr[tid], 1);
	__syncthreads();
	if (threadIdx.x < 10) {
		atomicAdd(&d_a[threadIdx.x], sh_arr[threadIdx.x]);
	}
}

首先就是一个叫gpu_increment_atomic的核函数，内部先获取当前调用这个核函数的block和thread，然后通过__sharde__修饰一个共享内存数组数组，

__syncthreads()表示等待同步，所有block和thread会等到都执行到此处再一起往下走，之后就是累加，由于共享内存属于block内通用，所以相当于在第二个__syncthreads()时，每个block都得到了共享内存数组累加的结果。在此同步，再将block内的结果累加到输出中。

这里使用atomicAdd()是为了防止线程同时占用，类似于线程锁，不然用i++之类的方式会有问题。（不理解的就自己再思考下，不太想写）

个人感觉CUDA编程最麻烦的就是共享的东西或者同步问题，反正调的很辛酸。

int main(int argc, char **argv)
{
	printf("%d total threads in %d blocks writing into %d array elements\n",
		NUM_THREADS, NUM_THREADS / BLOCK_WIDTH, SIZE);

	// declare and allocate host memory
	int h_a[SIZE];
	const int ARRAY_BYTES = SIZE * sizeof(int);

	// declare and allocate GPU memory
	int * d_a;

	clock_t start = clock();
	cudaMalloc((void **)&d_a, ARRAY_BYTES);
	// Initialize GPU memory to zero
	cudaMemset((void *)d_a, 0, ARRAY_BYTES);

	gpu_increment_atomic << <NUM_THREADS / BLOCK_WIDTH, BLOCK_WIDTH >> >(d_a);
	// cudaDeviceSynchronize();

	// copy back the array to host memory
	
	cudaMemcpy(h_a, d_a, ARRAY_BYTES, cudaMemcpyDeviceToHost);
	double time = (double)(clock() - start) / CLOCKS_PER_SEC;
	
	std::cout << "time" << time << std::endl;
	printf("Number of times a particular Array index has been incremented is: \n");

	for (int i = 0; i < SIZE; i++)
	{
		printf("index: %d --> %d times\n ", i, h_a[i]);
	}

	cudaFree(d_a);
	system("pause");
	return 0;
}

主程序中创建设备数据指针的d_a和主机数据变量h_a，通过cudaMalloc为的d_a开辟空间，再通过cudaMemcpy进行赋值，调用核函数后再通过cudaMemcpy拷贝结果。

调用核函数时<<<NUM_THREADS / BLOCK_WIDTH, BLOCK_WIDTH>>>前面表示使用的block数量，后面表示单个block内的thread数量，你可以理解为总共启动了block_num×thread_num数量的线程，一方面block内有些内存能公用，比较方便，一方面显卡对单个block内的thread数量有限制。稍新的显卡大约是1024。

运算完成之后需要通过cudaFree()释放资源。