code: def test_multichannel_linear_texture(self):
      mod = SourceModule(
          """
          #define CHANNELS 4
          texture<float4, 1, cudaReadModeElementType> mtx_tex;
  
          __global__ void copy_texture(float *dest)
          {
            int i = threadIdx.x+blockDim.x*threadIdx.y;
            float4 texval = tex1Dfetch(mtx_tex, i);
            dest[i*CHANNELS + 0] = texval.x;
            dest[i*CHANNELS + 1] = texval.y;
            dest[i*CHANNELS + 2] = texval.z;
            dest[i*CHANNELS + 3] = texval.w;
          }
          """
      )
  
      copy_texture = mod.get_function("copy_texture")
      mtx_tex = mod.get_texref("mtx_tex")
  
      shape = (16, 16)
      channels = 4
      a = np.random.randn(*(shape + (channels,))).astype(np.float32)
      a_gpu = drv.to_device(a)
      mtx_tex.set_address(a_gpu, a.nbytes)
      mtx_tex.set_format(drv.array_format.FLOAT, 4)
  
      dest = np.zeros(shape + (channels,), dtype=np.float32)
      copy_texture(drv.Out(dest), block=shape + (1,), texrefs=[mtx_tex])
      # print a
      # print dest
      assert la.norm(dest - a) == 0
  
imports: import numpy
  import numpy.linalg
  from pycuda.tools import mark_cuda_test, dtype_to_ctype
  import pytest
  import pycuda.gpuarray
  import pycuda.driver
  from pycuda.compiler import SourceModule
label_desc: Use image node_id 12 for calling the TestDriver obj's underlying member method code with example usage: obj.test_multichannel_linear_texture() without return types