| 496 | // Writes with boundary checking |
| 497 | |
| 498 | __device__ static void partialWriteOut128Bytes(uchar *out, const uint &index, |
| 499 | const uint &r1, const uint &r2, |
| 500 | const uint &r3, const uint &r4, |
| 501 | const uint &elements) { |
| 502 | if (index < elements) { out[index] = r1; } |
| 503 | if (index + blockDim.x < elements) { out[index + blockDim.x] = r1 >> 8; } |
| 504 | if (index + 2 * blockDim.x < elements) { |
| 505 | out[index + 2 * blockDim.x] = r1 >> 16; |
| 506 | } |
| 507 | if (index + 3 * blockDim.x < elements) { |
| 508 | out[index + 3 * blockDim.x] = r1 >> 24; |
| 509 | } |
| 510 | if (index + 4 * blockDim.x < elements) { out[index + 4 * blockDim.x] = r2; } |
| 511 | if (index + 5 * blockDim.x < elements) { |
| 512 | out[index + 5 * blockDim.x] = r2 >> 8; |
| 513 | } |
| 514 | if (index + 6 * blockDim.x < elements) { |
| 515 | out[index + 6 * blockDim.x] = r2 >> 16; |
| 516 | } |
| 517 | if (index + 7 * blockDim.x < elements) { |
| 518 | out[index + 7 * blockDim.x] = r2 >> 24; |
| 519 | } |
| 520 | if (index + 8 * blockDim.x < elements) { out[index + 8 * blockDim.x] = r3; } |
| 521 | if (index + 9 * blockDim.x < elements) { |
| 522 | out[index + 9 * blockDim.x] = r3 >> 8; |
| 523 | } |
| 524 | if (index + 10 * blockDim.x < elements) { |
| 525 | out[index + 10 * blockDim.x] = r3 >> 16; |
| 526 | } |
| 527 | if (index + 11 * blockDim.x < elements) { |
| 528 | out[index + 11 * blockDim.x] = r3 >> 24; |
| 529 | } |
| 530 | if (index + 12 * blockDim.x < elements) { |
| 531 | out[index + 12 * blockDim.x] = r4; |
| 532 | } |
| 533 | if (index + 13 * blockDim.x < elements) { |
| 534 | out[index + 13 * blockDim.x] = r4 >> 8; |
| 535 | } |
| 536 | if (index + 14 * blockDim.x < elements) { |
| 537 | out[index + 14 * blockDim.x] = r4 >> 16; |
| 538 | } |
| 539 | if (index + 15 * blockDim.x < elements) { |
| 540 | out[index + 15 * blockDim.x] = r4 >> 24; |
| 541 | } |
| 542 | } |
| 543 | |
| 544 | __device__ static void partialWriteOut128Bytes(schar *out, const uint &index, |
| 545 | const uint &r1, const uint &r2, |
no test coverage detected