接下来,我们给数组分配在设备上需要的内存,正如你所看到的,我们要处理一个二维数组,使用BLOCKS*BLOCKS分配内存:
cudaMalloc( (void**)&deviceArray, BLOCKS * BLOCKS * sizeof(int) );
cudaMemcpy( deviceArray,
hostArray, BLOCKS * BLOCKS * sizeof(int),
cudaMemcpyHostToDevice );
cudaMemcpy( deviceArray,
hostArray, BLOCKS * BLOCKS * sizeof(int),
cudaMemcpyHostToDevice );
获得设备上的空间后,我们开始启动内核,在GPU上执行计算任务。
generateArray<<<multiBlockArray,1>>>( deviceArray );
这里唯一的不同是,我们传递了multiBlockArray参数,它就是我们前面提到的我们想要运行的块数。
接下来我们拷贝GPU上的数组到主机,这样我们才能显示它:
cudaMemcpy( hostArray,
deviceArray,
BLOCKS * BLOCKS * sizeof(int),
cudaMemcpyDeviceToHost );
for (int i=0; i<BLOCKS; i++)
{
printf( “Thread ID running: %d”, hostArray[0][i] );
for (int j=1; j<BLOCKS; j++)
{
printf( ” %d”, hostArray[j][i] );
}
printf( “\n” );
}
cudaFree( deviceArray );
deviceArray,
BLOCKS * BLOCKS * sizeof(int),
cudaMemcpyDeviceToHost );
for (int i=0; i<BLOCKS; i++)
{
printf( “Thread ID running: %d”, hostArray[0][i] );
for (int j=1; j<BLOCKS; j++)
{
printf( ” %d”, hostArray[j][i] );
}
printf( “\n” );
}
cudaFree( deviceArray );
这里也没什么要说的,除了我们现在也从设备拷贝BLOCKS*BLOCKS外,我们也要遍历每个块,打印出它们的内容。
最后不要忘了main函数的结束括号“}”。
现在,我们添加内核:
__global__ void generateArray( int *hostArray )
{
int ThreadIndex = blockIdx.x + blockIdx.y * gridDim.x;
hostArray[ThreadIndex] = ThreadIndex;
}
{
int ThreadIndex = blockIdx.x + blockIdx.y * gridDim.x;
hostArray[ThreadIndex] = ThreadIndex;
}
看起来和前面的完全一样,blockId.x用于获得在x维度上工作的块,blockId.y用于获得在y维度上工作的块,girdDim是我们网格维度的最大长度,.x表示x轴,.y表示y轴,和创建块数组指定的数字是一样的。
如果运行这个例子,你将看到类似下图所示的结果:
▲图 2 运行结果