File indexing completed on 2026-04-09 07:49:29
0001 #pragma once
0002
0003
0004
0005
0006
0007
0008
0009
0010
0011
0012
0013
0014
0015
0016
0017
0018
0019
0020
0021
0022
0023
0024
0025
0026
0027
0028
0029
0030
0031
0032
0033
0034
0035
0036
0037
0038
0039
0040
0041
0042
0043
0044
0045
0046
0047
0048
0049
0050
0051
0052
0053
0054
0055
0056
0057
0058
0059
0060
0061
0062
0063
0064
0065
0066
0067
0068
0069
0070
0071
0072
0073
0074
0075
0076
0077
0078
0079
0080
0081
0082
0083
0084
0085
0086
0087
0088
0089
0090
0091
0092
0093
0094
0095
0096
0097
0098
0099
0100
0101
0102 #if defined(__CUDACC__) || defined(__CUDABE__)
0103 # define SPMT_FUNCTION __host__ __device__ __forceinline__
0104 #else
0105 # define SPMT_FUNCTION inline
0106 #endif
0107
0108
0109 #if defined(__CUDACC__) || defined(__CUDABE__)
0110 #else
0111 #include <cassert>
0112 #include <string>
0113 #include <iostream>
0114 #include <sstream>
0115 #include <iomanip>
0116 #endif
0117
0118
0119
0120
0121
0122
0123
0124
0125
0126
0127
0128
0129
0130
0131
0132
0133 #define WITH_MPMT 1
0134
0135 namespace s_pmt
0136 {
0137 enum
0138 {
0139 NUM_CAT = 3,
0140 NUM_LAYR = 4,
0141 NUM_PROP = 2,
0142 NUM_CD_LPMT = 17612,
0143 NUM_SPMT = 25600,
0144 NUM_WP = 2400,
0145 NUM_WP_ATM_LPMT = 348,
0146 NUM_WP_ATM_MPMT_ALREADY = 600,
0147 #ifdef WITH_MPMT
0148 NUM_WP_ATM_MPMT = 600,
0149 #else
0150 NUM_WP_ATM_MPMT = 0,
0151 #endif
0152 NUM_WP_WAL_PMT = 5,
0153 OFFSET_CD_LPMT = 0,
0154 OFFSET_CD_LPMT_END = 17612,
0155 OFFSET_CD_SPMT = 20000,
0156 OFFSET_CD_SPMT_END = 45600,
0157 OFFSET_WP_PMT = 50000,
0158 OFFSET_WP_PMT_END = 52400,
0159 OFFSET_WP_ATM_LPMT = 52400,
0160 OFFSET_WP_ATM_LPMT_END = 52748,
0161 OFFSET_WP_ATM_MPMT = 53000,
0162 OFFSET_WP_ATM_MPMT_END = 53600,
0163 OFFSET_WP_WAL_PMT = 54000,
0164 OFFSET_WP_WAL_PMT_END = 54005
0165 };
0166
0167
0168 enum
0169 {
0170 NUM_CD_LPMT_AND_WP = NUM_CD_LPMT + NUM_WP + 0 + 0 + 0 + 0,
0171 NUM_LPMTIDX = NUM_CD_LPMT + NUM_WP + NUM_WP_ATM_LPMT + NUM_WP_ATM_MPMT + NUM_WP_WAL_PMT + 0,
0172 NUM_ALL = NUM_CD_LPMT + NUM_WP + NUM_WP_ATM_LPMT + NUM_WP_ATM_MPMT + NUM_WP_WAL_PMT + NUM_SPMT,
0173 NUM_ALL_EXCEPT_MPMT = NUM_ALL - NUM_WP_ATM_MPMT,
0174 NUM_CONTIGUOUSIDX = NUM_CD_LPMT + NUM_WP + NUM_WP_ATM_LPMT + NUM_WP_ATM_MPMT + NUM_WP_WAL_PMT + NUM_SPMT,
0175 NUM_OLDCONTIGUOUSIDX = NUM_CD_LPMT + NUM_WP + NUM_WP_ATM_LPMT + NUM_WP_ATM_MPMT_ALREADY + NUM_WP_WAL_PMT + NUM_SPMT,
0176 };
0177
0178
0179 #if defined(__CUDACC__) || defined(__CUDABE__)
0180 #else
0181 SPMT_FUNCTION std::string desc()
0182 {
0183 std::stringstream ss ;
0184 ss
0185 << "[s_pmt::desc\n"
0186 #ifdef WITH_MPMT
0187 << " WITH_MPMT\n"
0188 #else
0189 << " NOT:WITH_MPMT\n"
0190 #endif
0191 << std::setw(25) << "NUM_CAT" << std::setw(7) << NUM_CAT << "\n"
0192 << std::setw(25) << "NUM_LAYR" << std::setw(7) << NUM_LAYR << "\n"
0193 << std::setw(25) << "NUM_PROP" << std::setw(7) << NUM_PROP << "\n"
0194 << std::setw(25) << "NUM_CD_LPMT" << std::setw(7) << NUM_CD_LPMT << "\n"
0195 << std::setw(25) << "NUM_SPMT" << std::setw(7) << NUM_SPMT << "\n"
0196 << std::setw(25) << "NUM_WP" << std::setw(7) << NUM_WP << "\n"
0197 << std::setw(25) << "NUM_WP_ATM_LPMT" << std::setw(7) << NUM_WP_ATM_LPMT << "\n"
0198 << std::setw(25) << "NUM_WP_ATM_MPMT_ALREADY" << std::setw(7) << NUM_WP_ATM_MPMT_ALREADY << "\n"
0199 << std::setw(25) << "NUM_WP_ATM_MPMT" << std::setw(7) << NUM_WP_ATM_MPMT << "\n"
0200 << std::setw(25) << "NUM_WP_WAL_PMT" << std::setw(7) << NUM_WP_WAL_PMT << "\n"
0201 << std::setw(25) << "NUM_CD_LPMT_AND_WP" << std::setw(7) << NUM_CD_LPMT_AND_WP << "\n"
0202 << std::setw(25) << "NUM_ALL" << std::setw(7) << NUM_ALL << "\n"
0203 << std::setw(25) << "NUM_LPMTIDX" << std::setw(7) << NUM_LPMTIDX << "\n"
0204 << std::setw(25) << "NUM_CONTIGUOUSIDX" << std::setw(7) << NUM_CONTIGUOUSIDX << "\n"
0205 << std::setw(25) << "NUM_OLDCONTIGUOUSIDX" << std::setw(7) << NUM_OLDCONTIGUOUSIDX << "\n"
0206 << std::setw(25) << "OFFSET_CD_LPMT" << std::setw(7) << OFFSET_CD_LPMT << "\n"
0207 << std::setw(25) << "OFFSET_CD_LPMT_END" << std::setw(7) << OFFSET_CD_LPMT_END << "\n"
0208 << std::setw(25) << "OFFSET_CD_SPMT" << std::setw(7) << OFFSET_CD_SPMT << "\n"
0209 << std::setw(25) << "OFFSET_CD_SPMT_END" << std::setw(7) << OFFSET_CD_SPMT_END << "\n"
0210 << std::setw(25) << "OFFSET_WP_PMT" << std::setw(7) << OFFSET_WP_PMT << "\n"
0211 << std::setw(25) << "OFFSET_WP_PMT_END" << std::setw(7) << OFFSET_WP_PMT_END << "\n"
0212 << std::setw(25) << "OFFSET_WP_ATM_LPMT" << std::setw(7) << OFFSET_WP_ATM_LPMT << "\n"
0213 << std::setw(25) << "OFFSET_WP_ATM_LPMT_END" << std::setw(7) << OFFSET_WP_ATM_LPMT_END << "\n"
0214 << std::setw(25) << "OFFSET_WP_ATM_MPMT" << std::setw(7) << OFFSET_WP_ATM_MPMT << "\n"
0215 << std::setw(25) << "OFFSET_WP_ATM_MPMT_END" << std::setw(7) << OFFSET_WP_ATM_MPMT_END << "\n"
0216 << std::setw(25) << "OFFSET_WP_WAL_PMT" << std::setw(7) << OFFSET_WP_WAL_PMT << "\n"
0217 << std::setw(25) << "OFFSET_WP_WAL_PMT_END" << std::setw(7) << OFFSET_WP_WAL_PMT_END << "\n"
0218 << "]s_pmt::desc\n"
0219 ;
0220
0221 std::string str = ss.str() ;
0222 return str ;
0223 }
0224
0225 SPMT_FUNCTION void check_pmtid( int pmtid )
0226 {
0227 assert( pmtid >= 0 && pmtid < OFFSET_WP_WAL_PMT_END );
0228 assert(!(pmtid >= OFFSET_CD_LPMT_END && pmtid < OFFSET_CD_SPMT ));
0229 assert(!(pmtid >= OFFSET_CD_SPMT_END && pmtid < OFFSET_WP_PMT ));
0230
0231 assert( OFFSET_CD_LPMT_END - OFFSET_CD_LPMT == NUM_CD_LPMT );
0232 assert( OFFSET_CD_SPMT_END - OFFSET_CD_SPMT == NUM_SPMT );
0233 assert( OFFSET_WP_PMT_END - OFFSET_WP_PMT == NUM_WP );
0234
0235 #ifdef WITH_MPMT
0236 assert( OFFSET_WP_ATM_MPMT_END - OFFSET_WP_ATM_MPMT == NUM_WP_ATM_MPMT ) ;
0237 #endif
0238
0239 assert( OFFSET_WP_WAL_PMT_END - OFFSET_WP_WAL_PMT == NUM_WP_WAL_PMT ) ;
0240
0241 assert( NUM_LPMTIDX == NUM_CD_LPMT + NUM_WP + NUM_WP_ATM_LPMT + NUM_WP_ATM_MPMT + NUM_WP_WAL_PMT );
0242 assert( NUM_CONTIGUOUSIDX == NUM_CD_LPMT + NUM_WP + NUM_WP_ATM_LPMT + NUM_WP_ATM_MPMT + NUM_WP_WAL_PMT + NUM_SPMT ) ;
0243 assert( NUM_ALL == NUM_CD_LPMT + NUM_WP + NUM_WP_ATM_LPMT + NUM_WP_ATM_MPMT + NUM_WP_WAL_PMT + NUM_SPMT ) ;
0244
0245 assert( NUM_OLDCONTIGUOUSIDX == NUM_CD_LPMT + NUM_WP + NUM_WP_ATM_LPMT + NUM_WP_ATM_MPMT_ALREADY + NUM_WP_WAL_PMT + NUM_SPMT ) ;
0246
0247
0248
0249 }
0250
0251
0252
0253
0254 #endif
0255
0256
0257
0258
0259
0260 SPMT_FUNCTION bool in_range( int ix, int ix0, int ix1 ){ return ix >= ix0 && ix < ix1 ; }
0261
0262
0263 SPMT_FUNCTION bool id_CD_LPMT( int id ){ return in_range( id, OFFSET_CD_LPMT , OFFSET_CD_LPMT_END ) ; }
0264 SPMT_FUNCTION bool id_CD_SPMT( int id ){ return in_range( id, OFFSET_CD_SPMT , OFFSET_CD_SPMT_END ) ; }
0265 SPMT_FUNCTION bool id_WP_PMT( int id ){ return in_range( id, OFFSET_WP_PMT , OFFSET_WP_PMT_END ) ; }
0266 SPMT_FUNCTION bool id_WP_ATM_LPMT( int id ){ return in_range( id, OFFSET_WP_ATM_LPMT , OFFSET_WP_ATM_LPMT_END ) ; }
0267 SPMT_FUNCTION bool id_WP_ATM_MPMT( int id ){ return in_range( id, OFFSET_WP_ATM_MPMT , OFFSET_WP_ATM_MPMT_END ) ; }
0268 SPMT_FUNCTION bool id_WP_WAL_PMT( int id ){ return in_range( id, OFFSET_WP_WAL_PMT , OFFSET_WP_WAL_PMT_END ) ; }
0269
0270
0271
0272
0273
0274
0275
0276
0277
0278
0279
0280
0281
0282
0283
0284
0285
0286
0287
0288
0289
0290
0291
0292
0293
0294
0295
0296
0297
0298
0299
0300
0301
0302
0303
0304
0305
0306
0307
0308
0309
0310
0311
0312
0313
0314
0315
0316
0317
0318
0319
0320
0321
0322
0323
0324
0325
0326
0327
0328
0329
0330
0331
0332
0333
0334
0335
0336
0337
0338
0339
0340
0341
0342
0343
0344
0345
0346
0347
0348
0349
0350
0351
0352
0353
0354
0355
0356
0357
0358
0359
0360 SPMT_FUNCTION bool ix_CD_LPMT( int ix ){ return in_range(ix, 0 , NUM_CD_LPMT ) ; }
0361 SPMT_FUNCTION bool ix_CD_SPMT( int ix ){ return in_range(ix, NUM_CD_LPMT , NUM_CD_LPMT + NUM_SPMT) ; }
0362 SPMT_FUNCTION bool ix_WP_PMT( int ix ){ return in_range(ix, NUM_CD_LPMT + NUM_SPMT , NUM_CD_LPMT + NUM_SPMT + NUM_WP) ; }
0363 SPMT_FUNCTION bool ix_WP_ATM_LPMT( int ix ){ return in_range(ix, NUM_CD_LPMT + NUM_SPMT + NUM_WP , NUM_CD_LPMT + NUM_SPMT + NUM_WP + NUM_WP_ATM_LPMT) ; }
0364 SPMT_FUNCTION bool ix_WP_ATM_MPMT( int ix ){ return in_range(ix, NUM_CD_LPMT + NUM_SPMT + NUM_WP + NUM_WP_ATM_LPMT , NUM_CD_LPMT + NUM_SPMT + NUM_WP + NUM_WP_ATM_LPMT + NUM_WP_ATM_MPMT_ALREADY) ; }
0365 SPMT_FUNCTION bool ix_WP_WAL_PMT( int ix ){ return in_range(ix, NUM_CD_LPMT + NUM_SPMT + NUM_WP + NUM_WP_ATM_LPMT + NUM_WP_ATM_MPMT_ALREADY , NUM_CD_LPMT + NUM_SPMT + NUM_WP + NUM_WP_ATM_LPMT + NUM_WP_ATM_MPMT_ALREADY + NUM_WP_WAL_PMT ) ; }
0366
0367
0368
0369
0370
0371
0372
0373
0374
0375
0376
0377
0378
0379 SPMT_FUNCTION int oldcontiguousidx_from_pmtid( int id )
0380 {
0381 #if defined(__CUDACC__) || defined(__CUDABE__)
0382 #else
0383 check_pmtid(id);
0384 #endif
0385
0386 int ix = -1 ;
0387
0388 if( id_CD_LPMT(id) ) ix = id ;
0389 else if( id_CD_SPMT(id) ) ix = id - OFFSET_CD_SPMT + NUM_CD_LPMT ;
0390 else if( id_WP_PMT(id) ) ix = id - OFFSET_WP_PMT + NUM_CD_LPMT + NUM_SPMT ;
0391 else if( id_WP_ATM_LPMT(id)) ix = id - OFFSET_WP_ATM_LPMT + NUM_CD_LPMT + NUM_SPMT + NUM_WP ;
0392 else if( id_WP_ATM_MPMT(id)) ix = id - OFFSET_WP_ATM_MPMT + NUM_CD_LPMT + NUM_SPMT + NUM_WP + NUM_WP_ATM_LPMT ;
0393 else if( id_WP_WAL_PMT(id)) ix = id - OFFSET_WP_WAL_PMT + NUM_CD_LPMT + NUM_SPMT + NUM_WP + NUM_WP_ATM_LPMT + NUM_WP_ATM_MPMT_ALREADY ;
0394
0395 return ix ;
0396 }
0397
0398
0399
0400
0401
0402
0403
0404
0405
0406
0407
0408 SPMT_FUNCTION int pmtid_from_oldcontiguousidx( int ix )
0409 {
0410 #if defined(__CUDACC__) || defined(__CUDABE__)
0411 #else
0412 assert( ix >= 0 && ix < NUM_OLDCONTIGUOUSIDX );
0413 #endif
0414
0415 int id = -1 ;
0416
0417 if( ix_CD_LPMT(ix) ) id = OFFSET_CD_LPMT + ix ;
0418 else if( ix_CD_SPMT(ix) ) id = OFFSET_CD_SPMT + ix - ( NUM_CD_LPMT ) ;
0419 else if( ix_WP_PMT(ix) ) id = OFFSET_WP_PMT + ix - ( NUM_CD_LPMT + NUM_SPMT ) ;
0420 else if( ix_WP_ATM_LPMT(ix) ) id = OFFSET_WP_ATM_LPMT + ix - ( NUM_CD_LPMT + NUM_SPMT + NUM_WP ) ;
0421 else if( ix_WP_ATM_MPMT(ix) ) id = OFFSET_WP_ATM_MPMT + ix - ( NUM_CD_LPMT + NUM_SPMT + NUM_WP + NUM_WP_ATM_LPMT ) ;
0422 else if( ix_WP_WAL_PMT(ix) ) id = OFFSET_WP_WAL_PMT + ix - ( NUM_CD_LPMT + NUM_SPMT + NUM_WP + NUM_WP_ATM_LPMT + NUM_WP_ATM_MPMT_ALREADY ) ;
0423
0424 return id ;
0425 }
0426
0427
0428
0429
0430
0431
0432
0433
0434
0435
0436
0437
0438
0439
0440
0441
0442
0443
0444
0445
0446
0447
0448
0449
0450
0451
0452
0453
0454
0455
0456
0457
0458
0459
0460
0461
0462
0463
0464
0465
0466
0467
0468
0469
0470
0471
0472
0473
0474
0475
0476
0477
0478
0479
0480
0481
0482
0483
0484
0485
0486
0487
0488
0489
0490
0491
0492
0493
0494 SPMT_FUNCTION int contiguousidx_from_pmtid( int id )
0495 {
0496 #if defined(__CUDACC__) || defined(__CUDABE__)
0497 #else
0498 check_pmtid(id);
0499 #endif
0500
0501 int iy = -1 ;
0502
0503 if( id_CD_LPMT(id) ) iy = id ;
0504 else if( id_WP_PMT(id) ) iy = id - OFFSET_WP_PMT + NUM_CD_LPMT ;
0505 else if( id_WP_ATM_LPMT(id)) iy = id - OFFSET_WP_ATM_LPMT + NUM_CD_LPMT + NUM_WP ;
0506 else if( id_WP_ATM_MPMT(id)) iy = id - OFFSET_WP_ATM_MPMT + NUM_CD_LPMT + NUM_WP + NUM_WP_ATM_LPMT ;
0507 else if( id_WP_WAL_PMT(id)) iy = id - OFFSET_WP_WAL_PMT + NUM_CD_LPMT + NUM_WP + NUM_WP_ATM_LPMT + NUM_WP_ATM_MPMT ;
0508 else if( id_CD_SPMT(id) ) iy = id - OFFSET_CD_SPMT + NUM_CD_LPMT + NUM_WP + NUM_WP_ATM_LPMT + NUM_WP_ATM_MPMT + NUM_WP_WAL_PMT ;
0509
0510 return iy ;
0511 }
0512
0513
0514
0515
0516
0517
0518
0519
0520
0521
0522
0523
0524
0525
0526
0527
0528
0529
0530
0531
0532
0533
0534
0535
0536
0537
0538
0539
0540
0541
0542
0543
0544
0545
0546
0547
0548
0549
0550
0551
0552
0553
0554
0555 SPMT_FUNCTION int lpmtidx_from_pmtid( int id )
0556 {
0557 int iy = -1 ;
0558
0559 if( id_CD_LPMT(id) ) iy = id ;
0560 else if( id_WP_PMT(id) ) iy = id - OFFSET_WP_PMT + NUM_CD_LPMT ;
0561 else if( id_WP_ATM_LPMT(id)) iy = id - OFFSET_WP_ATM_LPMT + NUM_CD_LPMT + NUM_WP ;
0562 else if( id_WP_ATM_MPMT(id)) iy = id - OFFSET_WP_ATM_MPMT + NUM_CD_LPMT + NUM_WP + NUM_WP_ATM_LPMT ;
0563 else if( id_WP_WAL_PMT(id)) iy = id - OFFSET_WP_WAL_PMT + NUM_CD_LPMT + NUM_WP + NUM_WP_ATM_LPMT + NUM_WP_ATM_MPMT ;
0564
0565 return iy ;
0566 }
0567
0568
0569
0570
0571
0572
0573
0574
0575
0576
0577
0578
0579
0580
0581
0582
0583
0584 SPMT_FUNCTION bool iy_CD_LPMT( int iy ){ return in_range(iy, 0 , NUM_CD_LPMT ) ; }
0585 SPMT_FUNCTION bool iy_WP_PMT( int iy ){ return in_range(iy, NUM_CD_LPMT , NUM_CD_LPMT + NUM_WP ) ; }
0586 SPMT_FUNCTION bool iy_WP_ATM_LPMT( int iy ){ return in_range(iy, NUM_CD_LPMT + NUM_WP , NUM_CD_LPMT + NUM_WP + NUM_WP_ATM_LPMT) ; }
0587 SPMT_FUNCTION bool iy_WP_ATM_MPMT( int iy ){ return in_range(iy, NUM_CD_LPMT + NUM_WP + NUM_WP_ATM_LPMT , NUM_CD_LPMT + NUM_WP + NUM_WP_ATM_LPMT + NUM_WP_ATM_MPMT) ; }
0588 SPMT_FUNCTION bool iy_WP_WAL_PMT( int iy ){ return in_range(iy, NUM_CD_LPMT + NUM_WP + NUM_WP_ATM_LPMT + NUM_WP_ATM_MPMT , NUM_CD_LPMT + NUM_WP + NUM_WP_ATM_LPMT + NUM_WP_ATM_MPMT + NUM_WP_WAL_PMT ) ; }
0589 SPMT_FUNCTION bool iy_CD_SPMT( int iy ){ return in_range(iy, NUM_CD_LPMT + NUM_WP + NUM_WP_ATM_LPMT + NUM_WP_ATM_MPMT + NUM_WP_WAL_PMT , NUM_CD_LPMT + NUM_WP + NUM_WP_ATM_LPMT + NUM_WP_ATM_MPMT + NUM_WP_WAL_PMT + NUM_SPMT) ; }
0590
0591 SPMT_FUNCTION int pmtid_from_contiguousidx( int iy )
0592 {
0593 #if defined(__CUDACC__) || defined(__CUDABE__)
0594 #else
0595 assert( iy >= 0 && iy < NUM_CONTIGUOUSIDX );
0596 #endif
0597
0598 int id = -1 ;
0599
0600 if( iy_CD_LPMT(iy) ) id = OFFSET_CD_LPMT + iy ;
0601 else if( iy_WP_PMT(iy) ) id = OFFSET_WP_PMT + iy - ( NUM_CD_LPMT ) ;
0602 else if( iy_WP_ATM_LPMT(iy) ) id = OFFSET_WP_ATM_LPMT + iy - ( NUM_CD_LPMT + NUM_WP ) ;
0603 else if( iy_WP_ATM_MPMT(iy) ) id = OFFSET_WP_ATM_MPMT + iy - ( NUM_CD_LPMT + NUM_WP + NUM_WP_ATM_LPMT ) ;
0604 else if( iy_WP_WAL_PMT(iy) ) id = OFFSET_WP_WAL_PMT + iy - ( NUM_CD_LPMT + NUM_WP + NUM_WP_ATM_LPMT + NUM_WP_ATM_MPMT ) ;
0605 else if( iy_CD_SPMT(iy) ) id = OFFSET_CD_SPMT + iy - ( NUM_CD_LPMT + NUM_WP + NUM_WP_ATM_LPMT + NUM_WP_ATM_MPMT + NUM_WP_WAL_PMT ) ;
0606
0607 return id ;
0608 }
0609
0610
0611
0612
0613
0614
0615
0616
0617
0618
0619
0620
0621
0622
0623
0624
0625
0626
0627
0628
0629
0630
0631
0632
0633
0634
0635
0636
0637
0638
0639
0640
0641
0642
0643
0644
0645
0646
0647
0648
0649
0650
0651
0652
0653
0654
0655
0656
0657
0658
0659
0660
0661
0662
0663
0664
0665
0666 SPMT_FUNCTION int pmtid_from_lpmtidx( int iy )
0667 {
0668 #if defined(__CUDACC__) || defined(__CUDABE__)
0669 #else
0670 assert( iy >= 0 && iy < NUM_LPMTIDX );
0671 #endif
0672
0673 int id = -1 ;
0674
0675 if( iy_CD_LPMT(iy) ) id = OFFSET_CD_LPMT + iy ;
0676 else if( iy_WP_PMT(iy) ) id = OFFSET_WP_PMT + iy - ( NUM_CD_LPMT ) ;
0677 else if( iy_WP_ATM_LPMT(iy) ) id = OFFSET_WP_ATM_LPMT + iy - ( NUM_CD_LPMT + NUM_WP ) ;
0678 else if( iy_WP_ATM_MPMT(iy) ) id = OFFSET_WP_ATM_MPMT + iy - ( NUM_CD_LPMT + NUM_WP + NUM_WP_ATM_LPMT ) ;
0679 else if( iy_WP_WAL_PMT(iy) ) id = OFFSET_WP_WAL_PMT + iy - ( NUM_CD_LPMT + NUM_WP + NUM_WP_ATM_LPMT + NUM_WP_ATM_MPMT ) ;
0680
0681 return id ;
0682 }
0683
0684
0685
0686
0687
0688
0689
0690
0691
0692
0693
0694
0695
0696
0697
0698 SPMT_FUNCTION bool is_spmtid( int pmtid )
0699 {
0700 return pmtid >= OFFSET_CD_SPMT && pmtid < OFFSET_CD_SPMT_END ;
0701 }
0702 SPMT_FUNCTION bool is_spmtidx( int pmtidx )
0703 {
0704 return pmtidx >= 0 && pmtidx < NUM_SPMT ;
0705 }
0706
0707
0708 SPMT_FUNCTION int pmtid_from_spmtidx( int spmtidx )
0709 {
0710 #if defined(__CUDACC__) || defined(__CUDABE__)
0711 #else
0712 assert( spmtidx >= 0 && spmtidx < NUM_SPMT );
0713 #endif
0714 return spmtidx + OFFSET_CD_SPMT ;
0715 }
0716
0717 SPMT_FUNCTION int spmtidx_from_pmtid( int pmtid )
0718 {
0719 #if defined(__CUDACC__) || defined(__CUDABE__)
0720 #else
0721 assert( pmtid >= OFFSET_CD_SPMT && pmtid < OFFSET_CD_SPMT_END );
0722 #endif
0723 return pmtid - OFFSET_CD_SPMT ;
0724 }
0725
0726 }
0727