最新帖子

页: [1] 2 3 ... 10
1
CUDA / 调用双GPU如何做到GPU0与GPU1之间是并行计算的?
« 最后发表 作者 100860xyz 一月 22, 2021, 05:58:18 pm »
我尝试将单卡的程序改成双卡,但是现在GPU0与GPU1是串行的,即现执行完GPU0上的任务再执行GPU1上的任务,下面附上相关的主机代码,请问这段程序需要怎么该呀?
   Ej_bi=0.
   Em_bi=0.
   Ej_bi_d=0.
   Em_bi_d=0.       
       
   istat=cudaGetDevice(0)    !第1个设备

   tBlock=dim3(4,16,1)
   grid=dim3(ceiling(real(Ntri)/tBlock%x),ceiling(real((n_scan+1)/2)/tBlock%y),1)
   call cuda_transmit_duoceng<<<grid,tBlock>>>(Ntri,(n_scan+1)/2,0)
       
   print*,"计算RCS"
      
    Ej_bi=Ej_bi_d
    Em_bi=Em_bi_d
   
   count=0   
    open(unit=10, fiLe="output\GPU_triple.csv") 
    do k = 1, (n_scan+1)/2
        temp2=0 
        do i = 1, Ntri
            E_bi(i,k,:)=Ej_bi(i,k,:)+Em_bi(i,k,:)
            temp2=temp2+E_bi(i,k,1)+E_bi(i,k,2)+E_bi(i,k,3)
        enddo
       
            print*, k,"flag"
            print*, abs(temp2),"E_bi_abs"
            print*, atan2(aimag(temp2),real(temp2)),"E_bi_angle"
            write(10,*) k,abs(temp2),atan2(aimag(temp2),real(temp2))
    enddo
    !close(10)
   Ej_bi=0.
   Em_bi=0.
   Ej_bi_d=0.
   Em_bi_d=0.
   istat=cudaGetDevice(1)      !第2个设备

   tBlock=dim3(4,16,1)
   grid=dim3(ceiling(real(Ntri)/tBlock%x),ceiling(real((n_scan-1)/2)/tBlock%y),1)
   call cuda_transmit_duoceng<<<grid,tBlock>>>(Ntri,(n_scan-1)/2,1)
       
   print*,"计算RCS"
      
    Ej_bi=Ej_bi_d
    Em_bi=Em_bi_d
     
    open(unit=10, fiLe="output\GPU_triple.csv") 
    do k = 1, (n_scan-1)/2
        temp2=0 
        do i = 1, Ntri
            E_bi(i,k,:)=Ej_bi(i,k,:)+Em_bi(i,k,:)
            temp2=temp2+E_bi(i,k,1)+E_bi(i,k,2)+E_bi(i,k,3)
        enddo
       
            print*, k+(n_scan+1)/2,"flag"
            print*, abs(temp2),"E_bi_abs"
            print*, atan2(aimag(temp2),real(temp2)),"E_bi_angle"
            write(10,*) k+(n_scan+1)/2,abs(temp2),atan2(aimag(temp2),real(temp2))
    enddo
   print*,count,"count"
   close(10)
   end
2
此外, 类似stdafx.h这种预编译头这种VC特有的东西, 尽量不要让nvcc看到(我假设你的.cuh扩展名是给nvcc看的).

如果上面说过的, 你总是可以通过适当的抽象逻辑, 让nvcc不看到这个的. 或者干脆不要使用任何stdafx.h这种VC的特有的预编译头(可以关闭不使用的, 不影响你的MFC项目, 当然, 可能会降低一点点编译速度). 正常情况下, nvcc也不应该看到他.
嗯嗯,好的我知道了,十分感谢!另外我最近尝试在win7上安装vs2015、cuda9,然后显卡用2080Ti,发现在运行模板工程addwithcuda时,cudasetdevice一直报错,错误提示cudaErrorunkonwn,在安装CUDA9时也提示了this graphic driver could not find compatible graphic hardware,然后我又试了试cuda10,因为cuda11好像没有win7的了,cuda10倒是没有提示找不到compatible graphic hardware,但是创建的模板工程addwithcuda依然会返回cudaErrorunkonwn的错误,win7是真的不能进行2080Ti的CUDA开发吗,是官方本身就不支持吗?真的没有什么办法吗?
3
此外, 类似stdafx.h这种预编译头这种VC特有的东西, 尽量不要让nvcc看到(我假设你的.cuh扩展名是给nvcc看的).

如果上面说过的, 你总是可以通过适当的抽象逻辑, 让nvcc不看到这个的. 或者干脆不要使用任何stdafx.h这种VC的特有的预编译头(可以关闭不使用的, 不影响你的MFC项目, 当然, 可能会降低一点点编译速度). 正常情况下, nvcc也不应该看到他.
4
你好, ATL和MFC我都不懂.

如果可能, 请尽量将GPU上的cuda kernel导出(例如通过一个C wrapper函数), 然后在你的项目的其他部分调用.

这样尽量维持你原始的CPU上的逻辑不动, 从而规避可能挂在ATL的.h中的某些(代码)行的情况.

很多时候这个常见(例如某些时候手工我们的用户报告过的OpenCV和CUDA配合, 挂在了N(值很大)层模板之类的代码行里头), 尚未知道原因, 解决方便一般都是反向操作(即上面说的), 尽量让NVCC编译较少数量的CUDA代码, 然后从其他部分调用.

欢迎试试这样看看能否解决.
5
更新:新建了一个MFC工程,添加了两个cu文件和两个cuh文件,两个cuh文件里面只要都包含stdafx.h或者包含atlstf.h等类似的头文件都会出现编译通过但是运行出错的问题,出错都是expression false&&"Too many categories defined";这样的错误,但是如果只有一个cuh文件包含,另一个cuh文件不包含就不会出现任何问题,所以推测是NVCC编译器和stdafx.h的问题
6
感觉像是预编译头stdafx.h的问题导致,但是还没有最终确定
7
后面我自己试着新建了一个MFC的工程,进行类似的包含操作,发现头文件中某些数据类型如CString也是可以正常运行的,但是当我把cuh文件的属性更改为CUDA c/C++时就会出现相同的情况,编译不报错,运行出错而且一模一样的错误,这时我去查看了下我的cuh文件属性是正常的c/c++标头,不是CUDA c/C++,这就很奇怪了,属性是对的运行起来还会报错
8
我的工程是VS2015,CUDA11,集运MFC的上位机程序,在这个程序中加入了CUDA程序,现在一共有三个cu和cuh文件分别是CHK.cu、PDW.cu、TFQ.cu、CHK.cuh、PDW.cuh、TFQ.cuh,还有两个头文件,分别是structCPU.h和structGPU.h,现在有一个问题,就是在PDW.cuh包含了structCPU.h和structGPU.h,另外两个只包含了structGPU.h,这种情况下编译可以通过,也可以运行,但是当CHK.cuh或者TFQ.cuh包含了structCPU.h时编译可以通过,但是运行会报错,如下所示:
D:\vs\VC\atlmfc\include\atltrace.h
line:337
expression false&&"Too many categories defined";
请问这是什么原因呢,一开始我以为是NVCC不支持structCPU.h[/size][size=78%]头文件中某些数据类型如CString,但是PDW.cuh也包含了[/size][/size]structCPU.h并且没有报错,而且我对比了这三个文件的属性,都是一样的设置,所以就很疑惑了,
[/size]希望[名词6]能分析下



9
CUDA / 如何使用cmake编译cuda的动态并行程序?
« 最后发表 作者 float风 一月 11, 2021, 11:14:28 am »
求助一个问题,我正在研究cuda的排序算法,需要对一组数据排序,然后返回排序后的索引。目前正在根据cuda案例cdpAdvancedQuicksort进行改编。
但是问题来了,我项目是采用cmake构建的,编译cdpAdvancedQuicksort时总会报错。查原因可能是动态并行的问题。
(是否有方便调用的API能实现这个功能呢,类似于numpy的argsort)

CMakeLists.txt如下:

cmake_minimum_required(VERSION 3.5)
project(quick_sort)

set(CMAKE_CXX_STANDARD 11)

find_package(CUDA)
SET(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -std=c++11")
SET(CUDA_NVCC_FLAGS ${CUDA_NVCC_FLAGS};-rdc=true -gencode arch=compute_75,code=sm_75;)

include_directories(/usr/local/cuda/include)  # .h

link_directories(/usr/local/cuda/lib64)
include_directories(/usr/local/cuda-10.0/samples/common/inc)

file(GLOB_RECURSE CURRENT_HEADERS *.h)
file(GLOB CURRENT_SOURCES *.cu)

source_group("Include" FILES ${CURRENT_HEADERS})
source_group("Source" FILES ${CURRENT_SOURCES})

cuda_add_executable(quick_sort ${CURRENT_HEADERS} ${CURRENT_SOURCES})

target_link_libraries(quick_sort ${CUDA_cudadevrt_LIBRARY})
target_link_libraries(quick_sort ${CUDA_cudart_LIBRARY})


报错为:
CMakeFiles/quick_sort.dir/quick_sort_generated_cdpAdvancedQuicksort.cu.o: In function `__sti____cudaRegisterAll()':
/tmp/tmpxft_00001604_00000000-5_cdpAdvancedQuicksort.cudafe1.stub.c:7: undefined reference to `__cudaRegisterLinkedBinary_55_tmpxft_00001604_00000000_6_cdpAdvancedQuicksort_cpp1_ii_1c7c59e6'
CMakeFiles/quick_sort.dir/quick_sort_generated_cdpBitonicSort.cu.o: In function `__sti____cudaRegisterAll()':
/tmp/tmpxft_000015e8_00000000-5_cdpBitonicSort.cudafe1.stub.c:2: undefined reference to `__cudaRegisterLinkedBinary_49_tmpxft_000015e8_00000000_6_cdpBitonicSort_cpp1_ii_e829267e'
collect2: error: ld returned 1 exit status
make[3]: *** [quick_sort] Error 1
make[2]: *** [CMakeFiles/quick_sort.dir/all] Error 2
make[1]: *** [CMakeFiles/quick_sort.dir/rule] Error 2
make: *** [quick_sort] Error 2
10
CUDA / Re: 我是做cuda实时信号处理...
« 最后发表 作者 sisiy 一月 08, 2021, 02:16:48 pm »
答复:这个问题比较有意思, 涉及到了GPU的吞吐率和延迟的问题. 楼主说的40Mhz采样, 得到40MB/s, 等于说这个是8-bit采样, 每1ms内采样40K次. 同时, 楼主说了需要1ms就完成一个反馈或者控制, 等于说, GPU需要每批次只处理40KB数据, 还需要在1ms内处理完. 这是问题的前提.

那么我们就需要从吞吐率和延迟两个角度, 来分析楼主的"GPU很难做到实时"的问题. 如果楼主的GPU实现, 每次kernel调用需要较大的延迟, 例如25ms, 但是吞吐率足够(例如等效每秒能完成几个GB的数据的运算), 那么移动"CPU上的控制逻辑到GPU"是没用的. 因为你的延迟在那里. 此时不如考虑损失一定的GPU的吞吐率来换取处理的延迟了, 这个需要具体问题具体分析.

而如果楼主本身的GPU一次kernel调用的延迟足够, 则他并不需要将"CPU上的逻辑移动到GPU". 哪怕一次处理KB级别的数据降低了GPU的有效利用率, 但是满足任务要求就好.最坑爹的是kernel的延迟满足不了要求, 吞吐率性能也不能满足要求, 那样我估计楼主得好好重新设计, 整体改写贵kernel了....
页: [1] 2 3 ... 10