code: def test_fp_textures(self):
      if drv.Context.get_device().compute_capability() < (1, 3):
          return
  
      for tp in [np.float32, np.float64]:
          tp_cstr = dtype_to_ctype(tp)
          mod = SourceModule(
              """
              #include <pycuda-helpers.hpp>
  
              texture<fp_tex_%(tp)s, 1, cudaReadModeElementType> my_tex;
  
              __global__ void copy_texture(%(tp)s *dest)
              {
                int i = threadIdx.x;
                dest[i] = fp_tex1Dfetch(my_tex, i);
              }
              """
              % {"tp": tp_cstr}
          )
  
          copy_texture = mod.get_function("copy_texture")
          my_tex = mod.get_texref("my_tex")
  
          shape = (384,)
          a = np.random.randn(*shape).astype(tp)
          a_gpu = gpuarray.to_gpu(a)
          a_gpu.bind_to_texref_ext(my_tex, allow_double_hack=True)
  
          dest = np.zeros(shape, dtype=tp)
          copy_texture(
              drv.Out(dest),
              block=shape
              + (
                  1,
                  1,
              ),
              texrefs=[my_tex],
          )
  
          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 27 for calling the TestDriver obj's underlying member method code with example usage: obj.test_fp_textures() without return types