最新帖子

页: 1 2 [3] 4 5 ... 10
21
CUDA / Re: 同一个CUDA流操作
« 最后发表 作者 屠戮人神 于 三月 22, 2022, 09:19:59 pm »
引用自: jinyer 于 三月 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();吗？

你在偷换概念， 每个流中的任务（kernel计算任务，或者数据传输任务）的确是按照顺序进行的。但是这不代表，夹杂在中间的其他CPU上的操作，也要按照GPU的流中的顺序。

实际上你的这个例子是，每次kernel启动后，CPU不必等待kernel执行完毕，就可以继续往下执行了（kernel的启动是异步的）。所以你看到这样的是正常的。

只有计算和传输任务才会在同1个流中（你的例子中是默认流，我建议你使用非默认流避免你遭遇其他情况的问题）按照你发布命令的顺序执行的，例如你的第二问中的，连续2个kernel，依次顺序执行，这点是没有问题的。
22
CUDA / 请教：现在较新的GPU规定一个寄存器内存大小了吗？
« 最后发表 作者 cudaxiaox 于 三月 21, 2022, 08:53:38 am »
请问像TITAN V、RTX3090这些显卡有没有规定CUDA编程时一个线程内可分配的存储在寄存器上的变量最多占多大内存呢？
23
CUDA / Re: 同一个CUDA流操作
« 最后发表 作者 jinyer 于 三月 16, 2022, 02:58:50 pm »
是异步调度，但实际执行的时候还是要等第一个kernel执行完再会执行第二个kernel？或者是说 并行资源足够的情况下，两个核函数会一起执行呢
24
CUDA / 同一个CUDA流操作
« 最后发表 作者 jinyer 于 三月 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();吗？
25
CUDA / Re: 关于threadfence、volatile、threadfence_block、threadfence_system、syncthreads的一些见解和问题
« 最后发表 作者 LibAndLab 于 三月 10, 2022, 05:10:29 pm »
引用自: 屠戮人神 于 三月 10, 2022, 02:56:56 pm
有可能是这个原因。但是去掉了volatile后永远不会检测到正确的flag了。

实际上有没有volatile是两种指令，
一种是：
LDG.STRONG.SYS
另外一种：LDG
所以可能L2也会缓冲系统内存（因为后一种多次读取都检测不到，可能是kernel运行期间将旧值缓冲了）。
嗯嗯从结果来看是会缓冲的
26
CUDA / Re: atomiccas和其他原子操作的区别
« 最后发表 作者 LibAndLab 于 三月 10, 2022, 04:25:55 pm »
引用自: 屠戮人神 于 三月 10, 2022, 01:11:37 pm
我建议你总是先cuobjdump --dump-sass看下编译出来的8.6上的结果， 然后再考虑解释结果。因为已知有的时候，编译出来的执行的操作，和你的原始写法不同。

不过对于你的atomicCAS(&input[threadIdx.x], 100, threadIdx.x)倒是编译出来的和写出来的是一样的，大致是：
R5 = threadIdx.x;
R2 = 4;
R4 = 100;
PTR = P + R5 * R2
ATOMG...CAS...GPU(PTR, R4, R5);
所以还是符合你的原意的。

那只能说，这里的request可能最多代表一个64B的范围内， 或者干脆你的原始指针没有对齐到32B的边界（如果你是cudaMalloc来的，可以无视，因为它有最低对齐性保证）。但是具体的原因我不知道。

不妨试试看看atomicCAS(&threadIdx.x[32 * threadIdx.x], ...)时候的效果。
测试了不同的间隔：
程序代码: [选择]
__global__ void LtsTRequestsApertureDeviceOpAtomDotCasKernel0(unsigned int *input, unsigned int *output) {
    output[threadIdx.x] = atomicMin(&input[1 * threadIdx.x], threadIdx.x);
    atomicCAS(&input[1 * threadIdx.x], 1000, threadIdx.x);
}

__global__ void LtsTRequestsApertureDeviceOpAtomDotCasKernel1(unsigned int *input, unsigned int *output) {
    output[threadIdx.x] = atomicMin(&input[2 * threadIdx.x], threadIdx.x);
    atomicCAS(&input[2 * threadIdx.x], 1000, threadIdx.x);
}

__global__ void LtsTRequestsApertureDeviceOpAtomDotCasKernel2(unsigned int *input, unsigned int *output) {
    output[threadIdx.x] = atomicMin(&input[4 * threadIdx.x], threadIdx.x);
    atomicCAS(&input[4 * threadIdx.x], 1000, threadIdx.x);
}

__global__ void LtsTRequestsApertureDeviceOpAtomDotCasKernel3(unsigned int *input, unsigned int *output) {
    output[threadIdx.x] = atomicMin(&input[8 * threadIdx.x], threadIdx.x);
    atomicCAS(&input[8 * threadIdx.x], 1000, threadIdx.x);
}

__global__ void LtsTRequestsApertureDeviceOpAtomDotCasKernel4(unsigned int *input, unsigned int *output) {
    output[threadIdx.x] = atomicMin(&input[16 * threadIdx.x], threadIdx.x);
    atomicCAS(&input[16 * threadIdx.x], 1000, threadIdx.x);
}

__global__ void LtsTRequestsApertureDeviceOpAtomDotCasKernel5(unsigned int *input, unsigned int *output) {
    output[threadIdx.x] = atomicMin(&input[32 * threadIdx.x], threadIdx.x);
    atomicCAS(&input[32 * threadIdx.x], 1000, threadIdx.x);
}
结果如下
程序代码: [选择]
    Section: Command line profiler metrics
    ---------------------------------------------------------------------- --------------- ------------------------------
    lts__t_requests_aperture_device_op_atom_dot_alu.sum                            request                              1
    lts__t_requests_aperture_device_op_atom_dot_cas.sum                            request                              2
    lts__t_requests_aperture_device_op_atom_dot_cas_lookup_hit.sum                 request                              2
    lts__t_requests_aperture_device_op_atom_lookup_hit.sum                         request                              2
    lts__t_requests_aperture_device_op_atom_lookup_miss.sum                        request                              1
    ---------------------------------------------------------------------- --------------- ------------------------------

  LtsTRequestsApertureDeviceOpAtomDotCasKernel1(unsigned int *, unsigned int *), 2022-Mar-10 08:22:27, Context 1, Stream 7
    Section: Command line profiler metrics
    ---------------------------------------------------------------------- --------------- ------------------------------
    lts__t_requests_aperture_device_op_atom_dot_alu.sum                            request                              2
    lts__t_requests_aperture_device_op_atom_dot_cas.sum                            request                              4
    lts__t_requests_aperture_device_op_atom_dot_cas_lookup_hit.sum                 request                              4
    lts__t_requests_aperture_device_op_atom_lookup_hit.sum                         request                              4
    lts__t_requests_aperture_device_op_atom_lookup_miss.sum                        request                              2
    ---------------------------------------------------------------------- --------------- ------------------------------

  LtsTRequestsApertureDeviceOpAtomDotCasKernel2(unsigned int *, unsigned int *), 2022-Mar-10 08:22:27, Context 1, Stream 7
    Section: Command line profiler metrics
    ---------------------------------------------------------------------- --------------- ------------------------------
    lts__t_requests_aperture_device_op_atom_dot_alu.sum                            request                              4
    lts__t_requests_aperture_device_op_atom_dot_cas.sum                            request                              8
    lts__t_requests_aperture_device_op_atom_dot_cas_lookup_hit.sum                 request                              8
    lts__t_requests_aperture_device_op_atom_lookup_hit.sum                         request                              8
    lts__t_requests_aperture_device_op_atom_lookup_miss.sum                        request                              4
    ---------------------------------------------------------------------- --------------- ------------------------------

  LtsTRequestsApertureDeviceOpAtomDotCasKernel3(unsigned int *, unsigned int *), 2022-Mar-10 08:22:27, Context 1, Stream 7
    Section: Command line profiler metrics
    ---------------------------------------------------------------------- --------------- ------------------------------
    lts__t_requests_aperture_device_op_atom_dot_alu.sum                            request                              8
    lts__t_requests_aperture_device_op_atom_dot_cas.sum                            request                             16
    lts__t_requests_aperture_device_op_atom_dot_cas_lookup_hit.sum                 request                             16
    lts__t_requests_aperture_device_op_atom_lookup_hit.sum                         request                             16
    lts__t_requests_aperture_device_op_atom_lookup_miss.sum                        request                              8
    ---------------------------------------------------------------------- --------------- ------------------------------

  LtsTRequestsApertureDeviceOpAtomDotCasKernel4(unsigned int *, unsigned int *), 2022-Mar-10 08:22:27, Context 1, Stream 7
    Section: Command line profiler metrics
    ---------------------------------------------------------------------- --------------- ------------------------------
    lts__t_requests_aperture_device_op_atom_dot_alu.sum                            request                             16
    lts__t_requests_aperture_device_op_atom_dot_cas.sum                            request                             32
    lts__t_requests_aperture_device_op_atom_dot_cas_lookup_hit.sum                 request                             32
    lts__t_requests_aperture_device_op_atom_lookup_hit.sum                         request                             32
    lts__t_requests_aperture_device_op_atom_lookup_miss.sum                        request                             16
    ---------------------------------------------------------------------- --------------- ------------------------------

  LtsTRequestsApertureDeviceOpAtomDotCasKernel5(unsigned int *, unsigned int *), 2022-Mar-10 08:22:27, Context 1, Stream 7
    Section: Command line profiler metrics
    ---------------------------------------------------------------------- --------------- ------------------------------
    lts__t_requests_aperture_device_op_atom_dot_alu.sum                            request                             32
    lts__t_requests_aperture_device_op_atom_dot_cas.sum                            request                             32
    lts__t_requests_aperture_device_op_atom_dot_cas_lookup_hit.sum                 request                             32
    lts__t_requests_aperture_device_op_atom_lookup_hit.sum                         request                             32
    lts__t_requests_aperture_device_op_atom_lookup_miss.sum                        request                             32
    ---------------------------------------------------------------------- --------------- ------------------------------

有点奇怪 只有是32*threadIdx.x的时候两者才相同，其他情况下都是cas的request是非cas的2倍，有点不理解
27
CUDA / Re: 关于threadfence、volatile、threadfence_block、threadfence_system、syncthreads的一些见解和问题
« 最后发表 作者 屠戮人神 于 三月 10, 2022, 02:56:56 pm »
引用自: LibAndLab 于 三月 10, 2022, 02:15:53 pm
https://developer.download.nvidia.com/video/gputechconf/gtc/2020/presentations/s21819-optimizing-applications-for-nvidia-ampere-gpu-architecture.pdf这篇文章第七页中图示对于system memory不会在L2 上cache，上述的例子会不会是这个原因

有可能是这个原因。但是去掉了volatile后永远不会检测到正确的flag了。

实际上有没有volatile是两种指令，
一种是：
LDG.STRONG.SYS
另外一种：LDG
所以可能L2也会缓冲系统内存（因为后一种多次读取都检测不到，可能是kernel运行期间将旧值缓冲了）。
28
CUDA / Re: 关于threadfence、volatile、threadfence_block、threadfence_system、syncthreads的一些见解和问题
« 最后发表 作者 LibAndLab 于 三月 10, 2022, 02:15:53 pm »
引用自: 屠戮人神 于 三月 10, 2022, 01:37:17 pm
（1）如同我们用的CPU编译器（例如gcc好了）在特定的CPU硬件上的编译。任何对threadfence的指定都有2个效果，一个是软件效果（编译器效果），一个是硬件效果（CPU）。GPU也同理。

前者可能会导致编译器对threadfence前后的读写指令的妥当安排（例如可能不会跨越fence将后面的读取指令调到前面，等于说你用fence限制了编译器的对你的原始代码行的自由调整的发挥）。后者可能在特定的硬件上生成特定的指令，要求硬件确定特定的指令生效顺序（例如可能要求某些写入指令在某些单元上，至少执行到某些特定的结果，从而使得其他单元能够看到写入效果）。

（2）__syncthreads()隐含了fence效果。等于fence调用一次+线程块同步一次（+可选的块内数据迭代）（看后缀）。

（3）volatile我不懂。但有的时候是可以越过L2的。例如这个例子：
__global__ void labandlib(...., volatile int *flag)
{
    while(....)
    {
         ....//long running kernel code
         if (*flag == 0x12345678) return; //early return
         ....//long running kernel code   
    }
}
其中flag是通过cudaMallocHost分配的，并随时通过CPU设定值来打断长期运行的kernel的执行的。

这种代码可以正常执行。所以它在某些情况下至少得bypass L2. 更多的得你教晓我。
https://developer.download.nvidia.com/video/gputechconf/gtc/2020/presentations/s21819-optimizing-applications-for-nvidia-ampere-gpu-architecture.pdf这篇文章第七页中图示对于system memory不会在L2 上cache，上述的例子会不会是这个原因

29
CUDA / Re: 关于threadfence、volatile、threadfence_block、threadfence_system、syncthreads的一些见解和问题
« 最后发表 作者 屠戮人神 于 三月 10, 2022, 01:37:17 pm »
引用自: LibAndLab 于 二月 24, 2022, 05:01:21 pm
volatile：
读写数据都会bypass掉L1 cache；

syncthreads：
1. (threadblock-wide) execution barrier，当同一个block中的线程都执行到该处时才会继续往下执行；
2. (device-wide) memory fence， 保证该语句前的write的数据对同一个block中的其他线程可见，这个可见意味着数据被写入L2 cache，同时因为L1 cache是write through cache，所以L1 cache中的数据也会被更新；

threadfence：
1. 有了这个指令后, 线程则会暂停, 等待写入的过程完成到一定程度,保证线程自己的写入操作能被整个grid的线程"看"到, 才会继续执行，但是不会导致其他线程的停止运行，这里的看到也是指数据被写入L2 cache，同时因为L1 cache是write through cache，所以L1 cache中的数据也会被更新，但是仅限于已经执行该指令的SM，对于其他的SM该数据的缓存则不会改变或者失效，其他的线程仍然有可能会从L1 cache中读取到老旧的数据；
2. 将threadfence函数前后的访存隔离开, 编译器不会跨越threadfence的边界, 不会重新调整访存语句的顺序


但是threadfence_block和threadfence_system我是不太清楚的，这里的可见性具体是写入到什么程度呢？
threadfence_block也是写入到L2 cache才算对同一个block内的线程可见？ threadfence_system是写入到global memory才算对多个显卡的线程可见？

（1）如同我们用的CPU编译器（例如gcc好了）在特定的CPU硬件上的编译。任何对threadfence的指定都有2个效果，一个是软件效果（编译器效果），一个是硬件效果（CPU）。GPU也同理。

前者可能会导致编译器对threadfence前后的读写指令的妥当安排（例如可能不会跨越fence将后面的读取指令调到前面，等于说你用fence限制了编译器的对你的原始代码行的自由调整的发挥）。后者可能在特定的硬件上生成特定的指令，要求硬件确定特定的指令生效顺序（例如可能要求某些写入指令在某些单元上，至少执行到某些特定的结果，从而使得其他单元能够看到写入效果）。

（2）__syncthreads()隐含了fence效果。等于fence调用一次+线程块同步一次（+可选的块内数据迭代）（看后缀）。

（3）volatile我不懂。但有的时候是可以越过L2的。例如这个例子：
__global__ void labandlib(...., volatile int *flag)
{
    while(....)
    {
         ....//long running kernel code
         if (*flag == 0x12345678) return; //early return
         ....//long running kernel code   
    }
}
其中flag是通过cudaMallocHost分配的，并随时通过CPU设定值来打断长期运行的kernel的执行的。

这种代码可以正常执行。所以它在某些情况下至少得bypass L2. 更多的得你教晓我。
30
CUDA / Re: 原子操作、L2 cache、sysmem
« 最后发表 作者 屠戮人神 于 三月 10, 2022, 01:28:58 pm »
引用自: LibAndLab 于 二月 23, 2022, 10:37:00 am
RED 和ATOM指令的区别经过查阅文档，我已经大概清楚，RED是Reduction+atomic，而ATOM仅仅是Atomic，RED运行耗时更短

都是原子操作。前者不需要返回值，从而可能减少了一些执行上的成本，例如不需要设定对指令的写入结果值（例如旧值）的寄存器的结果写入完成的跟踪之类的（scoreboarding).
页: 1 2 [3] 4 5 ... 10