无法将具有相同优先级的对象插入ConcurrentSkipListSet(Unable to insert objects of equal priority into ConcurrentSkipListSet)

 事实证明，我没有正确理解Keppler架构。 正如Greg Smith上面的评论中指出的那样，Keppler可以配置为具有8个字节的32个共享存储体。 在这种情况下，使用cudaDeviceSetSharedMemConfig( cudaSharedMemBankSizeEightByte ) ，共享内存布局如下所示： 
 
bank:  B0   B1   B2   B3   B4    ..   B31
       ----------------------------------
index: D00  D01  D02  D03  D04   ..   D31
       D32  D33  D34  D35  D36   ..   D63   
 
 现在，对于我的简单示例（使用itot=16 ），在一个warp中处理对例如前两行（ threadIdx.y=0 ， threadIdx.y=1 ）的共享内存的写入/读取。 这意味着对于threadIdx.y=0值D00..D15存储在B0..B15 ，则存在两个双精度的填充，之后在相同的变形值D18..D33中存储在B18..B31+B00..B01 ，导致B00-B01发生银行冲突。 如果没有填充（ ng=0 ），第一行将写入D00..D15中的B00..B15 ，即D00..D15中B00..B15中的第二行，因此不会发生库冲突。 
 
 对于blockDim.x>=32的线程块，不应该出现问题。 例如，对于itot=32 ， blockDim.x=32 ， ng=2 ，第一行存储在存储体B00..B31 ，然后存储两个单元格填充，第二行存储在B02..B31+B00..B01等中。 

It turns out that I didn’t understand the Keppler architecture correctly. As pointed out in one of the comments above by Greg Smith, Keppler can be configured to have 32 shared memory banks of 8 bytes. In such a case, using cudaDeviceSetSharedMemConfig( cudaSharedMemBankSizeEightByte ), the shared memory layout looks like:
 
bank:  B0   B1   B2   B3   B4    ..   B31
       ----------------------------------
index: D00  D01  D02  D03  D04   ..   D31
       D32  D33  D34  D35  D36   ..   D63   
 
Now, for my simple example (using itot=16), the writing/reading to/from shared memory on e.g. the first two rows (threadIdx.y=0, threadIdx.y=1) is handled within one warp. This means that for threadIdx.y=0 values D00..D15 are stored in B0..B15, then there is a padding of two doubles, after which within the same warp values D18..D33 are stored in B18..B31+B00..B01, which causes a bank conflict on B00-B01. Without the padding (ng=0) the first row is written to D00..D15 in B00..B15, the second row in D16..D31 in B16..B31, so no bank conflict occurs. 
 
For a thread block of blockDim.x>=32 the problem shouldn’t occur. For example, for itot=32, blockDim.x=32, ng=2, the first row is stored in banks B00..B31, then two cells padding, second row in B02..B31+B00..B01, etc.