列出帖子

该操作将允许你查看该会员所有的帖子，注意你只能看到你有权限看到的板块的帖子。


显示所有帖子 - jinyer

页: [1] 2
1
CUDA / Re: 使用了共享内存最终时间没有没有减少
« 于: 十一月 01, 2022, 09:15:40 pm »
引用自: 屠戮人神 于 十一月 01, 2022, 01:34:45 pm
并不存在真正"全局可见"的shared memory的，例如以下代码，会给你造成幻觉：
__shared__ int fake_global_visible_shared_array[100];
__global__ void your_kernel()
{
   ....//use fake_global_visible_shared_array[] here
}

这种依然等效于在your_kernel的内部，定义一个shared memory的数组的。虽然它看起来像是"全局的shared memory", 然而它不是。

所以并不能跨kernel多次使用（例如你的一个kernel写入，另外一个kernel读取的）。

其次，因为每次kernel启动前， shared memory中的数据并不清零，你存在概率会读到之前的某kernel的某个block的残留内容的，但这种并不可靠（你也可以读不到，或者读到错误的内容）。

所以综合这两点，这种写法不可取。
那意思就是如果想要借助共享内存来减少存取延时，必须要在一个核函数范围内了吗？ 
假设想要把这3个kernel合并在一个kernel中，因为寄存器的原因导致launch失败【寄存器数量超出硬件限制】，即不能合并成一个kernel，那这种情况下只能把 U数组 存在global memory了吗？

2
CUDA / Re: 使用了共享内存最终时间没有没有减少
« 于: 十一月 01, 2022, 10:39:51 am »
引用自: 屠戮人神 于 十月 31, 2022, 04:42:51 pm
Shared Memory的使用，不能跨kernel生效。你在kernel 1中计算U，和在Kernel2、3中又使用到U，这种不能通过Shared memory进行缓冲（因为shared memory中的数据有效生存期是1次kernel启动中的1个block，不能跨block，更不能跨多次kernel启动）。

这种情况依赖于GPU的L1和L2来缓冲即可。正常情况下，他们可以跨多个kernel启动有效。特殊情况下（例如6.1上的某些情况中），至少L2会生效。
感谢回复！
但是代码中我是把共享内存U声明为了全局可见， kernel 1中用计算结果初始化共享数组U，kernel2和kernel3 从共享内存U中取数据来完成自己的计算，结果表示确实正确得到并使用了kernel 1的计算结果。相较于之前数组U是存放在global memory，现在存到了shared memory，这难道不能起到加速作用吗？【另外这是不是也说明了全局可见的共享内存可以跨kernel生效呀？】

3
CUDA / 使用了共享内存最终时间没有没有减少
« 于: 十月 19, 2022, 03:23:15 pm »
程序中 kernel 1得到u值，kernel 2和kernel 3都要使用这个u值，因此我申请了全局可见的共享内存用来存储u，以减少2、3步中读取u的延时，可是为什么使用共享内存优化之后runtime没减少呢？
如下图，可以明显看出访问全局内存的次数明显变少，但时间却从1.96ms变成了1.97ms【这是使用nsight compute记录的核函数执行时间】

4
CUDA / Re: GPU寄存器使用
« 于: 九月 15, 2022, 10:54:47 am »
好的，谢谢您的解答。我之前还以为如果寄存器被使用完，剩余的部分编译器会自动使用local memory来代替，从而不会出现限制occupancy的情况。

5
CUDA / GPU寄存器使用
« 于: 九月 07, 2022, 03:52:21 pm »
手册上说若寄存器不足，会使用本地内存，而本地内存位于全局内存上，理论上不会有限制。
3090卡1个SM理论可以同时调度1536个线程，可是在实际使用中之恩那个调度1152个线程，通过ncu查看当前函数执行时一个线程使用寄存器56个【65536/1152约等于56】.这说明寄存器限制了SM可同时调度的线程数量。
请问为什么没有使用本地内存呢？

6
CUDA / Re: 线程分化-以warp为分支单位
« 于: 五月 25, 2022, 09:35:56 am »
引用自: jinyer 于 五月 25, 2022, 09:09:00 am
感谢您的解答。
请问您说的异步载入是指这个吧？
程序代码: [选择]
__pipeline_memcpy_async(&shared[blockDim.x * i + threadIdx.x],
                            &global[blockDim.x * i + threadIdx.x], sizeof(T));这个代码的效果是不是和我上述使用if~else效果相同，只是异步载入使用起来会更方便？
但是我又有一个疑问，如果使用异步载入的话，执行计算是一个核函数，进行异步载入又是一个核函数，可是在单流中，两个核函数只能串行进行，那就达不到计算和搬运数据同时进行的效果了。

7
CUDA / Re: 线程分化-以warp为分支单位
« 于: 五月 25, 2022, 09:09:00 am »
感谢您的解答。
请问您说的异步载入是指这个吧？
程序代码: [选择]
__pipeline_memcpy_async(&shared[blockDim.x * i + threadIdx.x],
                            &global[blockDim.x * i + threadIdx.x], sizeof(T));这个代码的效果是不是和我上述使用if~else效果相同，只是异步载入使用起来会更方便？

8
CUDA / Re: 线程分化-以warp为分支单位
« 于: 五月 09, 2022, 11:18:22 am »
或者怎么看作用于if的和作用于else的不同warp是否同时执行了呢

9
CUDA / 线程分化-以warp为分支单位
« 于: 五月 06, 2022, 07:32:19 pm »
对开发者而言，应该尽量避免在同一个Warp中有不同的执行路径。当必须写分支语句时，尽量让分支宽度大于WarpSize，即不以Thread为分支单位，而是以Warp为分支单位，这样就能保证同一个Warp内部不会出现分支。
请问如果分配线程块 block(256)
核函数中是
程序代码: [选择]
int warpIdx = threadIdx.x/32;
if(warpIdx < 2){
    执行语句1；
}else{
    执行语句2；
}代码是以warp为分支单位，那这种情况下if和else可以同时运行吗？
即现在有一个需求，想让线程块中一部分线程来执行计算，另一部分线程来把数据从全局内存搬运到共享内存，请问我可以通过上述if-else达到这两种需求同时运行的效果吗？

10
CUDA / Re: Nsight Compute没有权限问题（linux系统上）
« 于: 四月 18, 2022, 10:04:58 am »
引用自: jinyer 于 四月 15, 2022, 09:43:08 pm
直接输入 ncu-ui 可以打开compute，就是launch时核函数会报没有权限的错误；
sudo ncu-ui 会出现说没有qt平台导致不能start the app【而上边直接ncu-ui软件是可以使用的，这是为什么呢】
命令执行如下：
dell@dell-Precision-5820-Tower-X-Series:~$ ncu-ui
********BP-RCM----nFrame:1, FrameLen=4096, mea=8192, iteration=1  [cudaEvent]***********
dell@dell-Precision-5820-Tower-X-Series:~$ sudo ncu-ui
No protocol specified
qt.qpa.xcb: could not connect to display :10.0
qt.qpa.plugin: Could not load the Qt platform plugin "xcb" in "" even though it was found.
This application failed to start because no Qt platform plugin could be initialized. Reinstalling the application may fix this problem.

Available platform plugins are: xcb.

/usr/local/cuda-11.2/bin/../nsight-compute-2020.3.0/host/linux-desktop-glibc_2_11_3-x64/ncu-ui：行 16: 3025862 已放弃               （核心已转储） "$NV_AGORA_PATH/CrashReporter" "NVIDIA Nsight Compute" "NVIDIA Nsight Compute" "2020.3.0.0 (build 29307467) (public-release)" "$NV_AGORA_PATH/ncu-ui.bin" "$@"
已解决，感谢回答

11
CUDA / Re: Nsight Compute没有权限问题（linux系统上）
« 于: 四月 15, 2022, 09:43:08 pm »
引用自: LibAndLab 于 四月 15, 2022, 06:39:27 pm
sudo ncu 也不行吗
直接输入 ncu-ui 可以打开compute，就是launch时核函数会报没有权限的错误；
sudo ncu-ui 会出现说没有qt平台导致不能start the app【而上边直接ncu-ui软件是可以使用的，这是为什么呢】
命令执行如下：
dell@dell-Precision-5820-Tower-X-Series:~$ ncu-ui
********BP-RCM----nFrame:1, FrameLen=4096, mea=8192, iteration=1  [cudaEvent]***********
dell@dell-Precision-5820-Tower-X-Series:~$ sudo ncu-ui
No protocol specified
qt.qpa.xcb: could not connect to display :10.0
qt.qpa.plugin: Could not load the Qt platform plugin "xcb" in "" even though it was found.
This application failed to start because no Qt platform plugin could be initialized. Reinstalling the application may fix this problem.

Available platform plugins are: xcb.

/usr/local/cuda-11.2/bin/../nsight-compute-2020.3.0/host/linux-desktop-glibc_2_11_3-x64/ncu-ui：行 16: 3025862 已放弃               （核心已转储） "$NV_AGORA_PATH/CrashReporter" "NVIDIA Nsight Compute" "NVIDIA Nsight Compute" "2020.3.0.0 (build 29307467) (public-release)" "$NV_AGORA_PATH/ncu-ui.bin" "$@"

12
CUDA / Nsight Compute没有权限问题（linux系统上）
« 于: 四月 13, 2022, 09:34:44 am »
遇到的问题：launch之后每个核函数都会出现ERR_NVGPUCTRPERM - The user does not have permission to profile on the target device这个错误
查阅解决方法大家一般都是两点：
1、使用管理员权限运行nsight compute，但这好像是针对windows系统，对linux不太适用；
2、https://developer.nvidia.com/nvidia-development-tools-solutions-err_nvgpuctrperm-permission-issue-performance-counters官网给出了启用该权限的方法，但我按步骤做的时候这一步会出错：modprobe -r nvidia_uvm nvidia_drm nvidia_modeset nvidia-vgpu-vfio nvidia[2] # Unload dependent modules
，不知道要怎么做，请教大家

13
CUDA / SM、block、warp
« 于: 四月 01, 2022, 03:56:39 pm »
一个Multi-Processor上面可以同时驻留多个warps，例如一个计算能力8.6的卡，它的1个SM里面，最多可能驻留1536个线程，也就是大约48个warps，在同时等待被调度执行。并且同时最多能上16个block。
请问一个SM是同时在执行16个block吗？
若是的话，SM一个时刻只能执行每个block里边的一个warp吗？即并行线程数为 82*16*32吗？【目前我的理解是这样，可3090卡的SP个数只有10496是82*16*32的四分之一，这又是怎么做到的呢】

14
CUDA / Re: 同一个CUDA流操作
« 于: 三月 16, 2022, 02:58:50 pm »
是异步调度，但实际执行的时候还是要等第一个kernel执行完再会执行第二个kernel？或者是说 并行资源足够的情况下，两个核函数会一起执行呢

15
CUDA / 同一个CUDA流操作
« 于: 三月 16, 2022, 11:37:08 am »
看到大家说：同一个CUDA流操作是严格顺序的。
但是我跑代码时却出现了不一样的结果
程序代码: [选择]
noisePdf1<<<gridNoiPdf1, blockNoiPdf1 >>>(noiPdf1, d_rx, sigma);
printf("After---------------noisePdf1\n");
convPy1<<<gridConvPy, blockConvPy >>>(convResult, u0, u1);
printf("After------------------convPy\n");
然后分别在两个核函数中有一个printf语句。
在默认流中先后调用这两个核函数，结果输出语句如下：
After---------------noisePdf1
After------------------convPy
run----noisePdf1
running------convPy
请问这是表示如果第二个核函数想用第一个核函数的结果，这两个核函数调用中间必须加cudaDeviceSynchronize();吗？

页: [1] 2