我正在写一篇关于使用CUDA的atomicAdd()在我的Nvidia GTX 560Ti卡上出现明显的内存损坏问题的文章。在开发代码时，我遇到了使用atomicAdd时出现的问题，它似乎会破坏内存。我设计了一个测试来确定这是否确实是情况，并且行为是否可以在我的应用程序条件之外被复制。我编写了一个测试程序，使用atomicAdd增加缓冲区中较少的位置。在我的560Ti上，测试显示atomicAdd 会破坏随机位于内存中的比特。特别地，少量（看起来）随机放置的未被故意访问或修改的位置上的位被翻转。内核代码很简单，只有一个atomicAdd。测试代码如下：
#include <stdio.h>
#include <stdlib.h>
#include <sys/time.h>
#define ANSI_RED "\e[0;41m\e[41;37m"
#define ANSI_BLACK "\e[0;30m"
__global__ void kernel( unsigned int *a, unsigned int *map, int M, int N )
{
// Add to buffer.
atomicAdd( a + map[ blockIdx.x * N + threadIdx.x ], 1 );
}
template < class T > void swap( T &a, T &b ) { T t; t = a; a = b; b = t; }
int main( void )
{
// Chooses 560Ti on my machine
cudaSetDevice( 1 );
srand( time( 0 ) );
unsigned int M = 1024, N = 256;
unsigned int L = M * N, K = N;
unsigned int *dev_buf, *dev_map;
unsigned int *buf = new unsigned int[ L ];
unsigned int *map = new unsigned int[ L ];
unsigned int *indices = new unsigned int[ K ];
bool *check = new bool[ L ];
// Use buffer to indicate which spots in buffer should have valid values.
for( int l = 0; l < L; l++ ) check[ l ] = false;
// Generate K random indices into an L-sized buffer, init "check"
for( int k = 0; k < K; k++ )
{
int i = rand( ) % L;
while( check[ i ] )
i = rand( ) % L;
indices[ k ] = i;
check[ i ] = true;
}
// Generate a random M (blocks) x N (threads) array "map" of indices that contains
// offsets into "buf" such that there are at most K locations in "buf" that
// should be written to.
for( int m = 0; m < M; m++ )
for( int n = 0; n < N; n++ ) // Init.
map[ m * N + n ] = indices[ n ];
for( int i = 0; i < L; i++ ) // Shuffle.
swap( map[ i ], map[ i + rand( ) % ( L - i ) ] );
// Allocate and initialize device memory.
cudaMalloc( &dev_buf, L * sizeof( unsigned int ) );
cudaMalloc( &dev_map, N * M * sizeof( unsigned int ) );
cudaMemset( dev_buf, 0, L * sizeof( unsigned int ) );
cudaMemcpy( dev_map, map, L * sizeof( unsigned int ), cudaMemcpyHostToDevice );
kernel<<< M, N >>>( dev_buf, dev_map, M, N );
// Copy back to host.
cudaMemcpy( buf, dev_buf, L * sizeof( unsigned int ), cudaMemcpyDeviceToHost );
// Print non-zero values. Highlight abnormalities.
int j = 0;
for( int i = 0; i < L; i++ )
{
if( buf[ i ] != 0 )
{
if( ( buf[ i ] == M ) || ( buf[ i ] == 2 * M ) )
printf( "%d @ %d [%s]\t",
buf[ i ], i, check[ i ] ? "true" : "false" );
else
printf( ANSI_RED "%d @ %d [%s]\t" ANSI_BLACK,
buf[ i ], i, check[ i ] ? "true" : "false" );
j++;
}
}
printf( "\nj = %d\n", j );
}
编译环境：
nvcc test_atomicadd_bug.cu -o test_atomicadd_bug -arch sm_21
所有内核调用应该做的事情（总体上）就是将所有K个位置增加M次，从而在每个位置上产生一个K * M = 1024的结果。因此，在运行代码时，它应该打印出非零值（1024）及其位置。然而，在下面的示例输出中，它却打印出了一个1023和一个1，以及255个1024的实例。在其他运行中，结果也是不同的。即使srand(0)替换了时间种子的RNG，结果也会因每次运行而异。我已经在GTX 560Ti和Tesla C2070上尝试过这一点。Tesla没有出现任何损坏。我无法访问另一个560Ti。
1024 @ 1228 [true] 1024 @ 1271 [true] 1024 @ 1842 [true] 1024 @ 2480 [true] 1024 @ 3012 [true]
1024 @ 3802 [true] 1024 @ 4649 [true] 1024 @ 5636 [true] 1024 @ 6988 [true] 1024 @ 9400 [true]
1024 @ 10912 [true] 1024 @ 10930 [true] 1024 @ 11550 [true] 1024 @ 11888 [true] 1024 @ 12047 [true]
1024 @ 12837 [true] 1024 @ 12868 [true] 1024 @ 12991 [true] 1024 @ 16294 [true] 1024 @ 16690 [true]
1024 @ 17396 [true] 1024 @ 17529 [true] 1024 @ 19857 [true] 1024 @ 20926 [true] 1024 @ 22189 [true]
1024 @ 22391 [true] 1024 @ 22613 [true] 1024 @ 22851 [true] 1024 @ 23562 [true] 1024 @ 23955 [true]
1024 @ 24598 [true] 1024 @ 26058 [true] 1024 @ 26441 [true] 1024 @ 26962 [true] 1024 @ 27141 [true]
1024 @ 28101 [true] 1024 @ 28332 [true] 1024 @ 29485 [true] 1024 @ 29487 [true] 1024 @ 29942 [true]
1024 @ 31213 [true] 1024 @ 31965 [true] 1024 @ 35774 [true] 1024 @ 39342 [true] 1024 @ 39883 [true]
1024 @ 39960 [true] 1024 @ 40252 [true] 1024 @ 41435 [true] 1024 @ 42975 [true] 1024 @ 43336 [true]
1024 @ 44527 [true] 1024 @ 44657 [true] 1 @ 45494 [false] 1024 @ 46940 [true] 1024 @ 46983 [true]
1024 @ 48660 [true] 1024 @ 49034 [true] 1024 @ 49420 [true] 1024 @ 49620 [true] 1024 @ 50813 [true]
1024 @ 53859 [true] 1024 @ 55527 [true] 1024 @ 56677 [true] 1024 @ 57048 [true] 1024 @ 57759 [true]
1024 @ 58505 [true] 1024 @ 59539 [true] 1024 @ 59856 [true] 1024 @ 60341 [true] 1024 @ 61556 [true]
1024 @ 61733 [true] 1023 @ 61878 [true] 1024 @ 62025 [true] 1024 @ 65333 [true] 1024 @ 66131 [true]
1024 @ 67196 [true] 1024 @ 69428 [true] 1024 @ 70555 [true] 1024 @ 73135 [true] 1024 @ 73696 [true]
1024 @ 76797 [true] 1024 @ 76947 [true] 1024 @ 79166 [true] 1024 @ 79301 [true] 1024 @ 80182 [true]
1024 @ 80348 [true] 1024 @ 80574 [true] 1024 @ 81386 [true] 1024 @ 84416 [true] 1024 @ 86472 [true]
1024 @ 88234 [true] 1024 @ 88622 [true] 1024 @ 89355 [true] 1024 @ 89571 [true] 1024 @ 90716 [true]
1024 @ 91386 [true] 1024 @ 94846 [true] 1024 @ 95779 [true] 1024 @ 99146 [true] 1024 @ 99569 [true]
1024 @ 100202 [true] 1024 @ 102972 [true] 1024 @ 103909 [true] 1024 @ 104373 [true] 1024 @ 107707 [true]
1024 @ 108543 [true] 1024 @ 108617 [true] 1024 @ 109212 [true] 1024 @ 109388 [true] 1024 @ 111836 [true]
1024 @ 113078 [true] 1024 @ 113343 [true] 1024 @ 114451 [true] 1024 @ 114849 [true] 1024 @ 115024 [true]
1024 @ 115338 [true] 1024 @ 116675 [true] 1024 @ 118624 [true] 1024 @ 119884 [true] 1024 @ 120807 [true]
1024 @ 121993 [true] 1024 @ 122050 [true] 1024 @ 124643 [true] 1024 @ 125161 [true] 1024 @ 125843 [true]
1024 @ 126890 [true] 1024 @ 127718 [true] 1024 @ 127810 [true] 1024 @ 129646 [true] 1024 @ 129907 [true]
1024 @ 132288 [true] 1024 @ 132706 [true] 1024 @ 135574 [true] 1024 @ 136913 [true] 1024 @ 137346 [true]
1024 @ 138326 [true] 1024 @ 138685 [true] 1024 @ 138939 [true] 1024 @ 140996 [true] 1024 @ 141304 [true]
1024 @ 143902 [true] 1024 @ 145723 [true] 1024 @ 146149 [true] 1024 @ 149696 [true] 1024 @ 149726 [true]
1024 @ 150294 [true] 1024 @ 152057 [true] 1024 @ 152198 [true] 1024 @ 152239 [true] 1024 @ 153002 [true]
1024 @ 153776 [true] 1024 @ 156081 [true] 1024 @ 156377 [true] 1024 @ 156654 [true] 1024 @ 158008 [true]
1024 @ 158677 [true] 1024 @ 159369 [true] 1024 @ 159996 [true] 1024 @ 160060 [true] 1024 @ 161456 [true]
1024 @ 161732 [true] 1024 @ 163269 [true] 1024 @ 163675 [true] 1024 @ 163684 [true] 1024 @ 164397 [true]
1024 @ 165077 [true] 1024 @ 166036 [true] 1024 @ 168301 [true] 1024 @ 168409 [true] 1024 @ 171499 [true]
1024 @ 171772 [true] 1024 @ 173353 [true] 1024 @ 175290 [true] 1024 @ 175573 [true] 1024 @ 177155 [true]
1024 @ 178142 [true] 1024 @ 178718 [true] 1024 @ 178822 [true] 1024 @ 179161 [true] 1024 @ 179654 [true]
1024 @ 180683 [true] 1024 @ 182432 [true] 1024 @ 183086 [true] 1024 @ 183695 [true] 1024 @ 184730 [true]
1024 @ 186884 [true] 1024 @ 187746 [true] 1024 @ 188603 [true] 1024 @ 188948 [true] 1024 @ 189124 [true]
1024 @ 190268 [true] 1024 @ 191208 [true] 1024 @ 192630 [true] 1024 @ 193617 [true] 1024 @ 195426 [true]
1024 @ 198352 [true] 1024 @ 201345 [true] 1024 @ 201416 [true] 1024 @ 203214 [true] 1024 @ 205418 [true]
1024 @ 207467 [true] 1024 @ 208763 [true] 1024 @ 208924 [true] 1024 @ 209269 [true] 1024 @ 210679 [true]
1024 @ 211622 [true] 1024 @ 212029 [true] 1024 @ 212135 [true] 1024 @ 213228 [true] 1024 @ 216151 [true]
1024 @ 216425 [true] 1024 @ 216432 [true] 1024 @ 218039 [true] 1024 @ 219445 [true] 1024 @ 219675 [true]
1024 @ 220504 [true] 1024 @ 220702 [true] 1024 @ 220716 [true] 1024 @ 222687 [true] 1024 @ 223582 [true]
1024 @ 223758 [true] 1024 @ 223917 [true] 1024 @ 224254 [true] 1024 @ 224825 [true] 1024 @ 224845 [true]
1024 @ 225372 [true] 1024 @ 226297 [true] 1024 @ 228158 [true] 1024 @ 228367 [true] 1024 @ 229494 [true]
1024 @ 229636 [true] 1024 @ 230722 [true] 1024 @ 232001 [true] 1024 @ 232693 [true] 1024 @ 234729 [true]
1024 @ 235132 [true] 1024 @ 242699 [true] 1024 @ 245103 [true] 1024 @ 245948 [true] 1024 @ 246903 [true]
1024 @ 247836 [true] 1024 @ 247871 [true] 1024 @ 248694 [true] 1024 @ 248801 [true] 1024 @ 250204 [true]
1024 @ 250899 [true] 1024 @ 250968 [true] 1024 @ 251738 [true] 1024 @ 251930 [true] 1024 @ 256221 [true]
1024 @ 258244 [true] 1024 @ 258908 [true] 1024 @ 259884 [true] 1024 @ 260318 [true] 1024 @ 260424 [true]
1024 @ 260884 [true] 1024 @ 260953 [true]
j = 257
我的问题是：我使用atomicAdd的方式有问题吗？这在其他Nvidia GPU上会发生吗？其他560 Ti会出现这种情况吗？我的显卡可能存在故障吗？560Ti上的atomicAdd可能真的不安全吗？
提前感谢任何帮助， Chris
编辑：我的显卡可能有问题。当我用普通的加法操作替换atomicAdd时，测试也失败了。（是的，由于操作不是原子性的，会发生竞争条件等等——尽管如此，在一些内存位置上仍然有非零值，那里本不应该执行任何操作）。即使重启后仍然存在，我通过ssh连接进入重新启动的系统，只运行登录界面（所以可能是X，但没有OpenGL？），系统正在运行Ubuntu 10.04和CUDA 4.0。GPU是GeForce GTX 560 Ti。有人知道这是否是一个常见的故障模式吗？