diff --git a/c_api/include/taichi/cpp/taichi.hpp b/c_api/include/taichi/cpp/taichi.hpp
index 3f673d7f1c091..1dd43f050cd6f 100644
--- a/c_api/include/taichi/cpp/taichi.hpp
+++ b/c_api/include/taichi/cpp/taichi.hpp
@@ -1,4 +1,5 @@
 // C++ wrapper of Taichi C-API
+#include <cstring>
 #include <list>
 #include <vector>
 #include <string>
@@ -153,13 +154,34 @@ class NdArray {
     return *this;
   }
 
-  void *map() {
+  inline void *map() {
     return ti_map_memory(runtime_, ndarray_.memory);
   }
-  void unmap() {
+  inline void unmap() {
     return ti_unmap_memory(runtime_, ndarray_.memory);
   }
 
+  inline void read(T *dst, size_t size) {
+    T *src = (T *)map();
+    if (src != nullptr) {
+      std::memcpy(dst, src, size);
+    }
+    unmap();
+  }
+  inline void read(std::vector<T> &dst) {
+    read(dst.data(), dst.size() * sizeof(T));
+  }
+  inline void write(const T *src, size_t size) {
+    T *dst = (T *)map();
+    if (dst != nullptr) {
+      std::memcpy(dst, src, size);
+    }
+    unmap();
+  }
+  inline void write(const std::vector<T> &src) {
+    write(src.data(), src.size() * sizeof(T));
+  }
+
   constexpr TiMemory memory() const {
     return ndarray_.memory;
   }
@@ -171,33 +193,89 @@ class NdArray {
   }
 };
 
-class Texture {
+class Image {
   TiRuntime runtime_{TI_NULL_HANDLE};
-  TiTexture texture_{TI_NULL_HANDLE};
+  TiImage image_{TI_NULL_HANDLE};
   bool should_destroy_{false};
 
  public:
   constexpr bool is_valid() const {
-    return texture_ != nullptr;
+    return image_ != nullptr;
   }
   inline void destroy() {
     if (should_destroy_) {
-      ti_free_texture(runtime_, texture_);
-      texture_ = TI_NULL_HANDLE;
+      ti_free_image(runtime_, image_);
+      image_ = TI_NULL_HANDLE;
       should_destroy_ = false;
     }
   }
 
+  Image() {
+  }
+  Image(const Image &b) = delete;
+  Image(Image &&b)
+      : runtime_(detail::move_handle(b.runtime_)),
+        image_(detail::move_handle(b.image_)),
+        should_destroy_(std::exchange(b.should_destroy_, false)) {
+  }
+  Image(TiRuntime runtime, TiImage image, bool should_destroy)
+      : runtime_(runtime), image_(image), should_destroy_(should_destroy) {
+  }
+  ~Image() {
+    destroy();
+  }
+
+  Image &operator=(const Image &) = delete;
+  Image &operator=(Image &&b) {
+    destroy();
+    runtime_ = detail::move_handle(b.runtime_);
+    image_ = detail::move_handle(b.image_);
+    should_destroy_ = std::exchange(b.should_destroy_, false);
+    return *this;
+  }
+
+  TiImageSlice slice(TiImageOffset offset,
+                     TiImageExtent extent,
+                     uint32_t mip_level) const {
+    TiImageSlice slice{};
+    slice.image = image_;
+    slice.extent = extent;
+    slice.offset = offset;
+    slice.mip_level = mip_level;
+    return slice;
+  }
+
+  constexpr TiImage image() const {
+    return image_;
+  }
+  constexpr operator TiImage() const {
+    return image_;
+  }
+};
+
+class Texture {
+  Image image_{};
+  TiTexture texture_{};
+
+ public:
+  constexpr bool is_valid() const {
+    return image_.is_valid();
+  }
+  inline void destroy() {
+    image_.destroy();
+  }
+
   Texture() {
   }
   Texture(const Texture &b) = delete;
   Texture(Texture &&b)
-      : runtime_(detail::move_handle(b.runtime_)),
-        texture_(detail::move_handle(b.texture_)),
-        should_destroy_(std::exchange(b.should_destroy_, false)) {
+      : image_(std::move(b.image_)), texture_(std::move(b.texture_)) {
   }
-  Texture(TiRuntime runtime, TiTexture texture, bool should_destroy)
-      : runtime_(runtime), texture_(texture), should_destroy_(should_destroy) {
+  Texture(Image &&image, const TiTexture &texture)
+      : image_(std::move(image)), texture_(texture) {
+    if (texture.image != image_) {
+      ti_set_last_error(TI_ERROR_INVALID_ARGUMENT, "texture.image != image");
+    }
   }
   ~Texture() {
     destroy();
@@ -206,9 +284,8 @@ class Texture {
   Texture &operator=(const Texture &) = delete;
   Texture &operator=(Texture &&b) {
     destroy();
-    runtime_ = detail::move_handle(b.runtime_);
-    texture_ = detail::move_handle(b.texture_);
-    should_destroy_ = std::exchange(b.should_destroy_, false);
+    image_ = std::move(b.image_);
+    texture_ = std::move(b.texture_);
     return *this;
   }
 
@@ -227,7 +304,8 @@ class ArgumentEntry {
  public:
   ArgumentEntry() = delete;
   ArgumentEntry(const ArgumentEntry &) = delete;
-  ArgumentEntry(ArgumentEntry &&) = delete;
+  ArgumentEntry(ArgumentEntry &&b) : arg_(b.arg_) {
+  }
   ArgumentEntry(TiArgument *arg) : arg_(arg) {
   }
 
@@ -250,7 +328,7 @@ class ArgumentEntry {
     arg_->value.ndarray = ndarray;
     return *this;
   }
-  inline ArgumentEntry &operator=(TiTexture texture) {
+  inline ArgumentEntry &operator=(const TiTexture &texture) {
     arg_->type = TI_ARGUMENT_TYPE_TEXTURE;
     arg_->value.texture = texture;
     return *this;
@@ -324,7 +402,7 @@ class ComputeGraph {
     return at(name);
   }
 
-  void launch(size_t argument_count, const TiNamedArgument *arguments) {
+  void launch(uint32_t argument_count, const TiNamedArgument *arguments) {
     ti_launch_compute_graph(runtime_, compute_graph_, argument_count,
                             arguments);
   }
@@ -385,7 +463,7 @@ class Kernel {
     return at(i);
   }
 
-  void launch(size_t argument_count, const TiArgument *arguments) {
+  void launch(uint32_t argument_count, const TiArgument *arguments) {
     ti_launch_kernel(runtime_, kernel_, argument_count, arguments);
   }
   void launch() {
@@ -595,23 +673,36 @@ class Runtime {
     return NdArray<T>(runtime_, std::move(ndarray), true);
   }
 
-  Texture allocate_texture(const TiTextureAllocateInfo &allocate_info) {
-    TiTexture texture = ti_allocate_texture(runtime_, &allocate_info);
-    return Texture(runtime_, texture, true);
+  Image allocate_image(const TiImageAllocateInfo &allocate_info) {
+    TiImage image = ti_allocate_image(runtime_, &allocate_info);
+    return Image(runtime_, image, true);
   }
   Texture allocate_texture2d(uint32_t width,
                              uint32_t height,
-                             TiTextureFormat format) {
-    TiTextureAllocateInfo allocate_info{};
-    allocate_info.dimension = TI_TEXTURE_DIMENSION_2D;
-    allocate_info.extent.array_layer_count = 1;
-    allocate_info.extent.width = width;
-    allocate_info.extent.height = height;
-    allocate_info.extent.depth = 1;
+                             TiFormat format,
+                             TiSampler sampler) {
+    TiImageExtent extent{};
+    extent.width = width;
+    extent.height = height;
+    extent.depth = 1;
+    extent.array_layer_count = 1;
+
+    TiImageAllocateInfo allocate_info{};
+    allocate_info.dimension = TI_IMAGE_DIMENSION_2D;
+    allocate_info.extent = extent;
     allocate_info.mip_level_count = 1;
     allocate_info.format = format;
-    allocate_info.usage = TI_TEXTURE_USAGE_STORAGE_BIT;
-    return allocate_texture(allocate_info);
+    allocate_info.usage =
+        TI_IMAGE_USAGE_STORAGE_BIT | TI_IMAGE_USAGE_SAMPLED_BIT;
+
+    Image image = allocate_image(allocate_info);
+    TiTexture texture{};
+    texture.image = image;
+    texture.dimension = TI_IMAGE_DIMENSION_2D;
+    texture.extent = extent;
+    texture.format = format;
+    texture.sampler = sampler;
+    return Texture(std::move(image), texture);
   }
 
   AotModule load_aot_module(const char *path) {
@@ -626,12 +717,12 @@ class Runtime {
                                     const TiMemorySlice &src_memory) {
     ti_copy_memory_device_to_device(runtime_, &dst_memory, &src_memory);
   }
-  void copy_texture_device_to_device(const TiTextureSlice &dst_texture,
-                                     const TiTextureSlice &src_texture) {
-    ti_copy_texture_device_to_device(runtime_, &dst_texture, &src_texture);
+  void copy_image_device_to_device(const TiImageSlice &dst_texture,
+                                   const TiImageSlice &src_texture) {
+    ti_copy_image_device_to_device(runtime_, &dst_texture, &src_texture);
   }
-  void transition_texture(TiTexture texture, TiTextureLayout layout) {
-    ti_transition_texture(runtime_, texture, layout);
+  void transition_image(TiImage image, TiImageLayout layout) {
+    ti_transition_image(runtime_, image, layout);
   }
 
   void submit() {
diff --git a/c_api/include/taichi/taichi_core.h b/c_api/include/taichi/taichi_core.h
index 34e5626be603c..4b23cceb33a19 100644
--- a/c_api/include/taichi/taichi_core.h
+++ b/c_api/include/taichi/taichi_core.h
@@ -32,8 +32,11 @@ typedef struct TiEvent_t *TiEvent;
 // handle.memory
 typedef struct TiMemory_t *TiMemory;
 
-// handle.texture
-typedef struct TiTexture_t *TiTexture;
+// handle.image
+typedef struct TiImage_t *TiImage;
+
+// handle.sampler
+typedef struct TiSampler_t *TiSampler;
 
 // handle.kernel
 typedef struct TiKernel_t *TiKernel;
@@ -67,9 +70,10 @@ typedef enum TiArch {
   TI_ARCH_METAL = 6,
   TI_ARCH_OPENGL = 7,
   TI_ARCH_DX11 = 8,
-  TI_ARCH_OPENCL = 9,
-  TI_ARCH_AMDGPU = 10,
-  TI_ARCH_VULKAN = 11,
+  TI_ARCH_DX12 = 9,
+  TI_ARCH_OPENCL = 10,
+  TI_ARCH_AMDGPU = 11,
+  TI_ARCH_VULKAN = 12,
   TI_ARCH_MAX_ENUM = 0xffffffff,
 } TiArch;
 
@@ -116,7 +120,7 @@ typedef struct TiMemoryAllocateInfo {
   TiBool host_write;
   TiBool host_read;
   TiBool export_sharing;
-  TiMemoryUsageFlagBits usage;
+  TiMemoryUsageFlags usage;
 } TiMemoryAllocateInfo;
 
 // structure.memory_slice
@@ -140,122 +144,154 @@ typedef struct TiNdArray {
   TiDataType elem_type;
 } TiNdArray;
 
-// bit_field.texture_usage
-typedef enum TiTextureUsageFlagBits {
-  TI_TEXTURE_USAGE_STORAGE_BIT = 1 << 0,
-  TI_TEXTURE_USAGE_SAMPLED_BIT = 1 << 1,
-  TI_TEXTURE_USAGE_ATTACHMENT_BIT = 1 << 2,
-} TiTextureUsageFlagBits;
-typedef TiFlags TiTextureUsageFlags;
-
-// enumeration.texture_dimension
-typedef enum TiTextureDimension {
-  TI_TEXTURE_DIMENSION_1D = 0,
-  TI_TEXTURE_DIMENSION_2D = 1,
-  TI_TEXTURE_DIMENSION_3D = 2,
-  TI_TEXTURE_DIMENSION_1D_ARRAY = 3,
-  TI_TEXTURE_DIMENSION_2D_ARRAY = 4,
-  TI_TEXTURE_DIMENSION_CUBE = 5,
-  TI_TEXTURE_DIMENSION_MAX_ENUM = 0xffffffff,
-} TiTextureDimension;
-
-// enumeration.texture_layout
-typedef enum TiTextureLayout {
-  TI_TEXTURE_LAYOUT_UNDEFINED = 0,
-  TI_TEXTURE_LAYOUT_SHADER_READ = 1,
-  TI_TEXTURE_LAYOUT_SHADER_WRITE = 2,
-  TI_TEXTURE_LAYOUT_SHADER_READ_WRITE = 3,
-  TI_TEXTURE_LAYOUT_COLOR_ATTACHMENT = 4,
-  TI_TEXTURE_LAYOUT_COLOR_ATTACHMENT_READ = 5,
-  TI_TEXTURE_LAYOUT_DEPTH_ATTACHMENT = 6,
-  TI_TEXTURE_LAYOUT_DEPTH_ATTACHMENT_READ = 7,
-  TI_TEXTURE_LAYOUT_TRANSFER_DST = 8,
-  TI_TEXTURE_LAYOUT_TRANSFER_SRC = 9,
-  TI_TEXTURE_LAYOUT_PRESENT_SRC = 10,
-  TI_TEXTURE_LAYOUT_MAX_ENUM = 0xffffffff,
-} TiTextureLayout;
-
-// enumeration.texture_format
-typedef enum TiTextureFormat {
-  TI_TEXTURE_FORMAT_UNKNOWN = 0,
-  TI_TEXTURE_FORMAT_R8 = 1,
-  TI_TEXTURE_FORMAT_RG8 = 2,
-  TI_TEXTURE_FORMAT_RGBA8 = 3,
-  TI_TEXTURE_FORMAT_RGBA8SRGB = 4,
-  TI_TEXTURE_FORMAT_BGRA8 = 5,
-  TI_TEXTURE_FORMAT_BGRA8SRGB = 6,
-  TI_TEXTURE_FORMAT_R8U = 7,
-  TI_TEXTURE_FORMAT_RG8U = 8,
-  TI_TEXTURE_FORMAT_RGBA8U = 9,
-  TI_TEXTURE_FORMAT_R8I = 10,
-  TI_TEXTURE_FORMAT_RG8I = 11,
-  TI_TEXTURE_FORMAT_RGBA8I = 12,
-  TI_TEXTURE_FORMAT_R16 = 13,
-  TI_TEXTURE_FORMAT_RG16 = 14,
-  TI_TEXTURE_FORMAT_RGB16 = 15,
-  TI_TEXTURE_FORMAT_RGBA16 = 16,
-  TI_TEXTURE_FORMAT_R16U = 17,
-  TI_TEXTURE_FORMAT_RG16U = 18,
-  TI_TEXTURE_FORMAT_RGB16U = 19,
-  TI_TEXTURE_FORMAT_RGBA16U = 20,
-  TI_TEXTURE_FORMAT_R16I = 21,
-  TI_TEXTURE_FORMAT_RG16I = 22,
-  TI_TEXTURE_FORMAT_RGB16I = 23,
-  TI_TEXTURE_FORMAT_RGBA16I = 24,
-  TI_TEXTURE_FORMAT_R16F = 25,
-  TI_TEXTURE_FORMAT_RG16F = 26,
-  TI_TEXTURE_FORMAT_RGB16F = 27,
-  TI_TEXTURE_FORMAT_RGBA16F = 28,
-  TI_TEXTURE_FORMAT_R32U = 29,
-  TI_TEXTURE_FORMAT_RG32U = 30,
-  TI_TEXTURE_FORMAT_RGB32U = 31,
-  TI_TEXTURE_FORMAT_RGBA32U = 32,
-  TI_TEXTURE_FORMAT_R32I = 33,
-  TI_TEXTURE_FORMAT_RG32I = 34,
-  TI_TEXTURE_FORMAT_RGB32I = 35,
-  TI_TEXTURE_FORMAT_RGBA32I = 36,
-  TI_TEXTURE_FORMAT_R32F = 37,
-  TI_TEXTURE_FORMAT_RG32F = 38,
-  TI_TEXTURE_FORMAT_RGB32F = 39,
-  TI_TEXTURE_FORMAT_RGBA32F = 40,
-  TI_TEXTURE_FORMAT_DEPTH16 = 41,
-  TI_TEXTURE_FORMAT_DEPTH24STENCIL8 = 42,
-  TI_TEXTURE_FORMAT_DEPTH32F = 43,
-  TI_TEXTURE_FORMAT_MAX_ENUM = 0xffffffff,
-} TiTextureFormat;
-
-// structure.texture_offset
-typedef struct TiTextureOffset {
+// bit_field.image_usage
+typedef enum TiImageUsageFlagBits {
+  TI_IMAGE_USAGE_STORAGE_BIT = 1 << 0,
+  TI_IMAGE_USAGE_SAMPLED_BIT = 1 << 1,
+  TI_IMAGE_USAGE_ATTACHMENT_BIT = 1 << 2,
+} TiImageUsageFlagBits;
+typedef TiFlags TiImageUsageFlags;
+
+// enumeration.image_dimension
+typedef enum TiImageDimension {
+  TI_IMAGE_DIMENSION_1D = 0,
+  TI_IMAGE_DIMENSION_2D = 1,
+  TI_IMAGE_DIMENSION_3D = 2,
+  TI_IMAGE_DIMENSION_1D_ARRAY = 3,
+  TI_IMAGE_DIMENSION_2D_ARRAY = 4,
+  TI_IMAGE_DIMENSION_CUBE = 5,
+  TI_IMAGE_DIMENSION_MAX_ENUM = 0xffffffff,
+} TiImageDimension;
+
+// enumeration.image_layout
+typedef enum TiImageLayout {
+  TI_IMAGE_LAYOUT_UNDEFINED = 0,
+  TI_IMAGE_LAYOUT_SHADER_READ = 1,
+  TI_IMAGE_LAYOUT_SHADER_WRITE = 2,
+  TI_IMAGE_LAYOUT_SHADER_READ_WRITE = 3,
+  TI_IMAGE_LAYOUT_COLOR_ATTACHMENT = 4,
+  TI_IMAGE_LAYOUT_COLOR_ATTACHMENT_READ = 5,
+  TI_IMAGE_LAYOUT_DEPTH_ATTACHMENT = 6,
+  TI_IMAGE_LAYOUT_DEPTH_ATTACHMENT_READ = 7,
+  TI_IMAGE_LAYOUT_TRANSFER_DST = 8,
+  TI_IMAGE_LAYOUT_TRANSFER_SRC = 9,
+  TI_IMAGE_LAYOUT_PRESENT_SRC = 10,
+  TI_IMAGE_LAYOUT_MAX_ENUM = 0xffffffff,
+} TiImageLayout;
+
+// enumeration.format
+typedef enum TiFormat {
+  TI_FORMAT_UNKNOWN = 0,
+  TI_FORMAT_R8 = 1,
+  TI_FORMAT_RG8 = 2,
+  TI_FORMAT_RGBA8 = 3,
+  TI_FORMAT_RGBA8SRGB = 4,
+  TI_FORMAT_BGRA8 = 5,
+  TI_FORMAT_BGRA8SRGB = 6,
+  TI_FORMAT_R8U = 7,
+  TI_FORMAT_RG8U = 8,
+  TI_FORMAT_RGBA8U = 9,
+  TI_FORMAT_R8I = 10,
+  TI_FORMAT_RG8I = 11,
+  TI_FORMAT_RGBA8I = 12,
+  TI_FORMAT_R16 = 13,
+  TI_FORMAT_RG16 = 14,
+  TI_FORMAT_RGB16 = 15,
+  TI_FORMAT_RGBA16 = 16,
+  TI_FORMAT_R16U = 17,
+  TI_FORMAT_RG16U = 18,
+  TI_FORMAT_RGB16U = 19,
+  TI_FORMAT_RGBA16U = 20,
+  TI_FORMAT_R16I = 21,
+  TI_FORMAT_RG16I = 22,
+  TI_FORMAT_RGB16I = 23,
+  TI_FORMAT_RGBA16I = 24,
+  TI_FORMAT_R16F = 25,
+  TI_FORMAT_RG16F = 26,
+  TI_FORMAT_RGB16F = 27,
+  TI_FORMAT_RGBA16F = 28,
+  TI_FORMAT_R32U = 29,
+  TI_FORMAT_RG32U = 30,
+  TI_FORMAT_RGB32U = 31,
+  TI_FORMAT_RGBA32U = 32,
+  TI_FORMAT_R32I = 33,
+  TI_FORMAT_RG32I = 34,
+  TI_FORMAT_RGB32I = 35,
+  TI_FORMAT_RGBA32I = 36,
+  TI_FORMAT_R32F = 37,
+  TI_FORMAT_RG32F = 38,
+  TI_FORMAT_RGB32F = 39,
+  TI_FORMAT_RGBA32F = 40,
+  TI_FORMAT_DEPTH16 = 41,
+  TI_FORMAT_DEPTH24STENCIL8 = 42,
+  TI_FORMAT_DEPTH32F = 43,
+  TI_FORMAT_MAX_ENUM = 0xffffffff,
+} TiFormat;
+
+// structure.image_offset
+typedef struct TiImageOffset {
   uint32_t x;
   uint32_t y;
   uint32_t z;
   uint32_t array_layer_offset;
-} TiTextureOffset;
+} TiImageOffset;
 
-// structure.texture_extent
-typedef struct TiTextureExtent {
+// structure.image_extent
+typedef struct TiImageExtent {
   uint32_t width;
   uint32_t height;
   uint32_t depth;
   uint32_t array_layer_count;
-} TiTextureExtent;
+} TiImageExtent;
 
-// structure.texture_allocate_info
-typedef struct TiTextureAllocateInfo {
-  TiTextureDimension dimension;
-  TiTextureExtent extent;
+// structure.image_allocate_info
+typedef struct TiImageAllocateInfo {
+  TiImageDimension dimension;
+  TiImageExtent extent;
   uint32_t mip_level_count;
-  TiTextureFormat format;
-  TiTextureUsageFlagBits usage;
-} TiTextureAllocateInfo;
-
-// structure.texture_slice
-typedef struct TiTextureSlice {
-  TiTexture texture;
-  TiTextureOffset offset;
-  TiTextureExtent extent;
+  TiFormat format;
+  TiImageUsageFlags usage;
+} TiImageAllocateInfo;
+
+// structure.image_slice
+typedef struct TiImageSlice {
+  TiImage image;
+  TiImageOffset offset;
+  TiImageExtent extent;
   uint32_t mip_level;
-} TiTextureSlice;
+} TiImageSlice;
+
+// enumeration.filter
+typedef enum TiFilter {
+  TI_FILTER_NEAREST = 0,
+  TI_FILTER_LINEAR = 1,
+  TI_FILTER_MAX_ENUM = 0xffffffff,
+} TiFilter;
+
+// enumeration.address_mode
+typedef enum TiAddressMode {
+  TI_ADDRESS_MODE_REPEAT = 0,
+  TI_ADDRESS_MODE_MIRRORED_REPEAT = 1,
+  TI_ADDRESS_MODE_CLAMP_TO_EDGE = 2,
+  TI_ADDRESS_MODE_MAX_ENUM = 0xffffffff,
+} TiAddressMode;
+
+// structure.sampler_create_info
+typedef struct TiSamplerCreateInfo {
+  TiFilter mag_filter;
+  TiFilter min_filter;
+  TiAddressMode address_mode;
+  float max_anisotropy;
+} TiSamplerCreateInfo;
+
+// structure.texture
+typedef struct TiTexture {
+  TiImage image;
+  TiSampler sampler;
+  TiImageDimension dimension;
+  TiImageExtent extent;
+  TiFormat format;
+} TiTexture;
 
 // union.argument_value
 typedef union TiArgumentValue {
@@ -308,14 +344,20 @@ TI_DLL_EXPORT void *TI_API_CALL ti_map_memory(TiRuntime runtime,
 TI_DLL_EXPORT void TI_API_CALL ti_unmap_memory(TiRuntime runtime,
                                                TiMemory memory);
 
-// function.allocate_texture
-TI_DLL_EXPORT TiTexture TI_API_CALL
-ti_allocate_texture(TiRuntime runtime,
-                    const TiTextureAllocateInfo *allocate_info);
+// function.allocate_image
+TI_DLL_EXPORT TiImage TI_API_CALL
+ti_allocate_image(TiRuntime runtime, const TiImageAllocateInfo *allocate_info);
+
+// function.free_image
+TI_DLL_EXPORT void TI_API_CALL ti_free_image(TiRuntime runtime, TiImage image);
+
+// function.create_sampler
+TI_DLL_EXPORT TiSampler TI_API_CALL
+ti_create_sampler(TiRuntime runtime, const TiSamplerCreateInfo *create_info);
 
-// function.free_texture
-TI_DLL_EXPORT void TI_API_CALL ti_free_texture(TiRuntime runtime,
-                                               TiTexture texture);
+// function.destroy_sampler
+TI_DLL_EXPORT void TI_API_CALL ti_destroy_sampler(TiRuntime runtime,
+                                                  TiSampler sampler);
 
 // function.create_event
 TI_DLL_EXPORT TiEvent TI_API_CALL ti_create_event(TiRuntime runtime);
@@ -329,16 +371,16 @@ ti_copy_memory_device_to_device(TiRuntime runtime,
                                 const TiMemorySlice *dst_memory,
                                 const TiMemorySlice *src_memory);
 
-// function.copy_texture_device_to_device
+// function.copy_image_device_to_device
 TI_DLL_EXPORT void TI_API_CALL
-ti_copy_texture_device_to_device(TiRuntime runtime,
-                                 const TiTextureSlice *dst_texture,
-                                 const TiTextureSlice *src_texture);
-
-// function.transition_texture
-TI_DLL_EXPORT void TI_API_CALL ti_transition_texture(TiRuntime runtime,
-                                                     TiTexture texture,
-                                                     TiTextureLayout layout);
+ti_copy_image_device_to_device(TiRuntime runtime,
+                               const TiImageSlice *dst_image,
+                               const TiImageSlice *src_image);
+
+// function.transition_image
+TI_DLL_EXPORT void TI_API_CALL ti_transition_image(TiRuntime runtime,
+                                                   TiImage image,
+                                                   TiImageLayout layout);
 
 // function.launch_kernel
 TI_DLL_EXPORT void TI_API_CALL ti_launch_kernel(TiRuntime runtime,
diff --git a/c_api/include/taichi/taichi_vulkan.h b/c_api/include/taichi/taichi_vulkan.h
index 373e2d58527b1..fc1e58ce58d3d 100644
--- a/c_api/include/taichi/taichi_vulkan.h
+++ b/c_api/include/taichi/taichi_vulkan.h
@@ -25,8 +25,8 @@ typedef struct TiVulkanMemoryInteropInfo {
   VkBufferUsageFlags usage;
 } TiVulkanMemoryInteropInfo;
 
-// structure.vulkan_texture_interop_info
-typedef struct TiVulkanTextureInteropInfo {
+// structure.vulkan_image_interop_info
+typedef struct TiVulkanImageInteropInfo {
   VkImage image;
   VkImageType image_type;
   VkFormat format;
@@ -36,7 +36,7 @@ typedef struct TiVulkanTextureInteropInfo {
   VkSampleCountFlagBits sample_count;
   VkImageTiling tiling;
   VkImageUsageFlags usage;
-} TiVulkanTextureInteropInfo;
+} TiVulkanImageInteropInfo;
 
 // structure.vulkan_event_interop_info
 typedef struct TiVulkanEventInteropInfo {
@@ -71,18 +71,18 @@ ti_export_vulkan_memory(TiRuntime runtime,
                         TiMemory memory,
                         TiVulkanMemoryInteropInfo *interop_info);
 
-// function.import_vulkan_texture
-TI_DLL_EXPORT TiTexture TI_API_CALL
-ti_import_vulkan_texture(TiRuntime runtime,
-                         const TiVulkanTextureInteropInfo *interop_info,
-                         VkImageViewType view_type,
-                         VkImageLayout layout);
+// function.import_vulkan_image
+TI_DLL_EXPORT TiImage TI_API_CALL
+ti_import_vulkan_image(TiRuntime runtime,
+                       const TiVulkanImageInteropInfo *interop_info,
+                       VkImageViewType view_type,
+                       VkImageLayout layout);
 
-// function.export_vulkan_texture
+// function.export_vulkan_image
 TI_DLL_EXPORT void TI_API_CALL
-ti_export_vulkan_texture(TiRuntime runtime,
-                         TiTexture texture,
-                         TiVulkanTextureInteropInfo *interop_info);
+ti_export_vulkan_image(TiRuntime runtime,
+                       TiImage image,
+                       TiVulkanImageInteropInfo *interop_info);
 
 // function.import_vulkan_event
 TI_DLL_EXPORT TiEvent TI_API_CALL
diff --git a/c_api/src/taichi_core_impl.cpp b/c_api/src/taichi_core_impl.cpp
index 0d2ad61744ab0..4002af00376f1 100644
--- a/c_api/src/taichi_core_impl.cpp
+++ b/c_api/src/taichi_core_impl.cpp
@@ -223,8 +223,8 @@ void ti_unmap_memory(TiRuntime runtime, TiMemory devmem) {
   runtime2->get().unmap(devmem2devalloc(*runtime2, devmem));
 }
 
-TiTexture ti_allocate_texture(TiRuntime runtime,
-                              const TiTextureAllocateInfo *allocate_info) {
+TiImage ti_allocate_image(TiRuntime runtime,
+                          const TiImageAllocateInfo *allocate_info) {
   TI_CAPI_ARGUMENT_NULL_RV(runtime);
   TI_CAPI_ARGUMENT_NULL_RV(allocate_info);
 
@@ -232,13 +232,13 @@ TiTexture ti_allocate_texture(TiRuntime runtime,
   TI_CAPI_NOT_SUPPORTED_IF_RV(allocate_info->extent.array_layer_count > 1);
 
   taichi::lang::ImageAllocUsage usage{};
-  if (allocate_info->usage & TI_TEXTURE_USAGE_STORAGE_BIT) {
+  if (allocate_info->usage & TI_IMAGE_USAGE_STORAGE_BIT) {
     usage = usage | taichi::lang::ImageAllocUsage::Storage;
   }
-  if (allocate_info->usage & TI_TEXTURE_USAGE_SAMPLED_BIT) {
+  if (allocate_info->usage & TI_IMAGE_USAGE_SAMPLED_BIT) {
     usage = usage | taichi::lang::ImageAllocUsage::Sampled;
   }
-  if (allocate_info->usage & TI_TEXTURE_USAGE_ATTACHMENT_BIT) {
+  if (allocate_info->usage & TI_IMAGE_USAGE_ATTACHMENT_BIT) {
     usage = usage | taichi::lang::ImageAllocUsage::Attachment;
   }
 
@@ -275,14 +275,23 @@ TiTexture ti_allocate_texture(TiRuntime runtime,
   params.export_sharing = false;
   params.usage = usage;
 
-  TiTexture devtex = ((Runtime *)runtime)->allocate_texture(params);
-  return devtex;
+  TiImage devimg = ((Runtime *)runtime)->allocate_image(params);
+  return devimg;
 }
-void ti_free_texture(TiRuntime runtime, TiTexture texture) {
+void ti_free_image(TiRuntime runtime, TiImage image) {
   TI_CAPI_ARGUMENT_NULL(runtime);
-  TI_CAPI_ARGUMENT_NULL(texture);
+  TI_CAPI_ARGUMENT_NULL(image);
+
+  ((Runtime *)runtime)->free_image(image);
+}
 
-  ((Runtime *)runtime)->free_texture(texture);
+TiSampler ti_create_sampler(TiRuntime runtime,
+                            const TiSamplerCreateInfo *create_info) {
+  TI_CAPI_NOT_SUPPORTED(ti_create_sampler);
+  return TI_NULL_HANDLE;
+}
+void ti_destroy_sampler(TiRuntime runtime, TiSampler sampler) {
+  TI_CAPI_NOT_SUPPORTED(ti_destroy_sampler);
 }
 
 TiEvent ti_create_event(TiRuntime runtime) {
@@ -319,13 +328,13 @@ void ti_copy_memory_device_to_device(TiRuntime runtime,
 }
 
 void ti_copy_texture_device_to_device(TiRuntime runtime,
-                                      const TiTextureSlice *dst_texture,
-                                      const TiTextureSlice *src_texture) {
+                                      const TiImageSlice *dst_texture,
+                                      const TiImageSlice *src_texture) {
   TI_CAPI_ARGUMENT_NULL(runtime);
   TI_CAPI_ARGUMENT_NULL(dst_texture);
-  TI_CAPI_ARGUMENT_NULL(dst_texture->texture);
+  TI_CAPI_ARGUMENT_NULL(dst_texture->image);
   TI_CAPI_ARGUMENT_NULL(src_texture);
-  TI_CAPI_ARGUMENT_NULL(src_texture->texture);
+  TI_CAPI_ARGUMENT_NULL(src_texture->image);
   TI_CAPI_INVALID_ARGUMENT(src_texture->extent.width !=
                            dst_texture->extent.width);
   TI_CAPI_INVALID_ARGUMENT(src_texture->extent.height !=
@@ -336,8 +345,8 @@ void ti_copy_texture_device_to_device(TiRuntime runtime,
                            dst_texture->extent.array_layer_count);
 
   Runtime *runtime2 = (Runtime *)runtime;
-  auto dst = devtex2devalloc(*runtime2, dst_texture->texture);
-  auto src = devtex2devalloc(*runtime2, src_texture->texture);
+  auto dst = devimg2devalloc(*runtime2, dst_texture->image);
+  auto src = devimg2devalloc(*runtime2, src_texture->image);
 
   taichi::lang::ImageCopyParams params{};
   params.width = dst_texture->extent.width;
@@ -346,13 +355,13 @@ void ti_copy_texture_device_to_device(TiRuntime runtime,
   runtime2->copy_image(dst, src, params);
 }
 void ti_transition_texture(TiRuntime runtime,
-                           TiTexture texture,
-                           TiTextureLayout layout) {
+                           TiImage texture,
+                           TiImageLayout layout) {
   TI_CAPI_ARGUMENT_NULL(runtime);
   TI_CAPI_ARGUMENT_NULL(texture);
 
   Runtime *runtime2 = (Runtime *)runtime;
-  auto image = devtex2devalloc(*runtime2, texture);
+  auto image = devimg2devalloc(*runtime2, texture);
   auto layout2 = (taichi::lang::ImageLayout)layout;
 
   switch ((taichi::lang::ImageLayout)layout) {
@@ -461,6 +470,10 @@ void ti_launch_kernel(TiRuntime runtime,
         devallocs.emplace_back(std::move(devalloc));
         break;
       }
+      case TI_ARGUMENT_TYPE_TEXTURE: {
+        ti_set_last_error(TI_ERROR_NOT_SUPPORTED, "TI_ARGUMENT_TYPE_TEXTURE");
+        break;
+      }
       default: {
         ti_set_last_error(TI_ERROR_ARGUMENT_OUT_OF_RANGE,
                           ("args[" + std::to_string(i) + "].type").c_str());
@@ -485,6 +498,8 @@ void ti_launch_compute_graph(TiRuntime runtime,
   std::unordered_map<std::string, taichi::lang::aot::IValue> arg_map{};
   std::vector<taichi::lang::Ndarray> ndarrays;
   ndarrays.reserve(arg_count);
+  std::vector<taichi::lang::Texture> textures;
+  textures.reserve(arg_count);
 
   for (uint32_t i = 0; i < arg_count; ++i) {
     TI_CAPI_ARGUMENT_NULL(args[i].name);
@@ -574,6 +589,23 @@ void ti_launch_compute_graph(TiRuntime runtime,
             arg.name, taichi::lang::aot::IValue::create(ndarrays.back())));
         break;
       }
+      case TI_ARGUMENT_TYPE_TEXTURE: {
+        TI_CAPI_ARGUMENT_NULL(args[i].argument.value.texture.image);
+
+        taichi::lang::DeviceAllocation devalloc =
+            devimg2devalloc(runtime2, arg.argument.value.texture.image);
+        taichi::lang::BufferFormat format =
+            (taichi::lang::BufferFormat)arg.argument.value.texture.format;
+        uint32_t width = arg.argument.value.texture.extent.width;
+        uint32_t height = arg.argument.value.texture.extent.height;
+        uint32_t depth = arg.argument.value.texture.extent.depth;
+
+        textures.emplace_back(
+            taichi::lang::Texture(devalloc, format, width, height, depth));
+        arg_map.emplace(std::make_pair(
+            arg.name, taichi::lang::aot::IValue::create(textures.back())));
+        break;
+      }
       default: {
         ti_set_last_error(
             TI_ERROR_ARGUMENT_OUT_OF_RANGE,
diff --git a/c_api/src/taichi_core_impl.h b/c_api/src/taichi_core_impl.h
index a6445f432cad6..bc96cc95ba61f 100644
--- a/c_api/src/taichi_core_impl.h
+++ b/c_api/src/taichi_core_impl.h
@@ -32,15 +32,15 @@
     return TI_NULL_HANDLE;                         \
   }
 
-#define TI_CAPI_INVALID_ARGUMENT(x)                   \
-  if (x == TI_NULL_HANDLE) {                          \
-    ti_set_last_error(TI_ERROR_INVALID_ARGUMENT, #x); \
-    return;                                           \
+#define TI_CAPI_INVALID_ARGUMENT(pred)                   \
+  if (pred) {                                            \
+    ti_set_last_error(TI_ERROR_INVALID_ARGUMENT, #pred); \
+    return;                                              \
   }
-#define TI_CAPI_INVALID_ARGUMENT_RV(x)                \
-  if (x == TI_NULL_HANDLE) {                          \
-    ti_set_last_error(TI_ERROR_INVALID_ARGUMENT, #x); \
-    return TI_NULL_HANDLE;                            \
+#define TI_CAPI_INVALID_ARGUMENT_RV(pred)                \
+  if (pred) {                                            \
+    ti_set_last_error(TI_ERROR_INVALID_ARGUMENT, #pred); \
+    return TI_NULL_HANDLE;                               \
   }
 
 #define TI_CAPI_INVALID_INTEROP_ARCH(x, arch)                    \
@@ -78,10 +78,10 @@ class Runtime {
       const taichi::lang::Device::AllocParams &params);
   virtual void free_memory(TiMemory devmem);
 
-  virtual TiTexture allocate_texture(const taichi::lang::ImageParams &params) {
+  virtual TiImage allocate_image(const taichi::lang::ImageParams &params) {
     TI_NOT_IMPLEMENTED
   }
-  virtual void free_texture(TiTexture texture) {
+  virtual void free_image(TiImage image) {
     TI_NOT_IMPLEMENTED
   }
 
@@ -168,16 +168,16 @@ struct devalloc_cast_t {
   return devalloc_cast_t<TiMemory>::devalloc2handle(runtime, devalloc);
 }
 
-[[maybe_unused]] taichi::lang::DeviceAllocation devtex2devalloc(
+[[maybe_unused]] taichi::lang::DeviceAllocation devimg2devalloc(
     Runtime &runtime,
-    TiTexture devtex) {
-  return devalloc_cast_t<TiTexture>::handle2devalloc(runtime, devtex);
+    TiImage devimg) {
+  return devalloc_cast_t<TiImage>::handle2devalloc(runtime, devimg);
 }
 
-[[maybe_unused]] TiTexture devalloc2devtex(
+[[maybe_unused]] TiImage devalloc2devimg(
     Runtime &runtime,
     const taichi::lang::DeviceAllocation &devalloc) {
-  return devalloc_cast_t<TiTexture>::devalloc2handle(runtime, devalloc);
+  return devalloc_cast_t<TiImage>::devalloc2handle(runtime, devalloc);
 }
 
 }  // namespace
diff --git a/c_api/src/taichi_vulkan_impl.cpp b/c_api/src/taichi_vulkan_impl.cpp
index ae1eb9c2479b5..4081bcf5d1ad3 100644
--- a/c_api/src/taichi_vulkan_impl.cpp
+++ b/c_api/src/taichi_vulkan_impl.cpp
@@ -105,13 +105,13 @@ taichi::lang::gfx::GfxRuntime &VulkanRuntimeOwned::get_gfx_runtime() {
   return gfx_runtime_;
 }
 
-TiTexture VulkanRuntime::allocate_texture(
-    const taichi::lang::ImageParams &params) {
-  taichi::lang::DeviceAllocation devalloc = get_vk().create_image(params);
-  return devalloc2devtex(*this, devalloc);
+TiImage VulkanRuntime::allocate_image(const taichi::lang::ImageParams &params) {
+  taichi::lang::DeviceAllocation devalloc =
+      get_gfx_runtime().create_image(params);
+  return devalloc2devimg(*this, devalloc);
 }
-void VulkanRuntime::free_texture(TiTexture texture) {
-  get_vk().destroy_image(devtex2devalloc(*this, texture));
+void VulkanRuntime::free_image(TiImage image) {
+  get_vk().destroy_image(devimg2devalloc(*this, image));
 }
 
 // -----------------------------------------------------------------------------
@@ -221,11 +221,10 @@ void ti_export_vulkan_memory(TiRuntime runtime,
   interop_info->size = buffer.get()->size;
   interop_info->usage = buffer.get()->usage;
 }
-TiTexture ti_import_vulkan_texture(
-    TiRuntime runtime,
-    const TiVulkanTextureInteropInfo *interop_info,
-    VkImageViewType view_type,
-    VkImageLayout layout) {
+TiImage ti_import_vulkan_image(TiRuntime runtime,
+                               const TiVulkanImageInteropInfo *interop_info,
+                               VkImageViewType view_type,
+                               VkImageLayout layout) {
   TI_CAPI_ARGUMENT_NULL_RV(runtime);
   TI_CAPI_ARGUMENT_NULL_RV(interop_info);
   TI_CAPI_ARGUMENT_NULL_RV(interop_info->image);
@@ -269,32 +268,32 @@ TiTexture ti_import_vulkan_texture(
 
   taichi::lang::DeviceAllocation image2 =
       vk_runtime.import_vk_image(image, image_view, layout);
-  return devalloc2devtex(*runtime2, image2);
+  return devalloc2devimg(*runtime2, image2);
 }
 
-void ti_export_vulkan_texture(TiRuntime runtime,
-                              TiTexture texture,
-                              TiVulkanTextureInteropInfo *interop_info) {
+void ti_export_vulkan_image(TiRuntime runtime,
+                            TiImage image,
+                            TiVulkanImageInteropInfo *interop_info) {
   TI_CAPI_ARGUMENT_NULL(runtime);
-  TI_CAPI_ARGUMENT_NULL(texture);
+  TI_CAPI_ARGUMENT_NULL(image);
   TI_CAPI_ARGUMENT_NULL(interop_info);
   TI_CAPI_INVALID_INTEROP_ARCH(((Runtime *)runtime)->arch, vulkan);
 
   VulkanRuntime *runtime2 = ((Runtime *)runtime)->as_vk();
 
-  taichi::lang::DeviceAllocation devalloc = devtex2devalloc(*runtime2, texture);
-  vkapi::IVkImage image =
+  taichi::lang::DeviceAllocation devalloc = devimg2devalloc(*runtime2, image);
+  vkapi::IVkImage image2 =
       std::get<0>(runtime2->get_vk().get_vk_image(devalloc));
-  interop_info->image = image->image;
-  interop_info->image_type = image->type;
-  interop_info->extent.width = image->width;
-  interop_info->extent.height = image->height;
-  interop_info->extent.depth = image->depth;
-  interop_info->mip_level_count = image->mip_levels;
-  interop_info->array_layer_count = image->array_layers;
+  interop_info->image = image2->image;
+  interop_info->image_type = image2->type;
+  interop_info->extent.width = image2->width;
+  interop_info->extent.height = image2->height;
+  interop_info->extent.depth = image2->depth;
+  interop_info->mip_level_count = image2->mip_levels;
+  interop_info->array_layer_count = image2->array_layers;
   interop_info->sample_count = VK_SAMPLE_COUNT_1_BIT;
   interop_info->tiling = VK_IMAGE_TILING_OPTIMAL;
-  interop_info->usage = image->usage;
+  interop_info->usage = image2->usage;
 }
 
 TiEvent ti_import_vulkan_event(TiRuntime runtime,
diff --git a/c_api/src/taichi_vulkan_impl.h b/c_api/src/taichi_vulkan_impl.h
index 67ee479f30e05..25a149abf40e2 100644
--- a/c_api/src/taichi_vulkan_impl.h
+++ b/c_api/src/taichi_vulkan_impl.h
@@ -18,9 +18,9 @@ class VulkanRuntime : public GfxRuntime {
   VulkanRuntime();
 
   taichi::lang::vulkan::VulkanDevice &get_vk();
-  virtual TiTexture allocate_texture(
+  virtual TiImage allocate_image(
       const taichi::lang::ImageParams &params) override final;
-  virtual void free_texture(TiTexture texture) override final;
+  virtual void free_image(TiImage image) override final;
 };
 class VulkanRuntimeImported : public VulkanRuntime {
   // A dirty workaround to ensure the device is fully initialized before
diff --git a/c_api/taichi.json b/c_api/taichi.json
index 50130408f0fe5..d352b122cdf9b 100644
--- a/c_api/taichi.json
+++ b/c_api/taichi.json
@@ -61,7 +61,12 @@
                     "is_dispatchable": false
                 },
                 {
-                    "name": "texture",
+                    "name": "image",
+                    "type": "handle",
+                    "is_dispatchable": false
+                },
+                {
+                    "name": "sampler",
                     "type": "handle",
                     "is_dispatchable": false
                 },
@@ -94,12 +99,12 @@
                 {
                     "name": "arch",
                     "type": "enumeration",
-                    "inc_cases": "archs"
+                    "inc_cases": "PER_ARCH"
                 },
                 {
                     "name": "data_type",
                     "type": "enumeration",
-                    "inc_cases": "data_type"
+                    "inc_cases": "PER_TYPE"
                 },
                 {
                     "name": "argument_type",
@@ -202,7 +207,7 @@
                     ]
                 },
                 {
-                    "name": "texture_usage",
+                    "name": "image_usage",
                     "type": "bit_field",
                     "bits": {
                         "storage": 0,
@@ -211,7 +216,7 @@
                     }
                 },
                 {
-                    "name": "texture_dimension",
+                    "name": "image_dimension",
                     "type": "enumeration",
                     "cases": {
                         "1d": 0,
@@ -223,17 +228,17 @@
                     }
                 },
                 {
-                    "name": "texture_layout",
+                    "name": "image_layout",
                     "type": "enumeration",
-                    "inc_cases": "image_layout"
+                    "inc_cases": "PER_IMAGE_LAYOUT"
                 },
                 {
-                    "name": "texture_format",
+                    "name": "format",
                     "type": "enumeration",
-                    "inc_cases": "buffer_format"
+                    "inc_cases": "PER_BUFFER_FORMAT"
                 },
                 {
-                    "name": "texture_offset",
+                    "name": "image_offset",
                     "type": "structure",
                     "fields": [
                         {
@@ -255,7 +260,7 @@
                     ]
                 },
                 {
-                    "name": "texture_extent",
+                    "name": "image_extent",
                     "type": "structure",
                     "fields": [
                         {
@@ -277,16 +282,16 @@
                     ]
                 },
                 {
-                    "name": "texture_allocate_info",
+                    "name": "image_allocate_info",
                     "type": "structure",
                     "fields": [
                         {
                             "name": "dimension",
-                            "type": "enumeration.texture_dimension"
+                            "type": "enumeration.image_dimension"
                         },
                         {
                             "name": "extent",
-                            "type": "structure.texture_extent"
+                            "type": "structure.image_extent"
                         },
                         {
                             "name": "mip_level_count",
@@ -294,28 +299,28 @@
                         },
                         {
                             "name": "format",
-                            "type": "enumeration.texture_format"
+                            "type": "enumeration.format"
                         },
                         {
                             "name": "usage",
-                            "type": "bit_field.texture_usage"
+                            "type": "bit_field.image_usage"
                         }
                     ]
                 },
                 {
-                    "name": "texture_slice",
+                    "name": "image_slice",
                     "type": "structure",
                     "fields": [
                         {
-                            "type": "handle.texture"
+                            "type": "handle.image"
                         },
                         {
                             "name": "offset",
-                            "type": "structure.texture_offset"
+                            "type": "structure.image_offset"
                         },
                         {
                             "name": "extent",
-                            "type": "structure.texture_extent"
+                            "type": "structure.image_extent"
                         },
                         {
                             "name": "mip_level",
@@ -323,6 +328,63 @@
                         }
                     ]
                 },
+                {
+                    "name": "filter",
+                    "type": "enumeration",
+                    "inc_cases": "PER_FILTER"
+                },
+                {
+                    "name": "address_mode",
+                    "type": "enumeration",
+                    "inc_cases": "PER_ADDRESS_MODE"
+                },
+                {
+                    "name": "sampler_create_info",
+                    "type": "structure",
+                    "fields": [
+                        {
+                            "name": "mag_filter",
+                            "type": "enumeration.filter"
+                        },
+                        {
+                            "name": "min_filter",
+                            "type": "enumeration.filter"
+                        },
+                        {
+                            "name": "address_mode",
+                            "type": "enumeration.address_mode"
+                        },
+                        {
+                            "name": "max_anisotropy",
+                            "type": "float"
+                        }
+                    ]
+                },
+                {
+                    "name": "texture",
+                    "type": "structure",
+                    "fields": [
+                        {
+                            "name": "image",
+                            "type": "handle.image"
+                        },
+                        {
+                            "type": "handle.sampler"
+                        },
+                        {
+                            "name": "dimension",
+                            "type": "enumeration.image_dimension"
+                        },
+                        {
+                            "name": "extent",
+                            "type": "structure.image_extent"
+                        },
+                        {
+                            "name": "format",
+                            "type": "enumeration.format"
+                        }
+                    ]
+                },
                 {
                     "name": "argument_value",
                     "type": "union",
@@ -341,7 +403,7 @@
                         },
                         {
                             "name": "texture",
-                            "type": "handle.texture"
+                            "type": "structure.texture"
                         }
                     ]
                 },
@@ -486,32 +548,62 @@
                     ]
                 },
                 {
-                    "name": "allocate_texture",
+                    "name": "allocate_image",
                     "type": "function",
                     "parameters": [
                         {
                             "name": "@return",
-                            "type": "handle.texture"
+                            "type": "handle.image"
                         },
                         {
                             "type": "handle.runtime"
                         },
                         {
                             "name": "allocate_info",
-                            "type": "structure.texture_allocate_info",
+                            "type": "structure.image_allocate_info",
+                            "by_ref": true
+                        }
+                    ]
+                },
+                {
+                    "name": "free_image",
+                    "type": "function",
+                    "parameters": [
+                        {
+                            "type": "handle.runtime"
+                        },
+                        {
+                            "type": "handle.image"
+                        }
+                    ]
+                },
+                {
+                    "name": "create_sampler",
+                    "type": "function",
+                    "parameters": [
+                        {
+                            "name": "@return",
+                            "type": "handle.sampler"
+                        },
+                        {
+                            "type": "handle.runtime"
+                        },
+                        {
+                            "name": "create_info",
+                            "type": "structure.sampler_create_info",
                             "by_ref": true
                         }
                     ]
                 },
                 {
-                    "name": "free_texture",
+                    "name": "destroy_sampler",
                     "type": "function",
                     "parameters": [
                         {
                             "type": "handle.runtime"
                         },
                         {
-                            "type": "handle.texture"
+                            "type": "handle.sampler"
                         }
                     ]
                 },
@@ -558,7 +650,7 @@
                     ]
                 },
                 {
-                    "name": "copy_texture_device_to_device",
+                    "name": "copy_image_device_to_device",
                     "type": "function",
                     "is_device_command": true,
                     "parameters": [
@@ -566,19 +658,19 @@
                             "type": "handle.runtime"
                         },
                         {
-                            "name": "dst_texture",
-                            "type": "structure.texture_slice",
+                            "name": "dst_image",
+                            "type": "structure.image_slice",
                             "by_ref": true
                         },
                         {
-                            "name": "src_texture",
-                            "type": "structure.texture_slice",
+                            "name": "src_image",
+                            "type": "structure.image_slice",
                             "by_ref": true
                         }
                     ]
                 },
                 {
-                    "name": "transition_texture",
+                    "name": "transition_image",
                     "type": "function",
                     "is_device_command": true,
                     "parameters": [
@@ -586,11 +678,11 @@
                             "type": "handle.runtime"
                         },
                         {
-                            "type": "handle.texture"
+                            "type": "handle.image"
                         },
                         {
                             "name": "layout",
-                            "type": "enumeration.texture_layout"
+                            "type": "enumeration.image_layout"
                         }
                     ]
                 },
@@ -821,7 +913,7 @@
                     ]
                 },
                 {
-                    "name": "vulkan_texture_interop_info",
+                    "name": "vulkan_image_interop_info",
                     "type": "structure",
                     "fields": [
                         {
@@ -968,19 +1060,19 @@
                     ]
                 },
                 {
-                    "name": "import_vulkan_texture",
+                    "name": "import_vulkan_image",
                     "type": "function",
                     "parameters": [
                         {
                             "name": "@return",
-                            "type": "handle.texture"
+                            "type": "handle.image"
                         },
                         {
                             "type": "handle.runtime"
                         },
                         {
                             "name": "interop_info",
-                            "type": "structure.vulkan_texture_interop_info",
+                            "type": "structure.vulkan_image_interop_info",
                             "by_ref": true
                         },
                         {
@@ -994,18 +1086,18 @@
                     ]
                 },
                 {
-                    "name": "export_vulkan_texture",
+                    "name": "export_vulkan_image",
                     "type": "function",
                     "parameters": [
                         {
                             "type": "handle.runtime"
                         },
                         {
-                            "type": "handle.texture"
+                            "type": "handle.image"
                         },
                         {
                             "name": "interop_info",
-                            "type": "structure.vulkan_texture_interop_info",
+                            "type": "structure.vulkan_image_interop_info",
                             "by_mut": true
                         }
                     ]
diff --git a/c_api/tests/c_api_cgraph_test.cpp b/c_api/tests/c_api_cgraph_test.cpp
index d3c349fbd1037..edffb76a5dbf3 100644
--- a/c_api/tests/c_api_cgraph_test.cpp
+++ b/c_api/tests/c_api_cgraph_test.cpp
@@ -37,6 +37,44 @@ void graph_aot_test(TiArch arch) {
   arr_array.unmap();
 }
 
+void texture_aot_test(TiArch arch) {
+  const uint32_t width = 128;
+  const uint32_t height = 128;
+
+  const auto folder_dir = getenv("TAICHI_AOT_FOLDER_PATH");
+
+  std::stringstream aot_mod_ss;
+  aot_mod_ss << folder_dir;
+
+  ti::Runtime runtime(arch);
+
+  ti::AotModule aot_mod = runtime.load_aot_module(aot_mod_ss.str());
+  ti::ComputeGraph run_graph = aot_mod.get_compute_graph("run_graph");
+
+  ti::Texture tex0 =
+      runtime.allocate_texture2d(width, height, TI_FORMAT_R32F, TI_NULL_HANDLE);
+  ti::Texture tex1 =
+      runtime.allocate_texture2d(width, height, TI_FORMAT_R32F, TI_NULL_HANDLE);
+  ti::NdArray<float> arr =
+      runtime.allocate_ndarray<float>({width, height}, {}, true);
+
+  run_graph["tex0"] = tex0;
+  run_graph["rw_tex0"] = tex0;
+  run_graph["tex1"] = tex1;
+  run_graph["rw_tex1"] = tex1;
+  run_graph["arr"] = arr;
+  run_graph.launch();
+  runtime.wait();
+
+  EXPECT_GE(ti_get_last_error(0, nullptr), TI_ERROR_SUCCESS);
+
+  std::vector<float> arr_data(128 * 128);
+  arr.read(arr_data);
+  for (auto x : arr_data) {
+    EXPECT_GT(x, 0.5);
+  }
+}
+
 TEST(CapiGraphTest, CpuGraph) {
   TiArch arch = TiArch::TI_ARCH_X64;
   graph_aot_test(arch);
@@ -55,6 +93,12 @@ TEST(CapiGraphTest, VulkanGraph) {
     graph_aot_test(arch);
   }
 }
+TEST(CapiGraphTest, VulkanTextureGraph) {
+  if (capi::utils::is_vulkan_available()) {
+    TiArch arch = TiArch::TI_ARCH_VULKAN;
+    texture_aot_test(arch);
+  }
+}
 
 TEST(CapiGraphTest, OpenglGraph) {
   if (capi::utils::is_opengl_available()) {
diff --git a/c_api/tests/c_api_interface_test.cpp b/c_api/tests/c_api_interface_test.cpp
index 55deb82a693ad..b4110686a72d7 100644
--- a/c_api/tests/c_api_interface_test.cpp
+++ b/c_api/tests/c_api_interface_test.cpp
@@ -61,6 +61,18 @@ TEST(CapiDryRun, MemoryAllocation) {
   }
 }
 
+TEST(CapiDryRun, ImageAllocation) {
+  if (capi::utils::is_vulkan_available()) {
+    {
+      // Vulkan Runtime
+      TiArch arch = TiArch::TI_ARCH_VULKAN;
+      ti::Runtime runtime(arch);
+      ti::Texture texture =
+          runtime.allocate_texture2d(4, 4, TI_FORMAT_RGBA8, TI_NULL_HANDLE);
+    }
+  }
+}
+
 TEST(CapiDryRun, VulkanAotModule) {
   if (capi::utils::is_vulkan_available()) {
     const auto folder_dir = getenv("TAICHI_AOT_FOLDER_PATH");
diff --git a/misc/generate_c_api.py b/misc/generate_c_api.py
index 49e3e26f92c8d..36379f16b1281 100644
--- a/misc/generate_c_api.py
+++ b/misc/generate_c_api.py
@@ -12,7 +12,7 @@ def get_type_name(x: EntryBase):
     elif ty in [Alias, Handle, Enumeration, Structure, Union]:
         return x.name.upper_camel_case
     elif ty in [BitField]:
-        return x.name.extend('flag_bits').upper_camel_case
+        return x.name.extend('flags').upper_camel_case
     else:
         raise RuntimeError(f"'{x.id}' is not a type")
 
@@ -58,13 +58,14 @@ def get_declr(x: EntryBase):
         return '\n'.join(out)
 
     elif ty is BitField:
-        out = ["typedef enum " + get_type_name(x) + " {"]
+        bit_type_name = x.name.extend('flag_bits').upper_camel_case
+        out = ["typedef enum " + bit_type_name + " {"]
         for name, value in x.bits.items():
             out += [
                 f"  {name.extend('bit').screaming_snake_case} = 1 << {value},"
             ]
-        out += ["} " + get_type_name(x) + ";"]
-        out += [f"typedef TiFlags {x.name.extend('flags').upper_camel_case};"]
+        out += ["} " + bit_type_name + ";"]
+        out += [f"typedef TiFlags {get_type_name(x)};"]
         return '\n'.join(out)
 
     elif ty is Structure:
diff --git a/misc/generate_unity_language_binding.py b/misc/generate_unity_language_binding.py
index 2917f0a715776..0e68f04c479af 100644
--- a/misc/generate_unity_language_binding.py
+++ b/misc/generate_unity_language_binding.py
@@ -165,9 +165,21 @@ def get_declr(x: EntryBase):
         c_function_param_perm = []
         function_param_perm = []
         for param in x.params:
-            if isinstance(param.type,
-                          BuiltInType) and (param.type.id == "const void*"
-                                            or param.type.id == "void*"):
+            if isinstance(param.type, BuiltInType) and (
+                    param.type.id == "char") and (param.count is not None):
+                if param.by_mut:
+                    c_function_param_perm += [[
+                        f"  [MarshalAs(UnmanagedType.LPArray)] [In, Out] byte[] {param.name}"
+                    ]]
+                    function_param_perm += [[f"  byte[] {param.name}"]]
+                else:
+                    c_function_param_perm += [[
+                        f"  [MarshalAs(UnmanagedType.LPArray)] byte[] {param.name}"
+                    ]]
+                    function_param_perm += [[f"  byte[] {param.name}"]]
+            elif isinstance(param.type,
+                            BuiltInType) and (param.type.id == "const void*"
+                                              or param.type.id == "void*"):
                 perm = [
                     "byte", "sbyte", "short", "ushort", "int", "uint", "long",
                     "ulong", "IntPtr", "float", "double"
@@ -293,6 +305,7 @@ def generate_module_header(module):
         BuiltInType("int64_t", "long"),
         BuiltInType("uint64_t", "ulong"),
         BuiltInType("float", "float"),
+        BuiltInType("char", "byte"),
         BuiltInType("const char*", "string"),
         BuiltInType("void*", "IntPtr"),
         BuiltInType("const void*", "IntPtr"),
diff --git a/misc/taichi_json.py b/misc/taichi_json.py
index c1c034c63a1db..8550e2d861c6a 100644
--- a/misc/taichi_json.py
+++ b/misc/taichi_json.py
@@ -1,5 +1,7 @@
+import glob
 import json
 import re
+from collections import defaultdict
 
 
 class Name:
@@ -71,15 +73,20 @@ def set_current(declr_reg):
         DeclarationRegistry.current = declr_reg
 
 
-def load_inc_enums(name, inc_file_name):
-    path = "taichi/inc/" + inc_file_name + ".inc.h"
-    cases = {}
-    with open(path) as f:
-        for line in f.readlines():
-            m = re.match(r"\w+\((\w+)\).*", line)
-            if m:
-                case_name = name.extend(m[1])
-                cases[case_name] = len(cases)
+def load_inc_enums(name):
+    paths = glob.glob("taichi/inc/*.inc.h")
+    cases = defaultdict(dict)
+    for path in paths:
+        with open(path) as f:
+            for line in f.readlines():
+                m = re.match(r"(\w+)\((\w+)\).*", line)
+                if m:
+                    key = m[1]
+                    try:
+                        case_name = name.extend(m[2])
+                    except AssertionError:
+                        continue
+                    cases[key][case_name] = len(cases[key])
     return cases
 
 
@@ -143,7 +150,7 @@ class Enumeration(EntryBase):
     def __init__(self, j):
         super().__init__(j, "enumeration")
         if "inc_cases" in j:
-            self.cases = load_inc_enums(self.name, j["inc_cases"])
+            self.cases = load_inc_enums(self.name)[j["inc_cases"]]
         else:
             self.cases = dict((self.name.extend(name), value)
                               for name, value in j["cases"].items())
@@ -153,7 +160,7 @@ class BitField(EntryBase):
     def __init__(self, j):
         super().__init__(j, "bit_field")
         if "inc_cases" in j:
-            self.bits = load_inc_enums(self.name, j["inc_bits"])
+            self.bits = load_inc_enums(self.name)[j["inc_bits"]]
         else:
             self.bits = dict((self.name.extend(name), value)
                              for name, value in j["bits"].items())
diff --git a/taichi/inc/address_mode.inc.h b/taichi/inc/address_mode.inc.h
new file mode 100644
index 0000000000000..e38427693706c
--- /dev/null
+++ b/taichi/inc/address_mode.inc.h
@@ -0,0 +1,3 @@
+PER_ADDRESS_MODE(repeat)
+PER_ADDRESS_MODE(mirrored_repeat)
+PER_ADDRESS_MODE(clamp_to_edge)
diff --git a/taichi/inc/filter.inc.h b/taichi/inc/filter.inc.h
new file mode 100644
index 0000000000000..93dbffad5d304
--- /dev/null
+++ b/taichi/inc/filter.inc.h
@@ -0,0 +1,2 @@
+PER_FILTER(nearest)
+PER_FILTER(linear)
diff --git a/taichi/rhi/vulkan/vulkan_api.cpp b/taichi/rhi/vulkan/vulkan_api.cpp
index cab10ca167e38..d60d9ec19560b 100644
--- a/taichi/rhi/vulkan/vulkan_api.cpp
+++ b/taichi/rhi/vulkan/vulkan_api.cpp
@@ -111,7 +111,8 @@ IVkEvent create_event(VkDevice device,
   info.pNext = pnext;
   info.flags = flags;
 
-  vkCreateEvent(device, &info, nullptr, &obj->event);
+  VkResult res = vkCreateEvent(device, &info, nullptr, &obj->event);
+  BAIL_ON_VK_BAD_RESULT(res, "failed to create event");
   return obj;
 }
 
@@ -126,7 +127,8 @@ IVkSemaphore create_semaphore(VkDevice device,
   info.pNext = pnext;
   info.flags = flags;
 
-  vkCreateSemaphore(device, &info, nullptr, &obj->semaphore);
+  VkResult res = vkCreateSemaphore(device, &info, nullptr, &obj->semaphore);
+  BAIL_ON_VK_BAD_RESULT(res, "failed to create semaphore");
   return obj;
 }
 
@@ -139,7 +141,8 @@ IVkFence create_fence(VkDevice device, VkFenceCreateFlags flags, void *pnext) {
   info.pNext = pnext;
   info.flags = flags;
 
-  vkCreateFence(device, &info, nullptr, &obj->fence);
+  VkResult res = vkCreateFence(device, &info, nullptr, &obj->fence);
+  BAIL_ON_VK_BAD_RESULT(res, "failed to create fence");
   return obj;
 }
 
@@ -149,7 +152,9 @@ IVkDescriptorSetLayout create_descriptor_set_layout(
   IVkDescriptorSetLayout obj =
       std::make_shared<DeviceObjVkDescriptorSetLayout>();
   obj->device = device;
-  vkCreateDescriptorSetLayout(device, create_info, nullptr, &obj->layout);
+  VkResult res =
+      vkCreateDescriptorSetLayout(device, create_info, nullptr, &obj->layout);
+  BAIL_ON_VK_BAD_RESULT(res, "failed to create descriptor set layout");
   return obj;
 }
 
@@ -158,7 +163,9 @@ IVkDescriptorPool create_descriptor_pool(
     VkDescriptorPoolCreateInfo *create_info) {
   IVkDescriptorPool obj = std::make_shared<DeviceObjVkDescriptorPool>();
   obj->device = device;
-  vkCreateDescriptorPool(device, create_info, nullptr, &obj->pool);
+  VkResult res =
+      vkCreateDescriptorPool(device, create_info, nullptr, &obj->pool);
+  BAIL_ON_VK_BAD_RESULT(res, "failed to create descriptor pool");
   return obj;
 }
 
@@ -198,7 +205,8 @@ IVkCommandPool create_command_pool(VkDevice device,
   info.flags = flags;
   info.queueFamilyIndex = queue_family_index;
 
-  vkCreateCommandPool(device, &info, nullptr, &obj->pool);
+  VkResult res = vkCreateCommandPool(device, &info, nullptr, &obj->pool);
+  BAIL_ON_VK_BAD_RESULT(res, "failed to create command pool");
 
   return obj;
 }
@@ -222,7 +230,8 @@ IVkCommandBuffer allocate_command_buffer(IVkCommandPool pool,
     info.level = level;
     info.commandBufferCount = 1;
 
-    vkAllocateCommandBuffers(pool->device, &info, &cmdbuf);
+    VkResult res = vkAllocateCommandBuffers(pool->device, &info, &cmdbuf);
+    BAIL_ON_VK_BAD_RESULT(res, "failed to allocate command buffer");
   }
 
   IVkCommandBuffer obj = std::make_shared<DeviceObjVkCommandBuffer>();
@@ -238,7 +247,9 @@ IVkRenderPass create_render_pass(VkDevice device,
                                  VkRenderPassCreateInfo *create_info) {
   IVkRenderPass obj = std::make_shared<DeviceObjVkRenderPass>();
   obj->device = device;
-  vkCreateRenderPass(device, create_info, nullptr, &obj->renderpass);
+  VkResult res =
+      vkCreateRenderPass(device, create_info, nullptr, &obj->renderpass);
+  BAIL_ON_VK_BAD_RESULT(res, "failed to create render pass");
   return obj;
 }
 
@@ -265,7 +276,8 @@ IVkPipelineLayout create_pipeline_layout(
   info.pushConstantRangeCount = push_constant_range_count;
   info.pPushConstantRanges = push_constant_ranges;
 
-  vkCreatePipelineLayout(device, &info, nullptr, &obj->layout);
+  VkResult res = vkCreatePipelineLayout(device, &info, nullptr, &obj->layout);
+  BAIL_ON_VK_BAD_RESULT(res, "failed to create pipeline layout");
 
   return obj;
 }
@@ -284,7 +296,8 @@ IVkPipelineCache create_pipeline_cache(VkDevice device,
   info.initialDataSize = initial_size;
   info.pInitialData = initial_data;
 
-  vkCreatePipelineCache(device, &info, nullptr, &obj->cache);
+  VkResult res = vkCreatePipelineCache(device, &info, nullptr, &obj->cache);
+  BAIL_ON_VK_BAD_RESULT(res, "failed to create pipeline cache");
 
   return obj;
 }
@@ -314,8 +327,10 @@ IVkPipeline create_compute_pipeline(VkDevice device,
     info.basePipelineIndex = 0;
   }
 
-  vkCreateComputePipelines(device, cache ? cache->cache : VK_NULL_HANDLE, 1,
-                           &info, nullptr, &obj->pipeline);
+  VkResult res =
+      vkCreateComputePipelines(device, cache ? cache->cache : VK_NULL_HANDLE, 1,
+                               &info, nullptr, &obj->pipeline);
+  BAIL_ON_VK_BAD_RESULT(res, "failed to create compute pipeline");
 
   return obj;
 }
@@ -343,8 +358,10 @@ IVkPipeline create_graphics_pipeline(VkDevice device,
     create_info->basePipelineIndex = 0;
   }
 
-  vkCreateGraphicsPipelines(device, cache ? cache->cache : VK_NULL_HANDLE, 1,
-                            create_info, nullptr, &obj->pipeline);
+  VkResult res =
+      vkCreateGraphicsPipelines(device, cache ? cache->cache : VK_NULL_HANDLE,
+                                1, create_info, nullptr, &obj->pipeline);
+  BAIL_ON_VK_BAD_RESULT(res, "failed to create graphics pipeline");
 
   return obj;
 }
@@ -378,9 +395,10 @@ IVkPipeline create_raytracing_pipeline(
           taichi::lang::vulkan::VulkanLoader::instance().get_instance(),
           "vkCreateRayTracingPipelinesKHR"));
 
-  create_raytracing_pipeline_khr(device, deferredOperation,
-                                 cache ? cache->cache : VK_NULL_HANDLE, 1,
-                                 create_info, nullptr, &obj->pipeline);
+  VkResult res = create_raytracing_pipeline_khr(
+      device, deferredOperation, cache ? cache->cache : VK_NULL_HANDLE, 1,
+      create_info, nullptr, &obj->pipeline);
+  BAIL_ON_VK_BAD_RESULT(res, "failed to create raytracing pipeline");
 
   return obj;
 }
@@ -401,8 +419,9 @@ IVkImage create_image(VkDevice device,
   image->array_layers = image_info->arrayLayers;
   image->usage = alloc_info->usage;
 
-  vmaCreateImage(allocator, image_info, alloc_info, &image->image,
-                 &image->allocation, nullptr);
+  VkResult res = vmaCreateImage(allocator, image_info, alloc_info,
+                                &image->image, &image->allocation, nullptr);
+  BAIL_ON_VK_BAD_RESULT(res, "failed to create image");
 
   return image;
 }
@@ -475,7 +494,9 @@ IVkFramebuffer create_framebuffer(VkFramebufferCreateFlags flags,
   info.height = height;
   info.layers = layers;
 
-  vkCreateFramebuffer(renderpass->device, &info, nullptr, &obj->framebuffer);
+  VkResult res = vkCreateFramebuffer(renderpass->device, &info, nullptr,
+                                     &obj->framebuffer);
+  BAIL_ON_VK_BAD_RESULT(res, "failed to create framebuffer");
 
   return obj;
 }
@@ -490,8 +511,9 @@ IVkBuffer create_buffer(VkDevice device,
   buffer->size = buffer_info->size;
   buffer->usage = buffer_info->usage;
 
-  vmaCreateBuffer(allocator, buffer_info, alloc_info, &buffer->buffer,
-                  &buffer->allocation, nullptr);
+  VkResult res = vmaCreateBuffer(allocator, buffer_info, alloc_info,
+                                 &buffer->buffer, &buffer->allocation, nullptr);
+  BAIL_ON_VK_BAD_RESULT(res, "failed to create buffer");
 
   return buffer;
 }
@@ -530,7 +552,9 @@ IVkBufferView create_buffer_view(IVkBuffer buffer,
   info.offset = offset;
   info.range = range;
 
-  vkCreateBufferView(buffer->device, &info, nullptr, &view->view);
+  VkResult res =
+      vkCreateBufferView(buffer->device, &info, nullptr, &view->view);
+  BAIL_ON_VK_BAD_RESULT(res, "failed to create buffer view");
 
   return view;
 }
@@ -564,8 +588,9 @@ IVkAccelerationStructureKHR create_acceleration_structure(
           taichi::lang::vulkan::VulkanLoader::instance().get_instance(),
           "vkCreateAccelerationStructureKHR"));
 
-  create_acceleration_structure_khr(buffer->device, &info, nullptr,
-                                    &obj->accel);
+  VkResult res = create_acceleration_structure_khr(buffer->device, &info,
+                                                   nullptr, &obj->accel);
+  BAIL_ON_VK_BAD_RESULT(res, "failed to create acceleration structure");
 
   return obj;
 }
@@ -578,7 +603,9 @@ IVkQueryPool create_query_pool(VkDevice device) {
   info.queryType = VK_QUERY_TYPE_TIMESTAMP;
 
   VkQueryPool query_pool;
-  vkCreateQueryPool(device, &info, nullptr, &query_pool);
+  VkResult res = vkCreateQueryPool(device, &info, nullptr, &query_pool);
+  BAIL_ON_VK_BAD_RESULT(res, "failed to create query pool");
+
   IVkQueryPool obj = std::make_shared<DeviceObjVkQueryPool>();
   obj->device = device;
   obj->query_pool = query_pool;
diff --git a/tests/cpp/aot/python_scripts/texture_aot_test.py b/tests/cpp/aot/python_scripts/texture_aot_test.py
new file mode 100644
index 0000000000000..017f0e7d21b3c
--- /dev/null
+++ b/tests/cpp/aot/python_scripts/texture_aot_test.py
@@ -0,0 +1,90 @@
+import argparse
+import os
+
+import taichi as ti
+
+
+def compile_graph_aot(arch):
+    ti.init(arch=arch)
+
+    if ti.lang.impl.current_cfg().arch != arch:
+        return
+
+    @ti.kernel
+    def run0(rw_tex: ti.types.rw_texture(num_dimensions=2,
+                                         num_channels=1,
+                                         channel_format=ti.f32,
+                                         lod=0)):
+        for i, j in ti.ndrange(128, 128):
+            value = ti.cast((j * 129 + i) % 2, ti.f32)
+            rw_tex.store(ti.Vector([i, j]), ti.Vector([value, 0.0, 0.0, 0.0]))
+
+    @ti.kernel
+    def run1(tex: ti.types.texture(num_dimensions=2),
+             rw_tex: ti.types.rw_texture(num_dimensions=2,
+                                         num_channels=1,
+                                         channel_format=ti.f32,
+                                         lod=0)):
+        for i, j in ti.ndrange(128, 128):
+            value = tex.fetch(ti.Vector([i, j]), 0).x
+            rw_tex.store(ti.Vector([i, j]),
+                         ti.Vector([1.0 - value, 0.0, 0.0, 0.0]))
+
+    @ti.kernel
+    def run2(tex0: ti.types.texture(num_dimensions=2),
+             tex1: ti.types.texture(num_dimensions=2),
+             arr: ti.types.ndarray(field_dim=2)):
+        for i, j in arr:
+            value0 = tex0.fetch(ti.Vector([i, j]), 0)
+            value1 = tex1.fetch(ti.Vector([i, j]), 0)
+            arr[i, j] = value0.x + value1.x
+
+    _tex0 = ti.graph.Arg(ti.graph.ArgKind.TEXTURE,
+                         'tex0',
+                         channel_format=ti.f32,
+                         shape=(128, 128),
+                         num_channels=1)
+    _rw_tex0 = ti.graph.Arg(ti.graph.ArgKind.RWTEXTURE,
+                            'rw_tex0',
+                            channel_format=ti.f32,
+                            shape=(128, 128),
+                            num_channels=1)
+    _tex1 = ti.graph.Arg(ti.graph.ArgKind.TEXTURE,
+                         'tex1',
+                         channel_format=ti.f32,
+                         shape=(128, 128),
+                         num_channels=1)
+    _rw_tex1 = ti.graph.Arg(ti.graph.ArgKind.RWTEXTURE,
+                            'rw_tex1',
+                            channel_format=ti.f32,
+                            shape=(128, 128),
+                            num_channels=1)
+    _arr = ti.graph.Arg(ti.graph.ArgKind.NDARRAY,
+                        'arr',
+                        ti.f32,
+                        field_dim=2,
+                        element_shape=())
+
+    g_builder = ti.graph.GraphBuilder()
+    g_builder.dispatch(run0, _rw_tex0)
+    g_builder.dispatch(run1, _tex0, _rw_tex1)
+    g_builder.dispatch(run2, _tex0, _tex1, _arr)
+    run_graph = g_builder.compile()
+
+    assert "TAICHI_AOT_FOLDER_PATH" in os.environ.keys()
+    tmpdir = str(os.environ["TAICHI_AOT_FOLDER_PATH"])
+
+    mod = ti.aot.Module(arch)
+    mod.add_graph('run_graph', run_graph)
+    mod.save(tmpdir, '')
+
+
+if __name__ == "__main__":
+    parser = argparse.ArgumentParser()
+    parser.add_argument("--arch", type=str)
+    args = parser.parse_args()
+
+    if args.arch == "vulkan":
+        compile_graph_aot(arch=ti.vulkan)
+    else:
+        assert False
diff --git a/tests/test_utils.py b/tests/test_utils.py
index 641a421249db3..d0eecec78bc61 100644
--- a/tests/test_utils.py
+++ b/tests/test_utils.py
@@ -174,6 +174,10 @@
         os.path.join('cpp', 'aot', 'python_scripts', 'graph_aot_test.py'),
         "--arch=vulkan"
     ],
+    "CapiGraphTest.VulkanTextureGraph": [
+        os.path.join('cpp', 'aot', 'python_scripts', 'texture_aot_test.py'),
+        "--arch=vulkan"
+    ],
     "CapiGraphTest.OpenglGraph": [
         os.path.join('cpp', 'aot', 'python_scripts', 'graph_aot_test.py'),
         "--arch=opengl"