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