code: def test_fetch_float_texture(self):
      width, height, depth = self.dimensions
      dim = 3 if depth != 0 else 2 if height != 0 else 1
  
      if (self.mem_type == "linear" and dim != 1) or (
          self.mem_type == "pitch2D" and dim != 2
      ):
          pytest.skip(
              "The test case {0} is inapplicable for {1} and thus "
              "skipped.".format(self.dimensions, self.mem_type)
          )
  
      # generate input data and allocate output buffer
      shape = (
          (depth, height, width)
          if dim == 3
          else (height, width)
          if dim == 2
          else (width,)
      )
  
      # prepare input, output, and texture memory
      tex_data = cupy.random.random(shape, dtype=cupy.float32)
      real_output = cupy.zeros_like(tex_data)
      ch = ChannelFormatDescriptor(
          32, 0, 0, 0, runtime.cudaChannelFormatKindFloat
      )
      assert tex_data.flags["C_CONTIGUOUS"]
      assert real_output.flags["C_CONTIGUOUS"]
      if self.mem_type == "CUDAarray":
          arr = CUDAarray(ch, width, height, depth)
          expected_output = cupy.zeros_like(tex_data)
          assert expected_output.flags["C_CONTIGUOUS"]
          # test bidirectional copy
          arr.copy_from(tex_data)
          arr.copy_to(expected_output)
      else:  # linear are pitch2D are backed by ndarray
          arr = tex_data
          expected_output = tex_data
  
      # create resource and texture descriptors
      if self.mem_type == "CUDAarray":
          res = ResourceDescriptor(
              runtime.cudaResourceTypeArray, cuArr=arr
          )
      elif self.mem_type == "linear":
          res = ResourceDescriptor(
              runtime.cudaResourceTypeLinear,
              arr=arr,
              chDesc=ch,
              sizeInBytes=arr.size * arr.dtype.itemsize,
          )
      else:  # pitch2D
          # In this case, we rely on the fact that the hand-picked array
          # shape meets the alignment requirement. This is CUDA's limitation,
          # see CUDA Runtime API reference guide. "TexturePitchAlignment" is
          # assumed to be 32, which should be applicable for most devices.
          res = ResourceDescriptor(
              runtime.cudaResourceTypePitch2D,
              arr=arr,
              chDesc=ch,
              width=width,
              height=height,
              pitchInBytes=width * arr.dtype.itemsize,
          )
      address_mode = (
          runtime.cudaAddressModeClamp,
          runtime.cudaAddressModeClamp,
      )
      tex = TextureDescriptor(
          address_mode,
          runtime.cudaFilterModePoint,
          runtime.cudaReadModeElementType,
      )
  
      if self.target == "object":
          # create a texture object
          texobj = TextureObject(res, tex)
          mod = cupy.RawModule(code=source_texobj)
      else:
          assert False
  
      # get and launch the kernel
      ker_name = "copyKernel"
      ker_name += "3D" if dim == 3 else "2D" if dim == 2 else "1D"
      ker_name += "fetch" if self.mem_type == "linear" else ""
      ker = mod.get_function(ker_name)
      block = (4, 4, 2) if dim == 3 else (4, 4) if dim == 2 else (4,)
      grid = ()
      args = (real_output,)
      if self.target == "object":
          args = args + (texobj,)
      if dim >= 1:
          grid_x = (width + block[0] - 1) // block[0]
          grid = grid + (grid_x,)
          args = args + (width,)
      if dim >= 2:
          grid_y = (height + block[1] - 1) // block[1]
          grid = grid + (grid_y,)
          args = args + (height,)
      if dim == 3:
          grid_z = (depth + block[2] - 1) // block[2]
          grid = grid + (grid_z,)
          args = args + (depth,)
      ker(grid, block, args)
  
      # validate result
      assert (real_output == expected_output).all()
  
file_path: cupy/tests/cupy_tests/cuda_tests/test_texture.py
imports: import numpy
  import pytest
  import cupy
  from cupy import testing
  from cupy.cuda import runtime
  from cupy.cuda.texture import ChannelFormatDescriptor, CUDAarray, ResourceDescriptor, TextureDescriptor, TextureObject, SurfaceObject
label_desc: Use image node_id 1 for calling the TestTexture obj's underlying member method code with example usage: obj.test_fetch_float_texture() without return types