#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>
#include <iostream>
/*
线程块Block的数量限制是:65535
我的机器上maxThreadPerBlock = 1024.
也就是一个Block中最多1024个线程
介于前一篇中,是对长度很小的向量进行计算,所以直接在一个block中开几个并行线程即可
(前两篇分别是多Blcok单Thread--MBST和单Block多Thread--SBMT)
因此,如果我们要对一个长度超过1024的向量进行运算,该如何做呢?
在这种情况下,要将线程和block结合起来才行,本篇将会使用MBMT--多Block和多线程的方法
修改两个地方:核函数中的索引计算方法和核函数的调用方法
*/
/*
注意,计算核函数运行了几次要用原子操作
*/
/*
此例子是线程少,数据多的情况,具有典型性!!!
*/
#define N (33 * 1024)
__global__ void add(int *a , int *b , int* c ,int *n)
{
int tid = threadIdx.x + blockIdx.x * blockDim.x;//获得线程索引
// printf("tid = %d\n",tid);
atomicAdd(n,1); //计算函数运行了几次
while( tid < N) //必须检查是否在正确的偏移范围类内
{
c[ tid ] = a[ tid ] + b[ tid ];
tid += blockDim.x * gridDim.x; // 这一句不是很懂
}
}
int main()
{
int a[N] , b[N] , c[N];
int *dev_a, *dev_b , *dev_c;
int *n = 0;
int count = 2; //存放最后的计算此数结果
//在GPU上分配内存
cudaMalloc( (void**)&dev_a , N * sizeof(int ) );
cudaMalloc( (void**)&dev_b , N * sizeof(int ) );
cudaMalloc( (void**)&dev_c , N * sizeof(int ) );
cudaMalloc( (void**)&n , sizeof(int ) );
//在CPU上为数组a,b赋初值
for(int i = 0; i < N; i++)
{
a[i] = -i;
b[i] = i * i;
}
//将数组a, b 复制到GPU上去计算
cudaMemcpy( dev_a , a , N * sizeof(int) , cudaMemcpyHostToDevice);
cudaMemcpy( dev_b , b , N * sizeof(int) , cudaMemcpyHostToDevice);
add<<<128,128>>>( dev_a , dev_b , dev_c, n);//那么将会有128 * 128 个核函数运行
//计算完毕,将数组C从GPU复制到CPU
cudaMemcpy( c , dev_c , N * sizeof(int) , cudaMemcpyDeviceToHost);
//计算完毕,将n从GPU复制到CPU
cudaMemcpy( &count , n , sizeof(int) , cudaMemcpyDeviceToHost);
//验证GPU确实完成了我们要求的工作
bool success = true;
//显示结果
for(int i = 0; i < N ; i++)
{
if( ( a[i] + b[i]) != c[i] )
{
printf("error\n");
success = false;
}
//printf(" %d + %d = %d \n", a[i] , b[i] , c[i]);
}
if(success)
printf("we did it ! \n");
//最后记得释放在GPU上分配的内存
printf("n = %d", count);
cudaFree( dev_a );
cudaFree( dev_b );
cudaFree( dev_c );
cudaFree( n );
getchar();
return 0;
}
/*
使用线程与使用线程块相比存在哪些优势呢?几乎没有任何优势,
但是,线程块中的并行线程能够完成线程块无法完成的工作。
*/