CUDAC知识回顾1-3章课后题

1.1
是一个单纯的块在图像上的划分，划分结果如下图

1.2
hello.cu

#include<stdio.h>
__global__ void hellofromGPU(void)
{
printf("hello world from CPU\n");
}
int main()
{
printf("hello world from gpu\n");
hellofromGPU<<<1,10>>>();
cudaDeviceReset();
return 0;
}

cudaDeviceReset()函数是用来显示的释放和清空当前进程中于当前设备有关的所有资源
当重置函数移除后 编译运行则只输出

hello world from cpu

当printf在gpu上被调用，cudaDeviceReset()函数使这些来自gpu的输出发送到主机，然后在控制台输出。没有调用cudaDeviceReset()函数就不能保证这些可以被显示。

1.3
输出效果和hello.cu一样
cudaDeviceSynchronize()是另一个可以用来使gpu打印在用户可见控制台的函数

1.4
（个人感觉 现在的nvcc编译器可以直接编译 不用声明架构）
移除之后无法生成执行文件。当没有指定体系结构 标志时，nvcc将默认使用计算能力1.0 支持从gpu调用printf

1.5
nvcc支持：.cu .cup .c .cc .cxx .cpp .GPU .ptx .o .obj .a .lib .res .so
它可以作为通用编译器来构建更多的cuda源代码

1.6
使用条件来确保只有线程5执行打印，并将线程id添加到打印语句中

#include<stdio.h>
__global__ void hellofromGPU(void)
{
	if(threadIdx.x==5)
		{
		printf("hello world from GPU thread %d\n",threadIdx.x);
		}
}
int main()
{
	printf("hello world from cpu\n");
	
	hellofromGPU<<<1,10>>>();
	cudaDeviceReset();
	return 0;
}

2.1
sumarraysongpu-timer.cu
1024（用vs2017生成）

1023

很明显的看到不同块之间的性能差异 1024明显比1023要快
原因：
由于线程被安排在32组中，所以运行1023个线程的块会导致两个次优的应用特性。
每个线程块中的最后一个warp将有一个禁用的线程，不执行任何工作，因为1023不能被均匀的分割成warp
每个块的线程比较少，device需要更多的块来完全处理输入。
运行更多的线程块会导致更长的执行时间，因为只有固定数量的块可以并发运行。

2.2

__global__ void sumArrayOnGPU(float *A,*B,*C,const int n)
{
	int nthreads = gridDim.x*blockDim.x;
	int i = blockIdx.x*blockDim.x+threadIdx.x;
	int j=i+nthreads;
	if (i<n)
		C[i]=A[i]+B[i];
	if(j<n)
		C[i]=A[i]+B[i];
	}

改变内核的行为也需要改变内核的执行配置，需要一半的线程来完整的输入

dim3 grid ((nElm/2)+block.x-1)/block.x);

结果：

比1023性能要好
但没有达到1024的性能
个人感觉是有if判断分支 占用了时间和线程

2.3
整数矩阵的加法运算
sumMatrixOnGPU-2D-grid-2D-block.cu（1-3章code中会发布出来）
修改dimx和dimy求解最佳性能
（1024，2）的结果如上
通过对比host和device的运行时间 device并行要快一百倍
最佳性能应该是（128，2）
由此可得 最佳性能和块的分布无关

2.4
sumMatrixOnGPU-2D-grid-1D-block.cu

每个线程处理两个元素（2.2也是）

代码后面写在一个专门的code文档中
其中最佳性能应该是block.x=256

2.5
checkDeviceInfor.cu
直接运行此文档

我电脑的性能

3.1
展开循环，数据块或者线程束可以减少频繁的分支。此外展开可以使编译器发现的独立内存操作的数量增加。因此可以发出更多的并发读写操作，内存带宽利用率增加。

3.2
reduceUnrolling16.cu（代码写在code中）
运行结果如下

16相对于8有了近一倍的性能改进。
16的吞吐量相对于8来说会更好。

3.3（仅供参考）
reduceUnrolling8.cu（写在code中）
运行结果如下：
和原有cuda作比较
但，此改进取决于现在使用的cuda版本，不同的版本性能可能会被倒置。

3.4
由于我用的vs无法识别__syncthreads()函数 所以就不演示此题。

声明vmem为volatile会同时影响使用该指针执行的负载和存储。_syncthreads（）只需要确保块中的所有线程都知道同一块中的其他线程全局内存的更改。

3.5
（放在code里面）

3.6
运行结果：

根据架构的不同，性能可能会有差异。在旧的gpu体系结构中，浮点算术逻辑单元的争用可能导致性能损失。但现在很多gpu上，性能没有差异，因为int和float都有相同的字数字节。

3.7
当父级已经同步完成其孩子子级时，会继续通过其同步。

3.8/3.9
vs中无法进行父子级的嵌套循环。
代码放在code中