DMA FPGA to GPU, kernel and read back from RAM.

CMakeLists.txt

@@ -11,15 +11,13 @@ include_directories(include)
 #link_directories(${CUDA_LIBRARY_DIRS})
 
 set(CUDA_NVCC_FLAGS ${CUDA_NVCC_FLAGS}-Wno-deprecated-gpu-targets)
-set(CMAKE_C_COMPILER "/usr/bin/clang")
-set(CMAKE_C_FLAGS "-msse -msse4.1 -mavx")
 
 cuda_add_executable(gpufirstcomm
   src/main.cu
   src/common.cu
-  src/memcpy_avx.c
-  src/memcpy_sse41.c
-  src/memcpy_sse.c
-  src/kernels.cu)
-
+  #src/memcpy_avx.c
+ # src/memcpy_sse41.c
+ # src/memcpy_sse.c
+  src/kernels.cu
+)
 target_link_libraries(gpufirstcomm cuda pcilib gdrapi)

build/reload_cmake.sh

@@ -1,6 +0,0 @@
-# Used to reload the CMakeLists.txt file with clang as a compiler instead of gcc
-
-#!/bin/bash
-
-rm -r *
-CC=clang CXX=clang cmake ..

include/common.h

@@ -13,6 +13,6 @@
 void assert_cuda(cudaError_t err_id); /* for runtime api*/
 void assert_cu(CUresult res_id); /* for driver api */
 void assert_gdr(int gdr_id);
-void init_to_send(const void* dataPtr);
+void init_to_send(const void* dataPtr,size_t nb_bytes);
 
 #endif

kill_prog

@@ -0,0 +1,3 @@
+#!/bin/bash
+
+kill -SIGKILL `pidof gpufirstcomm`

launch.sh

@@ -6,5 +6,6 @@ echo "Resetting the FPGA"
 /usr/local/bin/pci -w 9040 0
 sleep 1
 echo "Putting the data on the FPGA. WARNING, hardcoded size"
-ddrio -v -i /home/mathiasb/sources/gpuFirstComm/to_send -s 0
+/home/mathiasb/sources/pciutils/ddrio -v -i /home/mathiasb/sources/gpuFirstComm/to_send -s 4096
+#/home/mathiasb/sources/pciutils/ddrio -v -o /home/mathiasb/sources/gpuFirstComm/written -s 4096
 

src/common.cu

@@ -31,7 +31,7 @@ void assert_gdr(int gdr_id)
     }
 }
 
-void init_to_send(const void* dataPtr)
+void init_to_send(const void* dataPtr, size_t nb_bytes)
 {
     FILE* filePtr = fopen("/home/mathiasb/sources/gpuFirstComm/to_send","wb");
     if( filePtr == NULL )
@@ -40,7 +40,7 @@ void init_to_send(const void* dataPtr)
 	exit( EXIT_FAILURE );
     }
     int errCheck;
-    errCheck = fwrite(dataPtr,4096,1,filePtr);
+    errCheck = fwrite(dataPtr,nb_bytes,1,filePtr);
     if( errCheck == 0 )
     {
 	printf("Could not write the items. Exiting...\n");

src/main.cu

@@ -42,11 +42,13 @@ int main()
     
     /* Try some stuff... */
     printf("Using binary data to feed FPGA...\n");
-    char* data=(char*)calloc(4096,sizeof(*data));
-    memset(data,0xCC,4096);
-    init_to_send(data);
+//    char* data=(char*)calloc(4096,sizeof(*data));
+//    memset(data,0x3C,4096);
+//    init_to_send(data,4096);
+    uint64_t nb=0xA10000001;
+    init_to_send(&nb,sizeof(nb));
     system("/home/mathiasb/sources/gpuFirstComm/launch.sh");
-
+    
     
     /* Allocate memory on the device, pin and map */
     uint8_t flagValueToSet = 1;
@@ -61,9 +63,9 @@ int main()
     assert_gdr( gdr_map(g,GPUMemHandlePage,&gpuPageVa,PAGE_SIZE) );
     gdr_info_t pageInfo;
     assert_gdr( gdr_get_info(g,GPUMemHandlePage,&pageInfo) );
-    printf("Bus ptr = %lx\nVA = %lx\nSize = %lu\n",pageInfo.bus_addr,pageInfo.va,pageInfo.mapped_size);
-    
+    printf("gpuPagePtr = %lx\nBus ptr = %lx\nVA = %lx\nSize = %lu\n",gpuPagePtr,pageInfo.bus_addr,pageInfo.va,pageInfo.mapped_size);
 
+    
     printf("Memory mapping with the GPU for descriptors\n");
     CUdeviceptr gpuDescPtr;
     assert_cu( cuMemAlloc(&gpuDescPtr,GPU_PAGE) );       
@@ -74,7 +76,9 @@ int main()
     assert_gdr( gdr_map(g,GPUMemHandleDesc,&gpuDescVa,GPU_PAGE) );
     gdr_info_t descInfo;
     assert_gdr( gdr_get_info(g,GPUMemHandleDesc,&descInfo) );
-    printf("Bus ptr = %lx\nVA = %lx\nSize = %lu\n",descInfo.bus_addr,descInfo.va,descInfo.mapped_size);
+    printf("gpuDescPtr = %lx\nBus ptr = %lx\nVA = %lx\nSize = %lu\n",gpuDescPtr,descInfo.bus_addr,descInfo.va,descInfo.mapped_size);
+    
+    
     
     printf("All set\n");
     
@@ -147,44 +151,62 @@ int main()
     }
     printf("pciMemPtrPage = %lx\npciMemPtrDesc = %lx\n",(uint64_t)pciMemPtrPage,(uint64_t)pciMemPtrDesc);
 
-    
+    volatile uint64_t* trackAddr = (uint64_t*)((char*)pciMemPtrDesc+DESCRIPTOR_OFFSET+2*sizeof(uint32_t));
+    printf("Data were written at %lx\n",trackAddr);
+    printf("trackAddr = %llx\n",*trackAddr);
     /* FPGA */
     printf("\nWorking on the FPGA\n");
     WR32(REG_RESET_DMA, 1);
     usleep(100000);
     WR32(REG_RESET_DMA, 0);
     usleep(100000);
-    WR32_sleep(REG_NUM_PACKETS_PER_DESCRIPTOR,16);
-    WR32_sleep(REG_PACKET_LENGTH,64);
+    WR32_sleep(REG_NUM_PACKETS_PER_DESCRIPTOR,1); //16);
+    WR32_sleep(REG_PACKET_LENGTH,0x80000 | 64); // added flag
     WR32_sleep(REG_TIMER_THRESHOLD, 0x1);
     WR32_sleep(REG_UPDATE_THRESHOLD, 0x1);
-/*    WR64_sleep(REG_UPDATE_COUNTER,pciBusDesc);*/
-    WR64_sleep(REG_UPDATE_ADDRESS,pciBusPage+DESCRIPTOR_OFFSET);
+    WR64_sleep(REG_UPDATE_COUNTER,descInfo.bus_addr);
+    WR64_sleep(REG_UPDATE_ADDRESS,pciBusDesc+DESCRIPTOR_OFFSET);
     WR32_sleep(REG_CONTROL,CONTROL_ENABLE_READ|CONTROL_SOURCE_RX_FIFO);
     /* WR64_sleep(REG_DESCRIPTOR_ADDRESS,descInfo.bus_addr); */
     WR32_sleep(REG_DMA,1);
-    WR32_sleep(REG_INTERCONNECT, 0x262);
+    WR32_sleep(REG_INTERCONNECT, 0x232); //0x262);
     WR32_sleep(REG_COUNTER,1);
     usleep(100000);
 
-    WR64_sleep(REG_DESCRIPTOR_ADDRESS,pciBusDesc);
-    /* usleep(100000); */
+
+    WR64_sleep(REG_DESCRIPTOR_ADDRESS,pageInfo.bus_addr);
+    usleep(100000);
     printf("pageInfo.bus_addr = %lx\ndescInfo.bus_addr = %lx\npciBusPage = %lx\npciBusDesc = %lx\n",pageInfo.bus_addr,descInfo.bus_addr,pciBusPage,pciBusDesc);
     printf("gpuDescVa = %lx\ngpuPageVa = %lx\n",(uint64_t)gpuDescVa,(uint64_t)gpuPageVa);
     printf("pciVa = %x\npciMemPtrPage = %lx\npciMemPtrDesc = %lx\n",pciVa,(uint64_t)pciMemPtrPage,(uint64_t)pciMemPtrDesc);
+    printf("gpuPagePtr = %lx\ngpuDescPtr = %lx\n",gpuPagePtr,gpuDescPtr);
 
-    /* assert_cu( cuMemcpyHtoD(gpuDescPtr,(const void*)data,4096) ); */
-    /* uint64_t* trackAddr = (uint64_t*); */
-    /* printf("Data were written at %lx\n",*trackAddr); */
-    char* getBack=(char*)calloc(4096,sizeof(*getBack));
-    memcpy(getBack,(const void*)pciMemPtrDesc,4096);
-    int i;
-    for(i=0;i<4096;i++)
-    {
-    	printf("%hhx",getBack[i]);
-    }
-    printf("\n");
-    free(getBack);
+//     assert_cu( cuMemcpyDtoH(gpuPageVa,gpuPagePtr,4096) );
+
+    printf("Data were written at %lx\n",trackAddr);
+    printf("trackAddr = %llx\n",*trackAddr);
+    
+    add_three_global<<< 1,1 >>>(gpuPagePtr);
+    assert_cu( cuCtxSynchronize() );
+    printf("Received : %lx\n",*(uint64_t*)gpuPageVa);
+    
+//    char* getBack=(char*)calloc(4096,sizeof(*getBack));
+//    memcpy(getBack,(const void*)gpuPageVa,4096);
+//    int i;
+//    for(i=254;i<4096;i++)
+//    {
+//    	printf("%hhx",getBack[i]);
+//    }
+//    printf("\n");
+//    printf("getBack info from Desc : %p\n%p\n",getBack,*(char**)getBack);
+//    memcpy(getBack,(const void*)(pciMemPtrDesc),4096);
+//    for(i=0;i<4096;i++)
+//    {
+//    	printf("%hhx",getBack[i]);
+//    }
+//    printf("\n");
+//    printf("getBack info from Pages%p\n%p\n",getBack,*(char**)getBack);
+//    free(getBack);
 
     
     /* Close everything */