我正在写关于在我的 Nvidia GTX 560Ti 卡上使用 CUDA 的atomicAdd()时发生的明显损坏。在开发一些代码期间,我遇到了 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) 及其位置。然而,在下面的示例输出中,除了 1024 的 255 个实例之外,它反而打印出一个 1023 和一个 1。在其他运行中,结果是不同的。即使 srand( 0 ) 替换了按时间播种的 RNG,每次运行的结果也不同。我在 GTX 560Ti 和 Tesla C2070 上都试过这个。特斯拉没有表现出任何腐败。我无法访问另一个 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?可以想象我的卡有问题吗?atomicAdd 真的有可能在 560Ti 上不安全吗?
在此先感谢您的帮助,克里斯
编辑:我的卡一定是坏的。当我用常规添加操作替换 atomicAdd 时,此测试也失败了。(是的,这些值不再是恒定的,因为操作不是原子的,竞争条件等 - 然而,在它们应该为零的地方有非零值,不应该在这些内存位置执行任何操作。)它在重新启动时也仍然存在,我正在进入仅运行登录屏幕的重新启动系统(所以可能是 X,但没有 OpenGL?)。系统正在运行 Ubuntu 10.04 和 CUDA 4.0。GPU 是 GeForce GTX 560 Ti。 有谁知道这是否是常见的故障模式?