之前我们说过,核函数除了没有返回值,还可以实现参数传递,我们来看这段代码
#include<cstdio>
#include<stdio.h>
__global__ void kernel(int *p)
{
*p = 20;
}
int main()
{
int *p;
int b;
cudaMalloc(&p,sizeof(int));
kernel<<<1,1>>>(p);
cudaMemcpy(&b,p,sizeof(int),cudaMemcpyDeviceToHost);
printf("%d\n",b);
return 0;
}
这里cudaMalloc函数是在设备上分配一块内存给p,也就是在gpu上分配一块显存给p。,值得注意的是cudaMalloc分配的指针可以传递给在在主机上执行的函数,也可以转递给设备上执行的函数,也可以在设备代码中使用指针进行内存(显存)读写操作,但是就是不能在主机内存中进行解引用读写操作。这里我们使用了cudaMemcpy函数来将显存的数据拷贝到cpu的内存中,这里我用b来接受p的值。当然,这个代码并不完美,我们还得释放显存中的地址p,可以使用cudaFree函数实现。完整代码
#include<cstdio>
#include<stdio.h>
__global__ void kernel(int *p)
{
*p = 20;
}
int main()
{
int *p;
int b;
cudaMalloc(&p,sizeof(int));
kernel<<<1,1>>>(p);
cudaMemcpy(&b,p,sizeof(int),cudaMemcpyDeviceToHost);
printf("%d\n",b);
cudaFree(p);
return 0;
}
值得注意的是,cudaMecpy函数执行后会自动调用cudaDeviceSynchronize()函数,这个函数的功能是让cpu等待直到所有的gpu程序执行完。对于一些追求性能的操作,cudaMecpy函数最好少用,多了容易导致gpu队列重新编排影响性能耗费时间。一般我们比较常用cudaMallocManaged函数来代替cudaMalloc函数,这个函数会同时在cpu和gpu上分配内存,可以实现共用数据,代码如下
#include<cstdio>
#include<stdio.h>
__global__ void kernel(int *p)
{
*p = 20;
}
int main()
{
int *p;
cudaMallocManaged(&p,sizeof(int));
kernel<<<1,1>>>(p);
printf("%d\n",*p);
cudaFree(p);
return 0;
}
现在直接可以在主机上进行读写操作解引用了。当然,你直接这样运行结果回是空的,之前可以打印出结果是因为后台调用了cudaDeviceSynchronize();来实现cpu等待gpu执行完才执行。因此为了显示结果我们可以加上来看,如图所示
点击此处登录后即可评论