compiler: Cast only product transfer bounds to long, update 01_gpu notebook

Applying as_long unconditionally also rewrote plain per-dimension
extents (IndexedPointer leaves), churning every non-linearized transfer
pragma and breaking the openacc codegen tests and the 01_gpu notebook.
Only a product of 32-bit sizes can overflow int, so restrict the
promotion to Mul bounds and update the notebook's reference output for
the linearized transfer clauses that legitimately gained the casts.

Author: binggao1230

devito/passes/iet/parpragma.py

@@ -495,10 +495,14 @@ def _generate(self):
         # is a product of the Function's 32-bit per-dimension sizes (e.g.
         # `size[0]*size[1]*size[2]`); for a Function with more than ~2**31
         # elements this product overflows `int` before being used as an array
-        # bound, producing a bogus transfer size (#2777). `as_long` promotes the
-        # size's symbolic leaves to 64-bit so the multiplication is carried out
-        # in 64-bit arithmetic.
-        sections = ''.join([f'[{ccode(i)}:{ccode(as_long(j))}]'
+        # bound, producing a bogus transfer size (#2777). Promote such products
+        # to 64-bit with `as_long` so the multiplication is carried out in
+        # 64-bit arithmetic. Non-product bounds (a single size, an offset, a
+        # constant) cannot overflow and are left untouched so that the
+        # generated code doesn't change where the cast serves no purpose.
+        def cast(e):
+            return as_long(e) if getattr(e, 'is_Mul', False) else e
+        sections = ''.join([f'[{ccode(cast(i))}:{ccode(cast(j))}]'
                             for i, j in self.sections])
         arguments = [ccode(i) for i in self.arguments]
         return self.pragma % (self.function.name, sections, *arguments)

examples/performance/01_gpu.ipynb

@@ -289,7 +289,7 @@
       "\n",
       "  float *u = (float *) u_vec->data;\n",
       "\n",
-      "  #pragma omp target enter data map(to: u[0:u_vec->size[0]*u_vec->size[1]*u_vec->size[2]])\n",
+      "  #pragma omp target enter data map(to: u[0:(long)(u_vec->size[2])*(long)(u_vec->size[1])*(long)(u_vec->size[0])])\n",
       "\n",
       "  const int x_fsz0 = u_vec->size[1];\n",
       "  const int y_fsz0 = u_vec->size[2];\n",
@@ -315,8 +315,8 @@
       "    STOP(section0,timers)\n",
       "  }\n",
       "\n",
-      "  #pragma omp target update from(u[0:u_vec->size[0]*u_vec->size[1]*u_vec->size[2]])\n",
-      "  #pragma omp target exit data map(release: u[0:u_vec->size[0]*u_vec->size[1]*u_vec->size[2]]) if(devicerm)\n",
+      "  #pragma omp target update from(u[0:(long)(u_vec->size[2])*(long)(u_vec->size[1])*(long)(u_vec->size[0])])\n",
+      "  #pragma omp target exit data map(release: u[0:(long)(u_vec->size[2])*(long)(u_vec->size[1])*(long)(u_vec->size[0])]) if(devicerm)\n",
       "\n",
       "  return 0;\n",
       "}\n",