|
@@ -51,23 +51,27 @@ int main()
|
|
|
printf("Support unified addressing? %d\n",pi);
|
|
|
|
|
|
/* Allocate memory on the device, pin and map */
|
|
|
- printf("\nMemory mapping with the GPU\n");
|
|
|
- CUdeviceptr dptr;
|
|
|
- assert_cu( cuMemAlloc(&dptr,(size_t) GPUProp.sharedMemPerBlock) );
|
|
|
- gdr_mh_t GPUMemHandle;
|
|
|
- assert_gdr( gdr_pin_buffer(g,dptr,(size_t) GPUProp.sharedMemPerBlock,0,0,&GPUMemHandle) );
|
|
|
- void* va;
|
|
|
- assert_gdr( gdr_map(g,GPUMemHandle,&va,(size_t) GPUProp.sharedMemPerBlock) );
|
|
|
- /*CHECK THE OFFSET*/
|
|
|
- gdr_info_t GPUInfo;
|
|
|
- int offset;
|
|
|
- assert_gdr( gdr_get_info(g,GPUMemHandle,&GPUInfo) );
|
|
|
- offset = (GPUInfo.va > dptr) ? GPUInfo.va - dptr:dptr - GPUInfo.va;
|
|
|
- uint32_t *buf_ptr = (uint32_t *)((char *)va + offset);
|
|
|
-
|
|
|
+ printf("\nMemory mapping with the GPU for pages\n");
|
|
|
+ CUdeviceptr gpuPagePtr;
|
|
|
+ assert_cu( cuMemAlloc(&gpuPagePtr,3*PAGE_SIZE) );
|
|
|
+ gdr_mh_t GPUMemHandlePage;
|
|
|
+ assert_gdr( gdr_pin_buffer(g,gpuPagePtr,3*PAGE_SIZE,0,0,&GPUMemHandlePage) );
|
|
|
+ void* gpuPageVa;
|
|
|
+ assert_gdr( gdr_map(g,GPUMemHandlePage,&gpuPageVa,3*PAGE_SIZE) );
|
|
|
+
|
|
|
+ printf("Memory mapping with the GPU for descriptors\n");
|
|
|
+ CUdeviceptr gpuDescPtr;
|
|
|
+ assert_cu( cuMemAlloc(&gpuDescPtr,PAGE_SIZE) );
|
|
|
+ gdr_mh_t GPUMemHandleDesc;
|
|
|
+ assert_gdr( gdr_pin_buffer(g,gpuDescPtr,PAGE_SIZE,0,0,&GPUMemHandleDesc) );
|
|
|
+ void* gpuDescVa;
|
|
|
+ assert_gdr( gdr_map(g,GPUMemHandleDesc,&gpuDescVa,PAGE_SIZE) );
|
|
|
+
|
|
|
printf("All set\n");
|
|
|
- printf("va: 0x%lx\ndptr: 0x%llx\nGPUInfo.va: 0x%lx\noffset: 0x%x\nbuf_ptr: 0x%lx\n",
|
|
|
- (uint64_t) va,dptr,GPUInfo.va,offset,buf_ptr);
|
|
|
+ printf("gpuPageVa: 0x%lx\ngpuPagePtr: 0x%llx\n",
|
|
|
+ (uint64_t) gpuPageVa,gpuPagePtr);
|
|
|
+ printf("gpuDesc: 0x%lx\ngpuDescPtr: 0x%llx\n",
|
|
|
+ (uint64_t) gpuDescVa,gpuDescPtr);
|
|
|
|
|
|
/* At this point the GPU's mem is mapped to a CPU buffer to enable DMA */
|
|
|
|
|
@@ -89,70 +93,61 @@ int main()
|
|
|
}
|
|
|
|
|
|
CUdeviceptr dBAR;
|
|
|
- int getdBAR = 10;
|
|
|
assert_cu( cuMemHostRegister((void*)pciVa,128,CU_MEMHOSTREGISTER_IOMEMORY) );
|
|
|
assert_cu( cuMemHostGetDevicePointer(&dBAR,(void*)pciVa, 0) );
|
|
|
- printf("pciVa = %p\ndBar = 0x%llx\n",pciVa,dBAR);
|
|
|
- assert_cu( cuMemcpyDtoH(&getdBAR,dBAR,sizeof(int)) );
|
|
|
- printf("getdBAR = %d\n",getdBAR);
|
|
|
-
|
|
|
- pcilib_kmem_handle_t* pciHandle;
|
|
|
- pciHandle = pcilib_alloc_kernel_memory(pciCtx,PCILIB_KMEM_TYPE_CONSISTENT,1,64,4096,PCILIB_KMEM_USE_STANDARD,PCILIB_KMEM_FLAG_REUSE);
|
|
|
- if( pciHandle == NULL )
|
|
|
+
|
|
|
+ /* Config PCI for Pages*/
|
|
|
+ pcilib_kmem_handle_t* pciHandlePage;
|
|
|
+ pciHandlePage = pcilib_alloc_kernel_memory(pciCtx,PCILIB_KMEM_TYPE_CONSISTENT,1,64,4096,PCILIB_KMEM_USE_STANDARD,PCILIB_KMEM_FLAG_REUSE);
|
|
|
+ if( pciHandlePage == NULL )
|
|
|
{
|
|
|
printf("Cannot allocate PCI kernel memory\n");
|
|
|
exit( EXIT_FAILURE );
|
|
|
}
|
|
|
- volatile void* pciMemPtr;
|
|
|
- uintptr_t pciBus;
|
|
|
- pciMemPtr = pcilib_kmem_get_ua(pciCtx,pciHandle);
|
|
|
- if( pciMemPtr == NULL )
|
|
|
+ volatile void* pciMemPtrPage;
|
|
|
+ uintptr_t pciBusPage;
|
|
|
+ pciMemPtrPage = pcilib_kmem_get_ua(pciCtx,pciHandlePage);
|
|
|
+ if( pciMemPtrPage == NULL )
|
|
|
{
|
|
|
printf("Cannot get PCI pointer to kernel memory\n");
|
|
|
exit( EXIT_FAILURE );
|
|
|
}
|
|
|
- pciBus = pcilib_kmem_get_ba(pciCtx,pciHandle);
|
|
|
- if( pciBus == 0 )
|
|
|
+ pciBusPage = pcilib_kmem_get_ba(pciCtx,pciHandlePage);
|
|
|
+ if( pciBusPage == 0 )
|
|
|
{
|
|
|
printf("Cannot get PCI Bus address on kernel memory\n");
|
|
|
exit( EXIT_FAILURE );
|
|
|
}
|
|
|
|
|
|
- printf("pciMemPtr = %p\npciBus = %p\n",pciMemPtr,pciBus);
|
|
|
+ printf("pciMemPtrPage = %p\npciBusPage = %p\n",pciMemPtrPage,pciBusPage);
|
|
|
|
|
|
- const pcilib_bar_info_t* bar_info;
|
|
|
- bar_info = pcilib_get_bar_info(pciCtx,0);
|
|
|
- if( bar_info == NULL )
|
|
|
+ /* Config PCI for Desc */
|
|
|
+ pcilib_kmem_handle_t* pciHandleDesc;
|
|
|
+ pciHandleDesc = pcilib_alloc_kernel_memory(pciCtx,PCILIB_KMEM_TYPE_CONSISTENT,1,64,4096,PCILIB_KMEM_USE_STANDARD,PCILIB_KMEM_FLAG_REUSE);
|
|
|
+ if( pciHandleDesc == NULL )
|
|
|
+ {
|
|
|
+ printf("Cannot allocate PCI kernel memory\n");
|
|
|
+ exit( EXIT_FAILURE );
|
|
|
+ }
|
|
|
+ volatile void* pciMemPtrDesc;
|
|
|
+ uintptr_t pciBusDesc;
|
|
|
+ pciMemPtrDesc = pcilib_kmem_get_ua(pciCtx,pciHandleDesc);
|
|
|
+ if( pciMemPtrDesc == NULL )
|
|
|
+ {
|
|
|
+ printf("Cannot get PCI pointer to kernel memory\n");
|
|
|
+ exit( EXIT_FAILURE );
|
|
|
+ }
|
|
|
+ pciBusDesc = pcilib_kmem_get_ba(pciCtx,pciHandleDesc);
|
|
|
+ if( pciBusDesc == 0 )
|
|
|
{
|
|
|
- printf("Cannot get BAR info\n");
|
|
|
+ printf("Cannot get PCI Bus address on kernel memory\n");
|
|
|
exit( EXIT_FAILURE );
|
|
|
}
|
|
|
- else
|
|
|
- printf("bar = %d\nsize = %lu\nphys_addr = %p\nvirt_addr = %p\n",bar_info->bar,bar_info->size,bar_info->phys_addr,bar_info->virt_addr);
|
|
|
|
|
|
+ printf("pciMemPtrDesc = %p\npciBusDesc = %p\n",pciMemPtrDesc,pciBusDesc);
|
|
|
+
|
|
|
+
|
|
|
/* Try some stuff... */
|
|
|
- printf("\nRunning some tests\n");
|
|
|
- *(int*) pciMemPtr = 65;
|
|
|
- printf("Try to read pciMem in another way : %d %d\n",*(int*)pciMemPtr,*(int*)dBAR);
|
|
|
- assert_cu( cuMemcpyHtoD(dptr,(void*)pciMemPtr,sizeof(int)) );
|
|
|
- assert_cu( cuMemcpyHtoD(dBAR,(void*)pciMemPtr,sizeof(int)) );
|
|
|
- assert_cu( cuMemcpyDtoH(&getdBAR,dBAR,sizeof(int)) );
|
|
|
- printf("getdBAR = %d\n",getdBAR);
|
|
|
- printf("Try to read pciMem in another way : %d %d\n",*(int*)pciMemPtr,*(int*)dBAR);
|
|
|
- printf("Try to read pciMem in another way : %d\n",*(int*)va);
|
|
|
- printf("A small computation\n");
|
|
|
- add_three_global<<< 1,1 >>>(dptr);
|
|
|
- assert_cuda( cudaDeviceSynchronize() );
|
|
|
- int foo;
|
|
|
- assert_gdr( gdr_copy_from_bar(&foo,va,sizeof(foo)) );
|
|
|
- printf("foo = %d\nva = %d\n",foo,*(int*) va);
|
|
|
- printf("Now the other way around\n");
|
|
|
- foo = 100;
|
|
|
- assert_gdr( gdr_copy_to_bar(va,&foo,sizeof(foo)) );
|
|
|
- add_three_global<<< 1,1 >>>(dptr);
|
|
|
- assert_cuda( cudaDeviceSynchronize() );
|
|
|
- assert_cu( cuMemcpyDtoH((void*)pciMemPtr,dptr,sizeof(int)) );
|
|
|
- printf("pciMem = %d\n",*(int*)pciMemPtr);
|
|
|
|
|
|
/* FPGA */
|
|
|
printf("\nWorking on the FPGA\n");
|
|
@@ -167,10 +162,13 @@ int main()
|
|
|
printf("\nClosing the connections\n");
|
|
|
assert_cu( cuMemHostUnregister((void*) pciVa) );
|
|
|
pcilib_close(pciCtx);
|
|
|
- assert_gdr( gdr_unmap(g,GPUMemHandle,va,(size_t) GPUProp.sharedMemPerBlock) );
|
|
|
- assert_gdr( gdr_unpin_buffer(g,GPUMemHandle) );
|
|
|
+ assert_gdr( gdr_unmap(g,GPUMemHandlePage,gpuPageVa,3*PAGE_SIZE) );
|
|
|
+ assert_gdr( gdr_unpin_buffer(g,GPUMemHandlePage) );
|
|
|
+ assert_gdr( gdr_unmap(g,GPUMemHandleDesc,gpuDescVa,3*PAGE_SIZE) );
|
|
|
+ assert_gdr( gdr_unpin_buffer(g,GPUMemHandleDesc) );
|
|
|
assert_gdr( gdr_close(g) );
|
|
|
- assert_cu( cuMemFree(dptr) );
|
|
|
+ assert_cu( cuMemFree(gpuPagePtr) );
|
|
|
+ assert_cu( cuMemFree(gpuDescPtr) );
|
|
|
|
|
|
printf("All Cleared\n");
|
|
|
|