code: def test_write_float_surface(self):
      width, height, depth = self.dimensions
      dim = 3 if depth != 0 else 2 if height != 0 else 1
  
      # 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 surface memory
      real_output = cupy.zeros(shape, dtype=cupy.float32)
      assert real_output.flags["C_CONTIGUOUS"]
      ch = ChannelFormatDescriptor(
          32, 0, 0, 0, runtime.cudaChannelFormatKindFloat
      )
      expected_output = cupy.arange(
          numpy.prod(shape), dtype=cupy.float32
      )
      expected_output = expected_output.reshape(shape) * 3.0
      assert expected_output.flags["C_CONTIGUOUS"]
  
      # create resource descriptor
      # note that surface memory only support CUDA array
      arr = CUDAarray(
          ch, width, height, depth, runtime.cudaArraySurfaceLoadStore
      )
      arr.copy_from(real_output)  # init to zero
      res = ResourceDescriptor(runtime.cudaResourceTypeArray, cuArr=arr)
  
      # create a surface object; currently we don't support surface reference
      surfobj = SurfaceObject(res)
      mod = cupy.RawModule(code=source_surfobj)
  
      # get and launch the kernel
      ker_name = "writeKernel"
      ker_name += "3D" if dim == 3 else "2D" if dim == 2 else "1D"
      ker = mod.get_function(ker_name)
      block = (4, 4, 2) if dim == 3 else (4, 4) if dim == 2 else (4,)
      grid = ()
      args = (surfobj,)
      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
      arr.copy_to(real_output)
      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 TestSurface obj's underlying member method code with example usage: obj.test_write_float_surface() without return types