前置命题，保序的命题：

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

测试两个命题：

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  代码的执行；