Edited NVCC definitions, working test of a kernel and/or write-read data.

CMakeLists.txt

@@ -10,7 +10,7 @@ include_directories(include)
 
 #link_directories(${CUDA_LIBRARY_DIRS})
 
-set(CUDA_NVCC_FLAGS ${CUDA_NVCC_FLAGS}-Wno-deprecated-gpu-targets)
+set(CUDA_NVCC_FLAGS ${CUDA_NVCC_FLAGS};-Wno-deprecated-gpu-targets;-gencode arch=compute_60,code=sm_60;-gencode arch=compute_35,code=sm_35)#;-rdc=true)
 
 cuda_add_executable(gpufirstcomm
   src/main.cu

launch.sh

@@ -6,6 +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"
-/home/mathiasb/sources/pciutils/ddrio -v -i /home/mathiasb/sources/gpuFirstComm/to_send -s 4096
+#/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
-
+ddrio -v -i /home/mathiasb/sources/gpuFirstComm/to_send -s 4096

src/.#main.cu

@@ -0,0 +1 @@
+root@ipepascal1.20424:1500564223

src/main.cu

@@ -29,12 +29,13 @@ int main()
     printf("\nInitialisation of the GPU\n");
     CUdevice GPU;
     CUdevprop GPUProp;
-    assert_cuda( cudaSetDevice(0) );
     assert_cu( cuDeviceGet(&GPU,0) );
     assert_cu( cuDeviceGetProperties(&GPUProp,GPU) );
-
+    char gpu_name[30] = {0};
+    assert_cu (cuDeviceGetName (gpu_name, 30, GPU));
+    printf("GPU: %s\n", gpu_name);
+    
     /* Check context */
-    assert_cu( cuCtxGetDevice(&GPU) );
     CUcontext cuCtx;
     assert_cu( cuCtxCreate(&cuCtx,CU_CTX_MAP_HOST|CU_CTX_SCHED_AUTO,GPU) );
     assert_cu( cuCtxSetCurrent(cuCtx) );
@@ -42,11 +43,11 @@ int main()
     
     /* Try some stuff... */
     printf("Using binary data to feed FPGA...\n");
-//    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));
+    char* data=(char*)calloc(4096,sizeof(*data));
+    memset(data,0xAB,4096);
+    init_to_send(data,4096);
+    /* uint64_t nb=0xA10000003; */
+    /* init_to_send(&nb,sizeof(nb)); */
     system("/home/mathiasb/sources/gpuFirstComm/launch.sh");
     
     
@@ -63,7 +64,7 @@ int main()
     assert_gdr( gdr_map(g,GPUMemHandlePage,&gpuPageVa,PAGE_SIZE) );
     gdr_info_t pageInfo;
     assert_gdr( gdr_get_info(g,GPUMemHandlePage,&pageInfo) );
-    printf("gpuPagePtr = %lx\nBus ptr = %lx\nVA = %lx\nSize = %lu\n",gpuPagePtr,pageInfo.bus_addr,pageInfo.va,pageInfo.mapped_size);
+    printf("gpuPagePtr = %llx\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");
@@ -76,11 +77,10 @@ int main()
     assert_gdr( gdr_map(g,GPUMemHandleDesc,&gpuDescVa,GPU_PAGE) );
     gdr_info_t descInfo;
     assert_gdr( gdr_get_info(g,GPUMemHandleDesc,&descInfo) );
-    printf("gpuDescPtr = %lx\nBus ptr = %lx\nVA = %lx\nSize = %lu\n",gpuDescPtr,descInfo.bus_addr,descInfo.va,descInfo.mapped_size);
-    
-    
+    printf("gpuDescPtr = %llx\nBus ptr = %lx\nVA = %lx\nSize = %lu\n",gpuDescPtr,descInfo.bus_addr,descInfo.va,descInfo.mapped_size);    
     
     printf("All set\n");
+
     
     /* PCI */
     printf("\nSetting up the PCI\n");
@@ -151,63 +151,66 @@ 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);
+    /* volatile uint64_t* trackAddr = (uint64_t*)((char*)pciMemPtrDesc+DESCRIPTOR_OFFSET+2*sizeof(uint32_t)); */
+    /* printf("Data were written at %p\n",trackAddr); */
+    /* printf("trackAddr = %lx\n",*trackAddr); */
+    *(uint32_t*)gpuDescVa=0;
+    printf("counter = %d\n",*(int*)gpuDescVa);
+
+    
     /* 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,1); //16);
+    WR32_sleep(REG_PERF_COUNTER,0);
+    WR32_sleep(REG_NUM_PACKETS_PER_DESCRIPTOR,2); //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,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, 0x232); //0x262);
-    WR32_sleep(REG_COUNTER,1);
+    WR32_sleep(REG_COUNTER,0x1);
     usleep(100000);
-
-
-    WR64_sleep(REG_DESCRIPTOR_ADDRESS,pageInfo.bus_addr);
+    printf("counter = %lx\n",*(uint64_t*)gpuDescVa);
+    WR32_sleep(REG_PERF_COUNTER,1);
     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( 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);
-
+    
+    WR64(REG_DESCRIPTOR_ADDRESS,pageInfo.bus_addr);
+    double lath = 4. *RD32 (0x14)/ 1000;
+    double lath20 = 4. * RD32 (0x20) / 1000;
+    printf("lath = %lf\nlath20 = %lf\n",lath,lath20);
+    usleep(50000);
+    lath = 4. *RD32 (0x14)/ 1000;
+    lath20 = 4. * RD32 (0x20) / 1000;
+    printf("lath = %lf\nlath20 = %lf\n",lath,lath20);
+    printf("counter = %lx\n",*(uint64_t*)gpuDescVa);
+
+    /* 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 = %p\npciMemPtrPage = %lx\npciMemPtrDesc = %lx\n",pciVa,(uint64_t)pciMemPtrPage,(uint64_t)pciMemPtrDesc); */
+    /* printf("gpuPagePtr = %llx\ngpuDescPtr = %llx\n",gpuPagePtr,gpuDescPtr); */
+
+    /* printf("Data were written at %p\n",trackAddr); */
+    /* printf("trackAddr = %lx\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);
+    for(int i=0;i<512;i++)
+    {
+    	printf("%hhx",getBack[i]);
+    }
+    printf("\n");
+    printf("counter = %lx\n",*(uint64_t*)gpuDescVa);
+    free(getBack);
     
     /* Close everything */
     printf("\nClosing the connections\n");