atomiccas和其他原子操作的区别

  • 2 replies
  • 1679 views
atomiccas和其他原子操作的区别
« 于: 二月 10, 2022, 03:38:12 pm »
atomicCAS函数是比较特别的原子操作,其他的原子操作都可以用这个原子操作函数实现,对于这个函数我只是大概知道是comapre and swap。
我在测试原子操作相关的metrics时候,发现ncu特意将CAS和其他的原子操作区分开了:
lts__t_requests_srcnode_gpc_aperture_device_op_atom_dot_alu # of LTS requests from node GPC accessing device memory (vidmem) for atomic ALU (non-CAS)
lts__t_sectors_srcnode_gpc_aperture_device_op_atom_dot_alu # of LTS sectors from node GPC accessing device memory (vidmem) for atomic ALU (non-CAS)
lts__t_requests_srcnode_gpc_aperture_device_op_atom_dot_cas # of LTS requests from node GPC accessing device memory (vidmem) for atomic CAS
lts__t_sectors_srcnode_gpc_aperture_device_op_atom_dot_cas # of LTS sectors from node GPC accessing device memory (vidmem) for atomic CAS
所以我做了一些测试,所有的测试均在global memory上,都只使用一个warp的线程:

测试1:对同一个位置的数据做原子加
程序代码: [选择]
atomicAdd(&input[0], threadIdx.x);结果:
 
引用
---------------------------------------------------------------------- --------------- ------------------------------
    lts__t_requests_srcnode_gpc_aperture_device_op_atom_dot_alu.sum                request                             32
    lts__t_requests_srcnode_gpc_aperture_device_op_atom_dot_cas.sum                request                              0
    lts__t_sectors_srcnode_gpc_aperture_device_op_atom_dot_alu.sum                  sector                             32
    lts__t_sectors_srcnode_gpc_aperture_device_op_atom_dot_cas.sum                  sector                              0
    ---------------------------------------------------------------------- --------------- ------------------------------
这个结果我认为是因为一个warp中所有线程进行原子加时串行执行,所以每一个线程都需要发出一个request,对应一个sector;

测试2:对不同位置的数据做原子加
程序代码: [选择]
atomicAdd(&input[threadIdx.x], threadIdx.x);结果:
引用
    ---------------------------------------------------------------------- --------------- ------------------------------
    lts__t_requests_srcnode_gpc_aperture_device_op_atom_dot_alu.sum                request                              1
    lts__t_requests_srcnode_gpc_aperture_device_op_atom_dot_cas.sum                request                              0
    lts__t_sectors_srcnode_gpc_aperture_device_op_atom_dot_alu.sum                  sector                              4
    lts__t_sectors_srcnode_gpc_aperture_device_op_atom_dot_cas.sum                  sector                              0
    ---------------------------------------------------------------------- --------------- ------------------------------
这个结果我认为是因为一个warp中所有线程进行原子加时可以合并,所以一个warp只需要发出一个request,对应4个sector;

测试3:对同一个位置的数据做atomicCAS
atomicCAS(&input[0], 100, threadIdx.x);
测试结果:
引用
    Section: Command line profiler metrics
    ---------------------------------------------------------------------- --------------- ------------------------------
    lts__t_requests_srcnode_gpc_aperture_device_op_atom_dot_alu.sum                request                              0
    lts__t_requests_srcnode_gpc_aperture_device_op_atom_dot_cas.sum                request                             32
    lts__t_sectors_srcnode_gpc_aperture_device_op_atom_dot_alu.sum                  sector                              0
    lts__t_sectors_srcnode_gpc_aperture_device_op_atom_dot_cas.sum                  sector                             32
    ---------------------------------------------------------------------- --------------- ------------------------------
这个结果感觉和测试1很类似,也能讲得通。
测试4:对不同位置的数据做atomicCAS
程序代码: [选择]
atomicCAS(&input[threadIdx.x], 100, threadIdx.x)测试结果:
引用
   Section: Command line profiler metrics
    ---------------------------------------------------------------------- --------------- ------------------------------
    lts__t_requests_srcnode_gpc_aperture_device_op_atom_dot_alu.sum                request                              0
    lts__t_requests_srcnode_gpc_aperture_device_op_atom_dot_cas.sum                request                              2
    lts__t_sectors_srcnode_gpc_aperture_device_op_atom_dot_alu.sum                  sector                              0
    lts__t_sectors_srcnode_gpc_aperture_device_op_atom_dot_cas.sum                  sector                              4
    ---------------------------------------------------------------------- --------------- ------------------------------
这个结果我有点不太理解,应该是可以合并的,但是为什么会发出了两个request,对应4个sector?


Re: atomiccas和其他原子操作的区别
« 回复 #1 于: 三月 10, 2022, 01:11:37 pm »
atomicCAS函数是比较特别的原子操作,其他的原子操作都可以用这个原子操作函数实现,对于这个函数我只是大概知道是comapre and swap。
我在测试原子操作相关的metrics时候,发现ncu特意将CAS和其他的原子操作区分开了:
lts__t_requests_srcnode_gpc_aperture_device_op_atom_dot_alu # of LTS requests from node GPC accessing device memory (vidmem) for atomic ALU (non-CAS)
lts__t_sectors_srcnode_gpc_aperture_device_op_atom_dot_alu # of LTS sectors from node GPC accessing device memory (vidmem) for atomic ALU (non-CAS)
lts__t_requests_srcnode_gpc_aperture_device_op_atom_dot_cas # of LTS requests from node GPC accessing device memory (vidmem) for atomic CAS
lts__t_sectors_srcnode_gpc_aperture_device_op_atom_dot_cas # of LTS sectors from node GPC accessing device memory (vidmem) for atomic CAS
所以我做了一些测试,所有的测试均在global memory上,都只使用一个warp的线程:

测试1:对同一个位置的数据做原子加
程序代码: [选择]
atomicAdd(&input[0], threadIdx.x);结果:
 这个结果我认为是因为一个warp中所有线程进行原子加时串行执行,所以每一个线程都需要发出一个request,对应一个sector;

测试2:对不同位置的数据做原子加
程序代码: [选择]
atomicAdd(&input[threadIdx.x], threadIdx.x);结果:这个结果我认为是因为一个warp中所有线程进行原子加时可以合并,所以一个warp只需要发出一个request,对应4个sector;

测试3:对同一个位置的数据做atomicCAS
atomicCAS(&input[0], 100, threadIdx.x);
测试结果:这个结果感觉和测试1很类似,也能讲得通。
测试4:对不同位置的数据做atomicCAS
程序代码: [选择]
atomicCAS(&input[threadIdx.x], 100, threadIdx.x)测试结果:这个结果我有点不太理解,应该是可以合并的,但是为什么会发出了两个request,对应4个sector?

我建议你总是先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], ...)时候的效果。


Re: atomiccas和其他原子操作的区别
« 回复 #2 于: 三月 10, 2022, 04:25:55 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倍,有点不理解