CUDA实验:数组相加

本文主要实现向量相加的cuda编程。
首先写出整体框架:

#include "error.cuh"
#include <stdio.h>
#include <cuda_runtime.h>
#define N   (33 * 1024)

__global__ void add( int *a, int *b, int *c ) {

// 此处为待填充区域

}

int main( void ) {
    //定义主机指针和设备指针
    int *a, *b, *c;
    int *dev_a, *dev_b, *dev_c;

    //CPU上动态分配内存
    a = (int*)malloc( N * sizeof(int) );
    b = (int*)malloc( N * sizeof(int) );
    c = (int*)malloc( N * sizeof(int) );

    //GPU上动态分配内存
    CHECK( cudaMalloc( (void**)&dev_a, N * sizeof(int) ) );
    CHECK( cudaMalloc( (void**)&dev_b, N * sizeof(int) ) );
    CHECK( cudaMalloc( (void**)&dev_c, N * sizeof(int) ) );

    //CPU上为a和b赋值
    for (int i=0; i<N; i++) {
        a[i] = i;
        b[i] = 2 * i;
    }

    //将a和b的数据从主机拷贝到设备
    CHECK( cudaMemcpy( dev_a, a, N * sizeof(int),
                              cudaMemcpyHostToDevice ) );
    CHECK( cudaMemcpy( dev_b, b, N * sizeof(int),
                              cudaMemcpyHostToDevice ) );
    //调用核函数,设置线程,网格大小1D,线程块大小1D,32
    add<<<128,128>>>( dev_a, dev_b, dev_c );

    //将数据从设备传递到主机上
    CHECK( cudaMemcpy( c, dev_c, N * sizeof(int),
                              cudaMemcpyDeviceToHost ) );

    //结果校验
    bool success = true;
    for (int i=0; i<N; i++) {
        if ((a[i] + b[i]) != c[i]) {
            success = false;
        }
    }
    if (success)
	printf( "We did it!\n" );
    else
	printf("Not Pass!Error!!!\n");

    //释放GPU内存空间
    CHECK( cudaFree( dev_a ) );
    CHECK( cudaFree( dev_b ) );
    CHECK( cudaFree( dev_c ) );

    //释放CPU内存空间
    free( a );
    free( b );
    free( c );

    return 0;
}

当数组大小<网格大小 * 线程块大小时的数组相加非常好理解,详见书本。本文主要分析当数组大小 >= 网格大小 * 线程块大小的情况。我们已经知道数据大小、网格大小和线程块大小分别为:33 * 1024 、128、128 。出现上述情况时如果依旧按照原先的方法会导致程序无法运行。因此我们需要对add函数进行改动。
常见的改法有两种:
改法一:

__global__ void add( int *a, int *b, int *c ) {
//获取线程id,网格大小1D,线程块大小1D,gridDim.x=128,blockDim.x=128
//数据量:33792,线程数:16384
	int tid = threadIdx.x + blockIdx.x * blockDim.x;
	for(int i=0;i<N;i++){
		if(tid==i%blockDim.x*gridDim.x){
			c[i] = a[i] + b[i];
		}
	}
	while (tid < N) {
		c[tid] = a[tid] + b[tid];
		tid += blockDim.x * gridDim.x;
	}
}

改法二:

__global__ void add( int *a, int *b, int *c ) {
//获取线程id,网格大小1D,线程块大小1D,gridDim.x=128,blockDim.x=128
//数据量:33792,线程数:16384
	int tid = threadIdx.x + blockIdx.x * blockDim.x;
	while (tid < N) {
		c[tid] = a[tid] + b[tid];
		tid += blockDim.x * gridDim.x;
	}
}

分析:
其中blockDim.x * gridDim.x可以理解为开的线程的总的数目。本文向量的大小为33 * 1024,先对[0,128128-1]内的数据进行并行计算,在计算完后,再对[128128,2128128-1]的数据进行计算,直到将大小为N的向量计算完。

  • 0
    点赞
  • 3
    收藏
    觉得还不错? 一键收藏
  • 0
    评论
评论
添加红包

请填写红包祝福语或标题

红包个数最小为10个

红包金额最低5元

当前余额3.43前往充值 >
需支付:10.00
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明:

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

余额充值