【CUDA编程笔记】(2)CPU与GPU之间的参数传递

利用GPU并行计算的的总体思路是:在CPU(Host)中创建数据,将数据传到GPU(Device)中进行计算,再将计算结果传回到CPU中。

最简单的例子:将CPU中的两个数字在GPU中进行相加,并在CPU中输出:


#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>
#include <iostream>


__device__ int add_gpu(int a, int b) {
	return a + b;
}


__global__ void add(int a, int b, int *c) {
	*c = add_gpu(a , b);
}


int main()
{
	//创建变量
	int a = 3, b = 5;
	int c,int *ptr;

	//分配GPU中的内存
	cudaMalloc((void **)&ptr, sizeof(int)); 

	//在GPU中进行计算
	add << <1, 1 >> > (a, b, ptr);

	//将GPU中的计算结果(ptr指针)复制到CPU主机中,赋给c
	cudaMemcpy(&c, ptr, sizeof(int), cudaMemcpyDeviceToHost);   
	printf("%d + %d = %d\n", a, b, c);

	//释放指针
	cudaFree(ptr);  
	return 0;
}


//3 + 5 = 8

上面的例子只有一个数相加,因此不需要多线程,下面我们加大数据量,引入多线程并行计算。

用10个线程对两个长度为10的数组相加:


#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>
#include <iostream>

const int N = 10;    //数组长度
const int Mem = N * sizeof(int);    //数组内存大小

__global__ void add(int *a, int *b, int *c) {
	int tid = threadIdx.x;
	if (tid < N)
		c[tid] = a[tid] + b[tid];
}


int main()
{
	//创建变量
	int a[N] = { 1,2,3,4,5,6,7,8,9,10 };
	int b[N] = { 1,3,5,7,9,11,13,15,17,19 };
	int c[N];
	int *dev_a, *dev_b, *dev_c;

	//分类GPU内存
	cudaMalloc((void **)&dev_a, Mem);
	cudaMalloc((void **)&dev_b, Mem);
	cudaMalloc((void **)&dev_c, Mem);

	//将数据传给GPU
	cudaMemcpy(dev_a, a, Mem, cudaMemcpyHostToDevice);
	cudaMemcpy(dev_b, b, Mem, cudaMemcpyHostToDevice);

	//在GPU中并行计算
	add << <1, 10 >> > (dev_a, dev_b, dev_c);

	//将计算结果传回CPU
	cudaMemcpy(c, dev_c, Mem, cudaMemcpyDeviceToHost);

	//输出计算结果
	for (int i = 0; i < N; i++)
		printf("%d + %d = %d\n", a[i], b[i], c[i]);

	//释放指针
	cudaFree(dev_a);
	cudaFree(dev_b);
	cudaFree(dev_c);

	return 0;
}


/*
1 + 1 = 2
2 + 3 = 5
3 + 5 = 8
4 + 7 = 11
5 + 9 = 14
6 + 11 = 17
7 + 13 = 20
8 + 15 = 23
9 + 17 = 26
10 + 19 = 29
*/

当然也可以用10个线程块,每个线程块分配一个线程的方式来实现,输出结果是一样的,代码如下:


#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>
#include <iostream>

#define N 1024
const int Mem = N * sizeof(int);    //数组内存大小

__global__ void add(int *a, int *b, int *c) {
	int bid = blockIdx.x;
	if (bid < N)
		c[bid] = a[bid] + b[bid];
}


int main()
{
	//创建变量
	int a[N] = { 1,2,3,4,5,6,7,8,9,10 };
	int b[N] = { 1,3,5,7,9,11,13,15,17,19 };
	int c[N];
	int *dev_a, *dev_b, *dev_c;

	//分类GPU内存
	cudaMalloc((void **)&dev_a, Mem);
	cudaMalloc((void **)&dev_b, Mem);
	cudaMalloc((void **)&dev_c, Mem);

	//将数据传给GPU
	cudaMemcpy(dev_a, a, Mem, cudaMemcpyHostToDevice);
	cudaMemcpy(dev_b, b, Mem, cudaMemcpyHostToDevice);

	//在GPU中并行计算
	add << <10, 1 >> > (dev_a, dev_b, dev_c);

	//将计算结果传回CPU
	cudaMemcpy(c, dev_c, Mem, cudaMemcpyDeviceToHost);

	//输出计算结果
	for (int i = 0; i < N; i++)
		printf("%d + %d = %d\n", a[i], b[i], c[i]);

	//释放指针
	cudaFree(dev_a);
	cudaFree(dev_b);
	cudaFree(dev_c);

	return 0;
}

实际情况中,数据量很大,我们不可能用一个线程只处理一个数据,而是用一个线程处理多个数据。


#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>
#include <iostream>

#define N 1024
const int Mem = N * sizeof(int);    //数组内存大小


__global__ void add(int *a, int *b, int *c) {
	int bid = blockIdx.x;
	while (bid < N) {
		c[bid] = a[bid] + b[bid];
		bid += gridDim.x;
	}
}


int main()
{
	//创建变量
	int a[N],b[N],c[N];
	int *dev_a, *dev_b, *dev_c;
	for (int i = 0; i < N; i++) {
		a[i] = i + 1;
		b[i] = 2 * i;
	}

	//分类GPU内存
	cudaMalloc((void **)&dev_a, Mem);
	cudaMalloc((void **)&dev_b, Mem);
	cudaMalloc((void **)&dev_c, Mem);

	//将数据传给GPU
	cudaMemcpy(dev_a, a, Mem, cudaMemcpyHostToDevice);
	cudaMemcpy(dev_b, b, Mem, cudaMemcpyHostToDevice);

	//在GPU中并行计算
	add << <10, 1 >> > (dev_a, dev_b, dev_c);

	//将计算结果传回CPU
	cudaMemcpy(c, dev_c, Mem, cudaMemcpyDeviceToHost);

	//输出计算结果
	for (int i = 0; i < N; i++)
		printf("%d + %d = %d\n", a[i], b[i], c[i]);

	//释放指针
	cudaFree(dev_a);
	cudaFree(dev_b);
	cudaFree(dev_c);

	return 0;
}
评论 5
添加红包

请填写红包祝福语或标题

红包个数最小为10个

红包金额最低5元

当前余额3.43前往充值 >
需支付:10.00
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包

打赏作者

ctrl A_ctrl C_ctrl V

你的鼓励将是我创作的最大动力

¥1 ¥2 ¥4 ¥6 ¥10 ¥20
扫码支付:¥1
获取中
扫码支付

您的余额不足,请更换扫码支付或充值

打赏作者

实付
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明:

1.余额是钱包充值的虚拟货币,按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载,可以购买VIP、付费专栏及课程。

余额充值