code: def sobelFilter(odata, iw, ih):
      global array, pixels, mode, scale
      if mode == 3:
          # Texture and shared memory with fixed BlockSize
          sm = SourceModule(
              """
              texture<unsigned char, 2> tex;
              extern __shared__ unsigned char LocalBlock[];
              #define RADIUS 1
              #define BlockWidth 80
              #define SharedPitch 384
              __device__ unsigned char
              ComputeSobel(unsigned char ul, // upper left
                           unsigned char um, // upper middle
                           unsigned char ur, // upper right
                           unsigned char ml, // middle left
                           unsigned char mm, // middle (unused)
                           unsigned char mr, // middle right
                           unsigned char ll, // lower left
                           unsigned char lm, // lower middle
                           unsigned char lr, // lower right
                           float fScale )
              {
                  short Horz = ur + 2*mr + lr - ul - 2*ml - ll;
                  short Vert = ul + 2*um + ur - ll - 2*lm - lr;
                  short Sum = (short) (fScale*(::abs(int(Horz))+::abs(int(Vert))));
                  if ( Sum < 0 ) return 0; else if ( Sum > 0xff ) return 0xff;
                  return (unsigned char) Sum;
              }
  
              __global__ void
              SobelShared( int* pSobelOriginal, unsigned short SobelPitch,
                           short w, short h, float fScale )
              {
                  short u = 4*blockIdx.x*BlockWidth;
                  short v = blockIdx.y*blockDim.y + threadIdx.y;
                  short ib;
  
                  int SharedIdx = threadIdx.y * SharedPitch;
  
                  for ( ib = threadIdx.x; ib < BlockWidth+2*RADIUS; ib += blockDim.x ) {
                      LocalBlock[SharedIdx+4*ib+0] = tex2D( tex,
                          (float) (u+4*ib-RADIUS+0), (float) (v-RADIUS) );
                      LocalBlock[SharedIdx+4*ib+1] = tex2D( tex,
                          (float) (u+4*ib-RADIUS+1), (float) (v-RADIUS) );
                      LocalBlock[SharedIdx+4*ib+2] = tex2D( tex,
                          (float) (u+4*ib-RADIUS+2), (float) (v-RADIUS) );
                      LocalBlock[SharedIdx+4*ib+3] = tex2D( tex,
                          (float) (u+4*ib-RADIUS+3), (float) (v-RADIUS) );
                  }
                  if ( threadIdx.y < RADIUS*2 ) {
                      //
                      // copy trailing RADIUS*2 rows of pixels into shared
                      //
                      SharedIdx = (blockDim.y+threadIdx.y) * SharedPitch;
                      for ( ib = threadIdx.x; ib < BlockWidth+2*RADIUS; ib += blockDim.x ) {
                          LocalBlock[SharedIdx+4*ib+0] = tex2D( tex,
                              (float) (u+4*ib-RADIUS+0), (float) (v+blockDim.y-RADIUS) );
                          LocalBlock[SharedIdx+4*ib+1] = tex2D( tex,
                              (float) (u+4*ib-RADIUS+1), (float) (v+blockDim.y-RADIUS) );
                          LocalBlock[SharedIdx+4*ib+2] = tex2D( tex,
                              (float) (u+4*ib-RADIUS+2), (float) (v+blockDim.y-RADIUS) );
                          LocalBlock[SharedIdx+4*ib+3] = tex2D( tex,
                              (float) (u+4*ib-RADIUS+3), (float) (v+blockDim.y-RADIUS) );
                      }
                  }
  
                  __syncthreads();
  
                  u >>= 2;    // index as uchar4 from here
                  uchar4 *pSobel = (uchar4 *) (((char *) pSobelOriginal)+v*SobelPitch);
                  SharedIdx = threadIdx.y * SharedPitch;
  
                  for ( ib = threadIdx.x; ib < BlockWidth; ib += blockDim.x ) {
  
                      unsigned char pix00 = LocalBlock[SharedIdx+4*ib+0*SharedPitch+0];
                      unsigned char pix01 = LocalBlock[SharedIdx+4*ib+0*SharedPitch+1];
                      unsigned char pix02 = LocalBlock[SharedIdx+4*ib+0*SharedPitch+2];
                      unsigned char pix10 = LocalBlock[SharedIdx+4*ib+1*SharedPitch+0];
                      unsigned char pix11 = LocalBlock[SharedIdx+4*ib+1*SharedPitch+1];
                      unsigned char pix12 = LocalBlock[SharedIdx+4*ib+1*SharedPitch+2];
                      unsigned char pix20 = LocalBlock[SharedIdx+4*ib+2*SharedPitch+0];
                      unsigned char pix21 = LocalBlock[SharedIdx+4*ib+2*SharedPitch+1];
                      unsigned char pix22 = LocalBlock[SharedIdx+4*ib+2*SharedPitch+2];
  
                      uchar4 out;
  
                      out.x = ComputeSobel(pix00, pix01, pix02,
                                           pix10, pix11, pix12,
                                           pix20, pix21, pix22, fScale );
  
                      pix00 = LocalBlock[SharedIdx+4*ib+0*SharedPitch+3];
                      pix10 = LocalBlock[SharedIdx+4*ib+1*SharedPitch+3];
                      pix20 = LocalBlock[SharedIdx+4*ib+2*SharedPitch+3];
                      out.y = ComputeSobel(pix01, pix02, pix00,
                                           pix11, pix12, pix10,
                                           pix21, pix22, pix20, fScale );
  
                      pix01 = LocalBlock[SharedIdx+4*ib+0*SharedPitch+4];
                      pix11 = LocalBlock[SharedIdx+4*ib+1*SharedPitch+4];
                      pix21 = LocalBlock[SharedIdx+4*ib+2*SharedPitch+4];
                      out.z = ComputeSobel( pix02, pix00, pix01,
                                            pix12, pix10, pix11,
                                            pix22, pix20, pix21, fScale );
  
                      pix02 = LocalBlock[SharedIdx+4*ib+0*SharedPitch+5];
                      pix12 = LocalBlock[SharedIdx+4*ib+1*SharedPitch+5];
                      pix22 = LocalBlock[SharedIdx+4*ib+2*SharedPitch+5];
                      out.w = ComputeSobel( pix00, pix01, pix02,
                                            pix10, pix11, pix12,
                                            pix20, pix21, pix22, fScale );
                      if ( u+ib < w/4 && v < h ) {
                          pSobel[u+ib] = out;
                      }
                  }
  
                  __syncthreads();
              }
          """
          )
          cuda_function = sm.get_function("SobelShared")
      elif mode == 2:
          # Texture and shared memory with variable BlockSize
          sm = SourceModule(
              """
          #define RADIUS 1
          texture<unsigned char, 2> tex;
          extern __shared__ unsigned char LocalBlock[];
          __device__ unsigned char
          ComputeSobel(unsigned char ul, // upper left
                       unsigned char um, // upper middle
                       unsigned char ur, // upper right
                       unsigned char ml, // middle left
                       unsigned char mm, // middle (unused)
                       unsigned char mr, // middle right
                       unsigned char ll, // lower left
                       unsigned char lm, // lower middle
                       unsigned char lr, // lower right
                       float fScale )
          {
              short Horz = ur + 2*mr + lr - ul - 2*ml - ll;
              short Vert = ul + 2*um + ur - ll - 2*lm - lr;
              short Sum = (short) (fScale*(::abs(int(Horz))+::abs(int(Vert))));
              if ( Sum < 0 ) return 0; else if ( Sum > 0xff ) return 0xff;
              return (unsigned char) Sum;
          }
  
          __global__ void
          SobelShared( int* pSobelOriginal, unsigned short SobelPitch,
                       short BlockWidth, short SharedPitch,
                       short w, short h, float fScale )
          {
              short u = 4*blockIdx.x*BlockWidth;
              short v = blockIdx.y*blockDim.y + threadIdx.y;
              short ib;
  
              int SharedIdx = threadIdx.y * SharedPitch;
  
              for ( ib = threadIdx.x; ib < BlockWidth+2*RADIUS; ib += blockDim.x ) {
                  LocalBlock[SharedIdx+4*ib+0] = tex2D( tex,
                      (float) (u+4*ib-RADIUS+0), (float) (v-RADIUS) );
                  LocalBlock[SharedIdx+4*ib+1] = tex2D( tex,
                      (float) (u+4*ib-RADIUS+1), (float) (v-RADIUS) );
                  LocalBlock[SharedIdx+4*ib+2] = tex2D( tex,
                      (float) (u+4*ib-RADIUS+2), (float) (v-RADIUS) );
                  LocalBlock[SharedIdx+4*ib+3] = tex2D( tex,
                      (float) (u+4*ib-RADIUS+3), (float) (v-RADIUS) );
              }
              if ( threadIdx.y < RADIUS*2 ) {
                  //
                  // copy trailing RADIUS*2 rows of pixels into shared
                  //
                  SharedIdx = (blockDim.y+threadIdx.y) * SharedPitch;
                  for ( ib = threadIdx.x; ib < BlockWidth+2*RADIUS; ib += blockDim.x ) {
                      LocalBlock[SharedIdx+4*ib+0] = tex2D( tex,
                          (float) (u+4*ib-RADIUS+0), (float) (v+blockDim.y-RADIUS) );
                      LocalBlock[SharedIdx+4*ib+1] = tex2D( tex,
                          (float) (u+4*ib-RADIUS+1), (float) (v+blockDim.y-RADIUS) );
                      LocalBlock[SharedIdx+4*ib+2] = tex2D( tex,
                          (float) (u+4*ib-RADIUS+2), (float) (v+blockDim.y-RADIUS) );
                      LocalBlock[SharedIdx+4*ib+3] = tex2D( tex,
                          (float) (u+4*ib-RADIUS+3), (float) (v+blockDim.y-RADIUS) );
                  }
              }
  
              __syncthreads();
  
              u >>= 2;    // index as uchar4 from here
              uchar4 *pSobel = (uchar4 *) (((char *) pSobelOriginal)+v*SobelPitch);
              SharedIdx = threadIdx.y * SharedPitch;
  
              for ( ib = threadIdx.x; ib < BlockWidth; ib += blockDim.x ) {
  
                  unsigned char pix00 = LocalBlock[SharedIdx+4*ib+0*SharedPitch+0];
                  unsigned char pix01 = LocalBlock[SharedIdx+4*ib+0*SharedPitch+1];
                  unsigned char pix02 = LocalBlock[SharedIdx+4*ib+0*SharedPitch+2];
                  unsigned char pix10 = LocalBlock[SharedIdx+4*ib+1*SharedPitch+0];
                  unsigned char pix11 = LocalBlock[SharedIdx+4*ib+1*SharedPitch+1];
                  unsigned char pix12 = LocalBlock[SharedIdx+4*ib+1*SharedPitch+2];
                  unsigned char pix20 = LocalBlock[SharedIdx+4*ib+2*SharedPitch+0];
                  unsigned char pix21 = LocalBlock[SharedIdx+4*ib+2*SharedPitch+1];
                  unsigned char pix22 = LocalBlock[SharedIdx+4*ib+2*SharedPitch+2];
  
                  uchar4 out;
  
                  out.x = ComputeSobel(pix00, pix01, pix02,
                                       pix10, pix11, pix12,
                                       pix20, pix21, pix22, fScale );
  
                  pix00 = LocalBlock[SharedIdx+4*ib+0*SharedPitch+3];
                  pix10 = LocalBlock[SharedIdx+4*ib+1*SharedPitch+3];
                  pix20 = LocalBlock[SharedIdx+4*ib+2*SharedPitch+3];
                  out.y = ComputeSobel(pix01, pix02, pix00,
                                       pix11, pix12, pix10,
                                       pix21, pix22, pix20, fScale );
  
                  pix01 = LocalBlock[SharedIdx+4*ib+0*SharedPitch+4];
                  pix11 = LocalBlock[SharedIdx+4*ib+1*SharedPitch+4];
                  pix21 = LocalBlock[SharedIdx+4*ib+2*SharedPitch+4];
                  out.z = ComputeSobel( pix02, pix00, pix01,
                                        pix12, pix10, pix11,
                                        pix22, pix20, pix21, fScale );
  
                  pix02 = LocalBlock[SharedIdx+4*ib+0*SharedPitch+5];
                  pix12 = LocalBlock[SharedIdx+4*ib+1*SharedPitch+5];
                  pix22 = LocalBlock[SharedIdx+4*ib+2*SharedPitch+5];
                  out.w = ComputeSobel( pix00, pix01, pix02,
                                        pix10, pix11, pix12,
                                        pix20, pix21, pix22, fScale );
                  if ( u+ib < w/4 && v < h ) {
                      pSobel[u+ib] = out;
                  }
              }
  
              __syncthreads();
          }
  
          """
          )
          cuda_function = sm.get_function("SobelShared")
      if mode == 1:
          # Just Texture
          sm = SourceModule(
              """
          texture<unsigned char, 2> tex;
          __device__ unsigned char ComputeSobel(unsigned char ul, // upper left
                       unsigned char um, // upper middle
                       unsigned char ur, // upper right
                       unsigned char ml, // middle left
                       unsigned char mm, // middle (unused)
                       unsigned char mr, // middle right
                       unsigned char ll, // lower left
                       unsigned char lm, // lower middle
                       unsigned char lr, // lower right
                       float fScale )
          {
              short Horz = ur + 2*mr + lr - ul - 2*ml - ll;
              short Vert = ul + 2*um + ur - ll - 2*lm - lr;
              short Sum = (short) (fScale*(::abs(int(Horz))+::abs(int(Vert))));
              if ( Sum < 0 ) return 0; else if ( Sum > 0xff ) return 0xff;
              return (unsigned char) Sum;
          }
          __global__ void SobelTex( int* pSobelOriginal, unsigned int Pitch,
                    int w, int h, float fScale )
          {
              unsigned char *pSobel =
                (unsigned char *) (((char *) pSobelOriginal)+blockIdx.x*Pitch);
              for ( int i = threadIdx.x; i < w; i += blockDim.x ) {
                  unsigned char pix00 = tex2D( tex, (float) i-1, (float) blockIdx.x-1 );
                  unsigned char pix01 = tex2D( tex, (float) i+0, (float) blockIdx.x-1 );
                  unsigned char pix02 = tex2D( tex, (float) i+1, (float) blockIdx.x-1 );
                  unsigned char pix10 = tex2D( tex, (float) i-1, (float) blockIdx.x+0 );
                  unsigned char pix11 = tex2D( tex, (float) i+0, (float) blockIdx.x+0 );
                  unsigned char pix12 = tex2D( tex, (float) i+1, (float) blockIdx.x+0 );
                  unsigned char pix20 = tex2D( tex, (float) i-1, (float) blockIdx.x+1 );
                  unsigned char pix21 = tex2D( tex, (float) i+0, (float) blockIdx.x+1 );
                  unsigned char pix22 = tex2D( tex, (float) i+1, (float) blockIdx.x+1 );
                  pSobel[i] = ComputeSobel(pix00, pix01, pix02,
                                           pix10, pix11, pix12,
                                           pix20, pix21, pix22, fScale );
              }
          }
          """
          )
          cuda_function = sm.get_function("SobelTex")
      elif mode == 0:
          # Just Copy
          sm = SourceModule(
              """
          texture<unsigned char, 2> tex;
          __global__ void SobelCopyImage(int* pSobelOriginal, unsigned int Pitch, int w, int h, float fscale )
          {
              unsigned char *pSobel =
                (unsigned char *) (((unsigned char *) pSobelOriginal)+blockIdx.x*Pitch);
              for ( int i = threadIdx.x; i < w; i += blockDim.x ) {
                  pSobel[i] = min( max((tex2D( tex, (float)i, (float)blockIdx.x ) * fscale), 0.f), 255.f);
              }
          }
          """
          )
          cuda_function = sm.get_function("SobelCopyImage")
      texref = sm.get_texref("tex")
      texref.set_array(array)
      texref.set_flags(cuda_driver.TRSA_OVERRIDE_FORMAT)
      if mode == 3:
          # fixed BlockSize Launch
          RADIUS = 1
          threads = (16, 4, 1)
          BlockWidth = 80  # Do not change!
          blocks = (
              iw / (4 * BlockWidth) + (0 != iw % (4 * BlockWidth)),
              ih / threads[1] + (0 != ih % threads[1]),
          )
          SharedPitch = ~0x3F & (4 * (BlockWidth + 2 * RADIUS) + 0x3F)
          sharedMem = SharedPitch * (threads[1] + 2 * RADIUS)
          iw = iw & ~3
          cuda_function(
              np.intp(odata),
              np.uint16(iw),
              np.int16(iw),
              np.int16(ih),
              np.float32(scale),
              texrefs=[texref],
              block=threads,
              grid=blocks,
              shared=sharedMem,
          )
      elif mode == 2:
          # variable BlockSize launch
          RADIUS = 1
          threads = (16, 4, 1)
          BlockWidth = 80  # Change only with divisible by 16 values!
          blocks = (
              iw / (4 * BlockWidth) + (0 != iw % (4 * BlockWidth)),
              ih / threads[1] + (0 != ih % threads[1]),
          )
          SharedPitch = ~0x3F & (4 * (BlockWidth + 2 * RADIUS) + 0x3F)
          sharedMem = SharedPitch * (threads[1] + 2 * RADIUS)
          iw = iw & ~3
          cuda_function(
              np.intp(odata),
              np.uint16(iw),
              np.int16(BlockWidth),
              np.int16(SharedPitch),
              np.int16(iw),
              np.int16(ih),
              np.float32(scale),
              texrefs=[texref],
              block=threads,
              grid=blocks,
              shared=sharedMem,
          )
      else:
          BlockWidth = 384
          cuda_function(
              np.intp(odata),
              np.uint32(iw),
              np.int32(iw),
              np.int32(ih),
              np.float32(scale),
              texrefs=[texref],
              block=(BlockWidth, 1, 1),
              grid=(ih, 1),
          )
  
imports: from OpenGL.GL import *
  from OpenGL.GLUT import *
  from OpenGL.GLU import *
  from OpenGL.GL.ARB.vertex_buffer_object import *
  import numpy, Image
  import sys, time, os
  import pycuda.driver
  import pycuda.gl
  import pycuda
  from pycuda.compiler import SourceModule
label_desc: Use image node_id 3 for calling a global function with example usage: sobelFilter(odata, iw, ih) without return types