最小限のコードであまり落ちてなかったので自分のメモ用としても

まずはCUDA版(nvccでコンパイルするやつ)
#include <stdio.h>
__global__ void VecAdd(float* A, float* B, float* C) {
	int id = 256 * blockIdx.x + threadIdx.x;
	C[id]=A[id]+B[id];
}

int main() {
	int N = 1024;
	// Allocate 3 arrays on GPU
	float *d_A, * d_B, * d_C;
	cudaMallocManaged(&d_A, N * sizeof(float));
	cudaMallocManaged(&d_B, N * sizeof(float));
	cudaMallocManaged(&d_C, N * sizeof(float));
	// CPU init
	for(int i=0;i<N;i++){
		d_A[i]=d_B[i]=1.0*i;
	}

	//gpu kernel
	VecAdd <<<N/256, 256>>> (d_A, d_B, d_C);
	cudaDeviceSynchronize();//wait

	for(int i = 0;i<N;i++){
		printf("%f ",d_C[i]);
	}
	
	cudaFree(d_A);
	cudaFree(d_B);
	cudaFree(d_C);
	return 0;
}

つぎにPyCUDA版(2020.1)
from pycuda.autoinit import context
import pycuda.driver as cuda
from pycuda.compiler import SourceModule
import numpy as np

mod = SourceModule("""
__global__ void doublify(float *a)
{
    a[threadIdx.x] *= 2;
}
""")
doublify = mod.get_function("doublify")

a = cuda.managed_empty(shape=12, dtype=np.float32, mem_flags=cuda.mem_attach_flags.GLOBAL)
a[:] = np.linspace(0, 11, len(a)) # Fill array on host
doublify(a, grid=(1,1), block=(len(a),1,1))
context.synchronize() # Wait for kernel completion before host access
print(a)
del a #これでガベージコレクションに削除してもらうことで、GPUメモリも解放される

ほんの19行でかけた aにCPUからもGPUからもアクセスできている。