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

2023-12-13 22:19:47

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

文章来源:https://blog.csdn.net/eloudy/article/details/134903469
本文来自互联网用户投稿,该文观点仅代表作者本人,不代表本站立场。本站仅提供信息存储空间服务,不拥有所有权,不承担相关法律责任。