Scheduling / Deadlock ?

Hello,

I’m currently assessing the possibility of using weird CUDA deadlocking behaviors in order to generate random numbers…

I’m using this kernel in order to generate deadlocks, on my GTX260… with a block size of 512 threads:

[codebox]

extern shared unsigned int sh_arr_a;

global void

kernel_dl( void* data,

				unsigned int dataSize)

{

const unsigned int tid = threadIdx.x;

const unsigned int bid = blockIdx.x;

const unsigned int bdim = blockDim.x;

const unsigned int ind = bid * bdim + tid;



if( ind > dataSize-1 )

	return;

unsigned int * arr = (unsigned int *)data;

arr[ind]=0;

(sh_arr_a)[ tid ] = 0;

for( int i = 0; i < 1; ++i )

{

	sh_arr_a[ tid ] = 	(sh_arr_a[ (tid+96)%512 ]+

						sh_arr_a[ (tid+64)%512 ]+

						sh_arr_a[ (tid+32)%512 ]);

			

	//__syncthreads();

}

__syncthreads();

arr[ind] = sh_arr_a[ tid ];

}[/codebox]

The output is :

[codebox]0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

1056964608

1056964608

1056964608

1056964608

1056964608

1056964608

1056964608

1056964608

1065353216

1065353216

1065353216

1065353216

1065353216

1065353216

1065353216

1065353216

1065353216

1065353216

1065353216

1065353216

1065353216

1065353216

1065353216

1065353216

1065353216

1065353216

1065353216

1065353216

1065353216

1065353216

1065353216

1065353216

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

999985384

1007740890

1141538816

1135050752

1111097344

1111097344

1111097344

1111097344

1022330940

1022330940

1025748521

1025748521

1025748521

1025748521

1025748521

1025748521

1053923125

1053923125

1131347968

1131347968

1132138424

1132138424

1132806144

1132806144

0

1154613248

0

1154613248

1154613248

0

0

1154613248

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

0

[/codebox]

As you can see, there are some weird values appearing here and there…

After 512 iterations… We obtain this:

[codebox]66011136

4196130816

2970542080

3135389696

2809077760

1940889600

73646080

941834240

3474243904

613957952

1658011968

223330624

1364082688

2798764032

3894689792

2460008448

3454509056

735035392

1111572480

3831046144

4276433260

1988473196

2365010284

358003052

478230784

3905001728

210868480

1079064832

1655570432

1894776832

3044540416

2805334016

3461005312

129785856

634413056

3965632512

1516290048

285556736

2475016192

3705749504

3747464896

3922446016

3602040512

3427059392

2219704320

2394685440

2326298624

2151317504

1398808576

20496384

1943355392

3321667584

1810238538

591957066

2514816074

3733097546

3725992832

3098100608

4089004928

421929856

1602879488

3937402880

2505572352

171048960

3818438656

145399808

3029311488

2407383040

4236763136

1931567104

3657252864

1667481600

4244882624

1923105984

2639040

2324415680

1669775360

3642966016

2141405184

168214528

1782218752

1948319744

2948988928

2782887936

1527507955

2191911923

3192581107

2528177139

2175619136

753307712

1395036224

2817347648

618921984

2017460224

333447168

3229876224

2075197440

2139455488

2145353728

2081095680

1191788544

1204502528

3116433408

3103719424

2528694720

985256384

3696677312

945148352

4088692736

2545254400

1884487680

3427926016

2716860416

614727680

2402942976

210108416

2922903097

3104306745

597554745

416151097

1773992128

354056384

119175360

1539111104

884473856

3439329280

2315255808

4055367680

1862336512

2139897856

4049682432

3772121088

382705664

1444454400

4291264512

3229515776

1268854848

1014083648

4197691456

157495360

756203520

501432320

2595373056

2850144256

1194696704

1752416256

886390784

328671232

1654794782

3885672990

3019647518

788769310

4071348864

1990736512

2574793344

360438400

99483648

2775580672

4151836672

1475739648

2779234304

2377080832

1693409280

2095562752

2073051136

786743296

2701115392

3987423232

3042325568

2796172352

338048064

584201280

1075445760

829292544

996933632

1243086848

2929819648

51183616

2290679808

874348544

249182745

1619540505

3859036697

2488678937

2624069824

3809943744

1104617664

4213711040

2742026240

2344615936

1822425088

2219835392

2196013056

1486225408

2774335488

3484123136

2992988160

612065280

2228158464

314114048

2536986048

3882669504

282842560

3232126400

1774010368

3119693824

2770288640

1424605184

3575980032

4183801856

875151360

267329536

101732004

2533019300

3519336100

1088048804

1308668672

2886939392

1411592960

4128289536

2826436608

2368733184

4166516736

329252864

1182056448

2571558912

3876904960

2487402496

3906600960

45645824

3656957952

3222945792

4279477824

2925176384

3576800832

636134976

2201714688

847413248

3123118080

182452224

3278831616

1812398080

2015526912

3481960448

1303880033

4028080481

4231209313

1507008865

138610368

3883787968

58582720

608372416

3120955392

2357198848

2457862144

3221618688

3696132096

2614263808

3632103424

419004416

3763396608

129032192

1110966272

450363392

520828992

283293760

773371968

1010907200

1394688000

1157152768

3693461504

3930996736

369975296

2644918272

3694968832

1420025856

3138538004

3648375316

403458580

4188588564

1176790784

1334183680

3929409280

3772016384

1089601536

1913651200

3787980800

2963931136

1971650560

3406807040

3963535360

2528378880

743890944

1717952512

2363785216

1389723648

3748266368

3478356352

2618655104

2888565120

4195368960

3925458944

467189760

737099776

539238400

625459200

3740975104

3654754304

245717888

590601088

3706116992

3361233792

1985060864

2757648384

4008599552

3236012032

2666266624

3417309184

3843555328

3092512768

3414441984

1795219456

1768022016

3387244544

2094522368

1368940544

4018913280

449527808

2028177344

2448558016

1875183552

1454802880

1928249344

2348630016

1705230336

1284849664

2907791360

2769174528

751616000

890232832

2720039511

2165572183

148013655

702480983

1676301120

4063097664

1457386304

3365557056

1938948096

1611661312

2831679488

3158966272

1965457408

1826816000

2164260864

2302902272

1287528448

2636652544

3356426240

2007302144

1137832320

1424126336

119697792

4128371072

2792669184

3078963200

2773827584

2487533568

2276884480

1915551744

1462239232

1823571968

964813707

3814450059

3361137547

511501195

3899773504

2094043712

3406860864

917623360

2946236416

2495610880

2489319424

2939944960

2614181888

444784640

1927405568

4096802816

48734208

3109822464

1472577536

2706456576

130945472

2307625408

635474368

2753761728

1995128832

4171808768

3569762304

1393082368

3002040320

454950912

608796672

3155886080

2355738424

757315384

911161144

2509584184

3699046912

1691474432

521263616

2528836096

2735079424

4173332480

3328180224

1889927168

112066560

569966592

369819648

4206886912

3155640320

2182823936

2685665280

3658481664

382582976

1919074496

2899689664

1363198144

3993796608

1235320832

2454454272

917962752

3452895232

667811840

2154299392

644415488

3789182309

1238783333

2725270885

980702565

1029636544

348488128

373653952

1054802368

1896349696

2139095040

2873098240

2630352896

3725557760

2124562432

3421782016

727810048

2955657216

1904918528

512000000

1562738688

3820273280

713408128

1061797504

4168662656

3751641088

644775936

4205084672

3016982528

285548544

3657547776

1190576128

2113544192

208679111

811774151

2639769799

2036674759

911624000

4053501760

1990952768

3144042304

2652110848

967835648

2102919168

3787194368

198541312

2093613056

4111794176

2216722432

2925256704

3973439488

2258305024

1210122240

893531776

471971456

96712320

518272640

167854080

4041261056

2279063552

2700623872

4022165504

2902016000

2430615552

3550765056

2711941905

2526311185

2054910737

2240541457

1002367680

1345366720

1986046656

1643047616

3688628224

3725590528

847773696

810811392

[/codebox]

And if we normalize these numbers to [0,1], and we plot them on a graph, we obtain this:

So I was actually thinking I could use this method as a basis for a PRNG…

My question is: Is this behavior normal ? Will it present on every CUDA-capable cards ? What’s actually happening ?

Thank you.

Why does this code generate deadlock??

Sorry, I meant race condition…