Skip to content

Instantly share code, notes, and snippets.

@yzhliu

yzhliu/slow.cu

Created December 10, 2019 00:08
Show Gist options
  • Save yzhliu/f1aa7c1f25cfd8788cda75a350c4b274 to your computer and use it in GitHub Desktop.
Save yzhliu/f1aa7c1f25cfd8788cda75a350c4b274 to your computer and use it in GitHub Desktop.
Download ZIP
Raw
slow.cu
extern "C" __global__ void tvmop_kernel0( float* __restrict__ buffer, float* __restrict__ buffer1, float* __restrict__ buffer2, int tindex, int tindex1, int tindex2, int stride, int stride1, int stride2, int stride3, int stride4, int stride5, int stride6, int stride7, int stride8) {
if (((int)blockIdx.x) < (((tindex * tindex1) * tindex2) >> 6)) {
if (((((tindex1 >= 0) && ((((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) % tindex1) >= 0)) || ((tindex1 < 0) && ((((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) % tindex1) <= 0))) ? (((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) / tindex1) : ((((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) / tindex1) - 1)) < tindex) {
if (((((tindex1 >= 0) && ((((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) % tindex1) >= 0)) || ((tindex1 < 0) && ((((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) % tindex1) <= 0))) ? (((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) % tindex1) : ((((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) % tindex1) + tindex1)) < tindex1) {
if (((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) + tindex2)) < tindex2) {
if (((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) < (tindex * tindex1)) {
if (0 <= ((((tindex1 >= 0) && ((((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) % tindex1) >= 0)) || ((tindex1 < 0) && ((((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) % tindex1) <= 0))) ? (((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) / tindex1) : ((((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) / tindex1) - 1))) {
if (((((tindex1 >= 0) && ((((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) % tindex1) >= 0)) || ((tindex1 < 0) && ((((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) % tindex1) <= 0))) ? (((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) / tindex1) : ((((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) / tindex1) - 1)) < tindex) {
if (0 <= ((((tindex1 >= 0) && ((((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) % tindex1) >= 0)) || ((tindex1 < 0) && ((((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) % tindex1) <= 0))) ? (((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) % tindex1) : ((((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) % tindex1) + tindex1))) {
if (((((tindex1 >= 0) && ((((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) % tindex1) >= 0)) || ((tindex1 < 0) && ((((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) % tindex1) <= 0))) ? (((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) % tindex1) : ((((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) % tindex1) + tindex1)) < tindex1) {
if (0 <= ((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) + tindex2))) {
if (((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) + tindex2)) < tindex2) {
buffer[(((((((tindex1 >= 0) && ((((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) % tindex1) >= 0)) || ((tindex1 < 0) && ((((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) % tindex1) <= 0))) ? (((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) / tindex1) : ((((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) / tindex1) - 1)) * stride6) + (((((tindex1 >= 0) && ((((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) % tindex1) >= 0)) || ((tindex1 < 0) && ((((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) % tindex1) <= 0))) ? (((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) % tindex1) : ((((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) % tindex1) + tindex1)) * stride7)) + (((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) + tindex2)) * stride8))] = (buffer1[(((((((tindex1 >= 0) && ((((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) % tindex1) >= 0)) || ((tindex1 < 0) && ((((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) % tindex1) <= 0))) ? (((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) / tindex1) : ((((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) / tindex1) - 1)) * stride) + (((((tindex1 >= 0) && ((((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) % tindex1) >= 0)) || ((tindex1 < 0) && ((((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) % tindex1) <= 0))) ? (((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) % tindex1) : ((((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) % tindex1) + tindex1)) * stride1)) + (((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) + tindex2)) * stride2))] + buffer2[(((((((tindex1 >= 0) && ((((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) % tindex1) >= 0)) || ((tindex1 < 0) && ((((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) % tindex1) <= 0))) ? (((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) / tindex1) : ((((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) / tindex1) - 1)) * stride3) + (((((tindex1 >= 0) && ((((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) % tindex1) >= 0)) || ((tindex1 < 0) && ((((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) % tindex1) <= 0))) ? (((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) % tindex1) : ((((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) % tindex1) + tindex1)) * stride4)) + (((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) + tindex2)) * stride5))]);
}
}
}
}
}
}
}
}
}
}
} else {
if (((((tindex1 >= 0) && ((((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) % tindex1) >= 0)) || ((tindex1 < 0) && ((((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) % tindex1) <= 0))) ? (((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) / tindex1) : ((((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) / tindex1) - 1)) < tindex) {
if (((((tindex1 >= 0) && ((((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) % tindex1) >= 0)) || ((tindex1 < 0) && ((((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) % tindex1) <= 0))) ? (((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) % tindex1) : ((((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) % tindex1) + tindex1)) < tindex1) {
if (((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) + tindex2)) < tindex2) {
if (((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) < (tindex * tindex1)) {
if (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) < ((tindex * tindex1) * tindex2)) {
if (0 <= ((((tindex1 >= 0) && ((((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) % tindex1) >= 0)) || ((tindex1 < 0) && ((((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) % tindex1) <= 0))) ? (((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) / tindex1) : ((((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) / tindex1) - 1))) {
if (((((tindex1 >= 0) && ((((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) % tindex1) >= 0)) || ((tindex1 < 0) && ((((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) % tindex1) <= 0))) ? (((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) / tindex1) : ((((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) / tindex1) - 1)) < tindex) {
if (0 <= ((((tindex1 >= 0) && ((((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) % tindex1) >= 0)) || ((tindex1 < 0) && ((((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) % tindex1) <= 0))) ? (((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) % tindex1) : ((((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) % tindex1) + tindex1))) {
if (((((tindex1 >= 0) && ((((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) % tindex1) >= 0)) || ((tindex1 < 0) && ((((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) % tindex1) <= 0))) ? (((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) % tindex1) : ((((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) % tindex1) + tindex1)) < tindex1) {
if (0 <= ((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) + tindex2))) {
if (((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) + tindex2)) < tindex2) {
buffer[(((((((tindex1 >= 0) && ((((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) % tindex1) >= 0)) || ((tindex1 < 0) && ((((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) % tindex1) <= 0))) ? (((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) / tindex1) : ((((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) / tindex1) - 1)) * stride6) + (((((tindex1 >= 0) && ((((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) % tindex1) >= 0)) || ((tindex1 < 0) && ((((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) % tindex1) <= 0))) ? (((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) % tindex1) : ((((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) % tindex1) + tindex1)) * stride7)) + (((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) + tindex2)) * stride8))] = (buffer1[(((((((tindex1 >= 0) && ((((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) % tindex1) >= 0)) || ((tindex1 < 0) && ((((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) % tindex1) <= 0))) ? (((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) / tindex1) : ((((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) / tindex1) - 1)) * stride) + (((((tindex1 >= 0) && ((((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) % tindex1) >= 0)) || ((tindex1 < 0) && ((((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) % tindex1) <= 0))) ? (((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) % tindex1) : ((((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) % tindex1) + tindex1)) * stride1)) + (((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) + tindex2)) * stride2))] + buffer2[(((((((tindex1 >= 0) && ((((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) % tindex1) >= 0)) || ((tindex1 < 0) && ((((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) % tindex1) <= 0))) ? (((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) / tindex1) : ((((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) / tindex1) - 1)) * stride3) + (((((tindex1 >= 0) && ((((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) % tindex1) >= 0)) || ((tindex1 < 0) && ((((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) % tindex1) <= 0))) ? (((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) % tindex1) : ((((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) / tindex2) - 1)) % tindex1) + tindex1)) * stride4)) + (((((tindex2 >= 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) >= 0)) || ((tindex2 < 0) && ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) <= 0))) ? (((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) : ((((((int)blockIdx.x) * 64) + ((int)threadIdx.x)) % tindex2) + tindex2)) * stride5))]);
}
}
}
}
}
}
}
}
}
}
}
}
}
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment