测试命题 cuda kernel 和 cudaMemcpy 是异步执行

0，前置命题

保序的命题：

同一个任意的stream中的gpu操作（memcpy和kernel），在gpu内部都是严格保序的，即，前一个gpu任务结束后才会执行下一个任务。

cudaMemcpy只能用于默认流，也注定是与cpu同步执行的；
cudaMemcpyAsync只能用于显式流，也注定是与cpu异步的；
故，也可以总结为，默认流的memcpy是与cpu代码同步的，显式流的memcpy是与cpu异步的；但在gpu内，各流自己内部是保序的。

然后才加入阻塞显式流，与非阻塞显式流；
阻塞显式流，表现形式是相对于默认流的低优先级流，其中的gpu操作会被默认流的gpu操作阻塞，让出计算资源给默认流优先使用；非阻塞显式流则是与默认流通优先级的流，谁也不必然阻塞谁；

cudaStreamCreate创建的流，是阻塞流；

cudaStreamCreateWithFlag可以用来创建非阻塞流；

测试两个命题：

1，cuda kernel 是异步执行，即，主机程序在调用kernel<<<,,,>>>(); 后，控制权会马上返回cpu程序，进而主机程序继续向下执行，对吗？ 对，会马上执行接下来的cpu代码。

2，如果启动了kernel<<<>>>()后，接下来马上执行到了 cudaMemcpy ，那么会等到到 gpu 端kernel执行完毕才执行真正的 gpu copy动作，还是会在不顾kernel<<<>>>是否执行完毕，不顾数据是否正确，就开始执行真正的硬件 gpu copy 了，从而导致数据错误呢？不会，同一个stream中肯定是保序的，是不会导致数据错误的。

3. cudaMemcpy内含隐式gpu和cpu的同步，异步的话要使用 cudaMemcpyAsync，同时要求所使用的内存是锁叶内存；
 

1, 测试命题1的主体代码

这个kernel在n=1024*1024*1024 这么大时，需要执行约 7.5S

需要测量的问题：

launch kernel 函数后，是否会立即进入cpu代码的执行中呢？

cudaMemcpy被调用后，是否会立即进入cpu代码的执行中呢？

#include <cuda_runtime.h>
#include <stdio.h>

__global__ void ke(float *A, int n)
{
	for(int i=0; i<n; i++)
		A[i] = n;

}


int main()
{
	float *Ad = nullptr;
	float *Ah = nullptr;

	int n = 1024*1024*4;
	

	cudaMalloc((void**)&Ad, n*sizeof(float));
	Ah = (float*)malloc(n*sizeof(float));

	ke<<<1,1>>>(Ad, n);

	for(int i=0; i<10; i++)
		printf("hello____\n");


	cudaMemcpy(Ah, Ad, n*sizeof(float), cudaMemcpyDeviceToHost);

	printf("Ah[111]=%f\n", Ah[111]);


	return 0;
}

2,  测试命题1的测量方法

  #include <stdio.h>
 #include <sys/time.h>


#include <cuda_runtime.h>
#include <stdio.h>

__global__ void ke(float *A, long  int n)
{

        for(long int i=0; i<n; i++)
                A[0] += 0.0001;

}


int main()
{
        struct timeval start;
        struct timeval end;
        unsigned long timer;

        float *Ad = nullptr;
        float *Ah = nullptr;

        long int n = 1024*1024*4;

        cudaMalloc((void**)&Ad, n*sizeof(float));
        Ah = (float*)malloc(n*sizeof(float));
gettimeofday(&start, NULL);
        ke<<<1,1>>>(Ad, n);
//      cudaDeviceSynchronize();
gettimeofday(&end, NULL);
timer = 1000000*(end.tv_sec - start.tv_sec) + end.tv_usec - start.tv_usec;
printf("time1 = %d us\n", timer);

        for(int i=0; i<10; i++)
                printf("hello____\n");

gettimeofday(&end, NULL);
timer = 1000000*(end.tv_sec - start.tv_sec) + end.tv_usec - start.tv_usec;
printf("time2 = %d us\n", timer);

        cudaMemcpy(Ah, Ad, sizeof(float), cudaMemcpyDeviceToHost);

gettimeofday(&end, NULL);
timer = 1000000*(end.tv_sec - start.tv_sec) + end.tv_usec -start.tv_usec;
printf("time3 = %d us\n", timer);

        printf("Ah[0]=%f\n", Ah[0]);

    return 0;
}

运行效果：

下面这张是没有注释掉 cudaDeviceSynchronize(); 的打印：

下面是注释掉了 cudaDeviceSynchronize(); 函数的打印：

这意味着，在调用玩ke<<<>>>() 之后，立即进入接下来的  cpu  代码的执行；