diff --git a/python_bindings/src/PyEnums.cpp b/python_bindings/src/PyEnums.cpp
index 1224b9659f62..cf45db94e813 100644
--- a/python_bindings/src/PyEnums.cpp
+++ b/python_bindings/src/PyEnums.cpp
@@ -36,7 +36,10 @@ void define_enums(py::module &m) {
         .value("Heap", MemoryType::Heap)
         .value("Stack", MemoryType::Stack)
         .value("Register", MemoryType::Register)
-        .value("GPUShared", MemoryType::GPUShared);
+        .value("GPUShared", MemoryType::GPUShared)
+        .value("GPUTexture", MemoryType::GPUTexture)
+        .value("LockedCache", MemoryType::LockedCache)
+        .value("VTCM", MemoryType::VTCM);
 
     py::enum_<NameMangling>(m, "NameMangling")
         .value("Default", NameMangling::Default)
diff --git a/python_bindings/src/PyImageParam.cpp b/python_bindings/src/PyImageParam.cpp
index 64f38ddfc74b..1a3e35f50a4d 100644
--- a/python_bindings/src/PyImageParam.cpp
+++ b/python_bindings/src/PyImageParam.cpp
@@ -31,6 +31,7 @@ void define_image_param(py::module &m) {
             .def("host_alignment", &OutputImageParam::host_alignment)
             .def("set_estimates", &OutputImageParam::set_estimates, py::arg("estimates"))
             .def("set_host_alignment", &OutputImageParam::set_host_alignment)
+            .def("store_in", &OutputImageParam::store_in, py::arg("memory_type"))
             .def("dimensions", &OutputImageParam::dimensions)
             .def("left", &OutputImageParam::left)
             .def("right", &OutputImageParam::right)
diff --git a/src/Closure.h b/src/Closure.h
index 4d8ebd59db59..858709f4205b 100644
--- a/src/Closure.h
+++ b/src/Closure.h
@@ -55,11 +55,14 @@ class Closure : public IRVisitor {
         /** The buffer is written to. */
         bool write;
 
+        /** The buffer is a texture */
+        bool texture;
+
         /** The size of the buffer if known, otherwise zero. */
         size_t size;
 
         Buffer()
-            : dimensions(0), read(false), write(false), size(0) {
+            : dimensions(0), read(false), write(false), texture(false), size(0) {
         }
     };
 
diff --git a/src/CodeGen_OpenCL_Dev.cpp b/src/CodeGen_OpenCL_Dev.cpp
index c5743785a5c2..0557bc0a0909 100644
--- a/src/CodeGen_OpenCL_Dev.cpp
+++ b/src/CodeGen_OpenCL_Dev.cpp
@@ -1,4 +1,5 @@
 #include <algorithm>
+#include <array>
 #include <sstream>
 #include <utility>
 
@@ -174,6 +175,21 @@ string CodeGen_OpenCL_Dev::CodeGen_OpenCL_C::get_memory_space(const string &buf)
     }
 }
 
+namespace {
+std::string image_type_suffix(const Type &type) {
+    if (type.is_int()) {
+        return "i";
+    } else if (type.is_uint()) {
+        return "ui";
+    } else if (type.is_float()) {
+        return "f";
+    } else {
+        internal_error << "Invalid type for image: " << type << "\n";
+    }
+    return "";
+}
+}  // namespace
+
 void CodeGen_OpenCL_Dev::CodeGen_OpenCL_C::visit(const Call *op) {
     if (op->is_intrinsic(Call::bool_to_mask)) {
         if (op->args[0].type().is_vector()) {
@@ -242,6 +258,107 @@ void CodeGen_OpenCL_Dev::CodeGen_OpenCL_C::visit(const Call *op) {
         } else {
             CodeGen_C::visit(op);
         }
+    } else if (op->is_intrinsic(Call::image_load_texture)) {
+        // image_load(<image name>, <buffer>, <x>, <x-extent>, <y>,
+        // <y-extent>, <z>, <z-extent>)
+        int dims = (op->args.size() - 2) / 2;
+        internal_assert(dims >= 1 && dims <= 3);
+        const StringImm *string_imm = op->args[0].as<StringImm>();
+        if (!string_imm) {
+            internal_assert(op->args[0].as<Broadcast>());
+            string_imm = op->args[0].as<Broadcast>()->value.as<StringImm>();
+        }
+        internal_assert(string_imm);
+        Type arg_type = op->args[2].type();
+        internal_assert(arg_type.lanes() <= 16);
+        internal_assert(arg_type.lanes() == op->type.lanes());
+
+        std::array<string, 3> coord;
+        for (int i = 0; i < dims; i++) {
+            coord[i] = print_expr(op->args[i * 2 + 2]);
+        }
+        vector<string> results(arg_type.lanes());
+        // For vectorized reads, codegen as a sequence of read_image calls
+        for (int i = 0; i < arg_type.lanes(); i++) {
+            ostringstream rhs;
+            rhs << "read_image" << image_type_suffix(op->type) << "(" << print_name(string_imm->value) << ", ";
+            string idx = arg_type.is_vector() ? string(".s") + vector_elements[i] : "";
+            switch (dims) {
+            case 1:
+                rhs << coord[0] << idx << ").s0";
+                break;
+            case 2:
+                rhs << "(int2)(" << coord[0] << idx << ", " << coord[1] << idx << ")).s0";
+                break;
+            case 3:
+                rhs << "(int4)(" << coord[0] << idx << ", " << coord[1] << idx
+                    << ", " << coord[2] << idx << ", 0)).s0";
+                break;
+            }
+            print_assignment(op->type.with_bits(32).with_lanes(1), rhs.str());
+            results[i] = id;
+        }
+
+        if (op->type.is_vector()) {
+            // Combine all results into a single vector
+            ostringstream rhs;
+            rhs << "(" << print_type(op->type) << ")(";
+            for (int i = 0; i < op->type.lanes(); i++) {
+                rhs << results[i];
+                if (i < op->type.lanes() - 1) {
+                    rhs << ", ";
+                }
+            }
+            rhs << ")";
+            print_assignment(op->type, rhs.str());
+        }
+        if (op->type.bits() != 32) {
+            // Widen to the correct type
+            print_assignment(op->type, "convert_" + print_type(op->type) + "(" + id + ")");
+        }
+    } else if (op->is_intrinsic(Call::image_store_texture)) {
+        // image_store(<image name>, <buffer>, <x>, <y>, <z>, <value>)
+        const StringImm *string_imm = op->args[0].as<StringImm>();
+        if (!string_imm) {
+            internal_assert(op->args[0].as<Broadcast>());
+            string_imm = op->args[0].as<Broadcast>()->value.as<StringImm>();
+        }
+        internal_assert(string_imm);
+        int dims = op->args.size() - 3;
+        internal_assert(dims >= 1 && dims <= 3);
+        Type arg_type = op->args[2].type();
+        internal_assert(arg_type.lanes() <= 16);
+        Type value_type = op->args.back().type();
+        internal_assert(arg_type.lanes() == value_type.lanes());
+
+        std::array<string, 3> coord;
+        for (int i = 0; i < dims; i++) {
+            coord[i] = print_expr(op->args[i + 2]);
+        }
+        string value = print_expr(op->args.back());
+        // For vectorized writes, codegen as a sequence of write_image calls
+        for (int i = 0; i < arg_type.lanes(); i++) {
+            ostringstream write_image;
+            write_image << "write_image" << image_type_suffix(op->type)
+                        << "(" << print_name(string_imm->value) << ", ";
+            string idx = arg_type.is_vector() ? string(".s") + vector_elements[i] : "";
+            switch (dims) {
+            case 1:
+                write_image << coord[0] << idx;
+                break;
+            case 2:
+                write_image << "(int2)(" << coord[0] << idx << ", " << coord[1] << idx << ")";
+                break;
+            case 3:
+                write_image << "(int4)(" << coord[0] << idx << ", " << coord[1] << idx
+                            << ", " << coord[2] << idx << ", 0)";
+                break;
+            }
+            write_image << ", (" << print_type(value_type.with_bits(32).with_lanes(4))
+                        << ")(" << value << idx << ", 0, 0, 0));\n";
+            //  do_indent();
+            stream << write_image.str();
+        }
     } else {
         CodeGen_C::visit(op);
     }
@@ -765,11 +882,27 @@ void CodeGen_OpenCL_Dev::CodeGen_OpenCL_C::add_kernel(Stmt s,
     stream << "__kernel void " << name << "(\n";
     for (size_t i = 0; i < args.size(); i++) {
         if (args[i].is_buffer) {
-            stream << " " << get_memory_space(args[i].name) << " ";
-            if (!args[i].write) stream << "const ";
-            stream << print_type(args[i].type) << " *"
-                   << "restrict "
-                   << print_name(args[i].name);
+            if (args[i].is_texture) {
+                int dims = args[i].dimensions;
+                internal_assert(dims >= 1 && dims <= 3) << "dims = " << dims << "\n";
+                if (args[i].read && args[i].write) {
+                    stream << " __read_write ";
+                } else if (args[i].read) {
+                    stream << " __read_only ";
+                } else if (args[i].write) {
+                    stream << " __write_only ";
+                } else {
+                    internal_error << "CL Image argument " << args[i].name
+                                   << " is neither read nor write";
+                }
+                stream << "image" << dims << "d_t ";
+            } else {
+                stream << " " << get_memory_space(args[i].name) << " ";
+                if (!args[i].write) stream << "const ";
+                stream << print_type(args[i].type) << " *"
+                       << "restrict ";
+            }
+            stream << print_name(args[i].name);
             Allocation alloc;
             alloc.type = args[i].type;
             allocations.push(args[i].name, alloc);
@@ -797,7 +930,7 @@ void CodeGen_OpenCL_Dev::CodeGen_OpenCL_C::add_kernel(Stmt s,
         void visit(const Allocate *op) override {
             if (op->memory_type == MemoryType::GPUShared) {
                 internal_assert(alloc == nullptr)
-                    << "Found multiple shared allocations in metal kernel\n";
+                    << "Found multiple shared allocations in opencl kernel\n";
                 alloc = op;
             }
         }
diff --git a/src/DeviceArgument.cpp b/src/DeviceArgument.cpp
index b5050eef6841..faefbf908013 100644
--- a/src/DeviceArgument.cpp
+++ b/src/DeviceArgument.cpp
@@ -15,15 +15,17 @@ std::vector<DeviceArgument> HostClosure::arguments() {
     std::vector<DeviceArgument> res;
     for (const auto &v : vars) {
         debug(2) << "var: " << v.first << "\n";
-        res.emplace_back(v.first, false, v.second, 0);
+        res.emplace_back(v.first, false, false, v.second, 0);
     }
     for (const auto &b : buffers) {
         debug(2) << "buffer: " << b.first << " " << b.second.size;
         if (b.second.read) debug(2) << " (read)";
         if (b.second.write) debug(2) << " (write)";
+        if (b.second.texture) debug(2) << " <texture>";
+        debug(2) << " dims=" << (int)b.second.dimensions;
         debug(2) << "\n";
 
-        DeviceArgument arg(b.first, true, b.second.type, b.second.dimensions, b.second.size);
+        DeviceArgument arg(b.first, true, b.second.texture, b.second.type, b.second.dimensions, b.second.size);
         arg.read = b.second.read;
         arg.write = b.second.write;
         res.push_back(arg);
@@ -34,8 +36,10 @@ std::vector<DeviceArgument> HostClosure::arguments() {
 void HostClosure::visit(const Call *op) {
     if (op->is_intrinsic(Call::glsl_texture_load) ||
         op->is_intrinsic(Call::image_load) ||
+        op->is_intrinsic(Call::image_load_texture) ||
         op->is_intrinsic(Call::glsl_texture_store) ||
-        op->is_intrinsic(Call::image_store)) {
+        op->is_intrinsic(Call::image_store) ||
+        op->is_intrinsic(Call::image_store_texture)) {
 
         // The argument to the call is either a StringImm or a broadcasted
         // StringImm if this is part of a vectorized expression
@@ -51,14 +55,19 @@ void HostClosure::visit(const Call *op) {
         std::string bufname = string_imm->value;
         Buffer &ref = buffers[bufname];
         ref.type = op->type;
-        // TODO: do we need to set ref.dimensions?
+        ref.texture = op->is_intrinsic(Call::image_load_texture) ||
+                      op->is_intrinsic(Call::image_store_texture);
 
         if (op->is_intrinsic(Call::glsl_texture_load) ||
-            op->is_intrinsic(Call::image_load)) {
+            op->is_intrinsic(Call::image_load) ||
+            op->is_intrinsic(Call::image_load_texture)) {
             ref.read = true;
+            ref.dimensions = (op->args.size() - 2) / 2;
         } else if (op->is_intrinsic(Call::glsl_texture_store) ||
-                   op->is_intrinsic(Call::image_store)) {
+                   op->is_intrinsic(Call::image_store) ||
+                   op->is_intrinsic(Call::image_store_texture)) {
             ref.write = true;
+            ref.dimensions = op->args.size() - 3;
         }
 
         // The Func's name and the associated .buffer are mentioned in the
diff --git a/src/DeviceArgument.h b/src/DeviceArgument.h
index f43d0ba9856c..8666650787d9 100644
--- a/src/DeviceArgument.h
+++ b/src/DeviceArgument.h
@@ -36,6 +36,12 @@ struct DeviceArgument {
      */
     bool is_buffer;
 
+    /** If is_buffer == true and is_texture == true, this argument should be
+     * passed and accessed through texture sampler operations instead of
+     * directly as a memory array
+     */
+    bool is_texture;
+
     /** If is_buffer is true, this is the dimensionality of the buffer.
      * If is_buffer is false, this value is ignored (and should always be set to zero) */
     uint8_t dimensions;
@@ -66,6 +72,7 @@ struct DeviceArgument {
 
     DeviceArgument()
         : is_buffer(false),
+          is_texture(false),
           dimensions(0),
           size(0),
           packed_index(0),
@@ -75,11 +82,13 @@ struct DeviceArgument {
 
     DeviceArgument(const std::string &_name,
                    bool _is_buffer,
+                   bool _is_texture,
                    Type _type,
                    uint8_t _dimensions,
                    size_t _size = 0)
         : name(_name),
           is_buffer(_is_buffer),
+          is_texture(_is_texture),
           dimensions(_dimensions),
           type(_type),
           size(_size),
diff --git a/src/DeviceInterface.cpp b/src/DeviceInterface.cpp
index 440b5ef9015a..643836843ba0 100644
--- a/src/DeviceInterface.cpp
+++ b/src/DeviceInterface.cpp
@@ -162,7 +162,7 @@ DeviceAPI get_default_device_api_for_target(const Target &target) {
 }
 
 namespace Internal {
-Expr make_device_interface_call(DeviceAPI device_api) {
+Expr make_device_interface_call(DeviceAPI device_api, bool texture) {
     if (device_api == DeviceAPI::Host) {
         return make_zero(type_of<const halide_device_interface_t *>());
     }
@@ -173,7 +173,11 @@ Expr make_device_interface_call(DeviceAPI device_api) {
         interface_name = "halide_cuda_device_interface";
         break;
     case DeviceAPI::OpenCL:
-        interface_name = "halide_opencl_device_interface";
+        if (texture) {
+            interface_name = "halide_opencl_image_device_interface";
+        } else {
+            interface_name = "halide_opencl_device_interface";
+        }
         break;
     case DeviceAPI::Metal:
         interface_name = "halide_metal_device_interface";
diff --git a/src/DeviceInterface.h b/src/DeviceInterface.h
index 1ba4c3773092..fb1c12028f6d 100644
--- a/src/DeviceInterface.h
+++ b/src/DeviceInterface.h
@@ -37,7 +37,7 @@ bool host_supports_target_device(const Target &t);
 
 namespace Internal {
 /** Get an Expr which evaluates to the device interface for the given device api at runtime. */
-Expr make_device_interface_call(DeviceAPI device_api);
+Expr make_device_interface_call(DeviceAPI device_api, bool texture = false);
 }  // namespace Internal
 
 }  // namespace Halide
diff --git a/src/Expr.h b/src/Expr.h
index 06ed94638b66..5f766bf6a68e 100644
--- a/src/Expr.h
+++ b/src/Expr.h
@@ -368,6 +368,10 @@ enum class MemoryType {
      * across GPU threads within the same block. */
     GPUShared,
 
+    /** Allocation is stored in GPU texture memory and accessed through
+     * hardware sampler */
+    GPUTexture,
+
     /** Allocate Locked Cache Memory to act as local memory */
     LockedCache,
     /** Vector Tightly Coupled Memory. HVX (Hexagon) local memory available on
diff --git a/src/FuseGPUThreadLoops.cpp b/src/FuseGPUThreadLoops.cpp
index 91b8e9a7361c..053b97660a1f 100644
--- a/src/FuseGPUThreadLoops.cpp
+++ b/src/FuseGPUThreadLoops.cpp
@@ -458,7 +458,8 @@ class ExtractSharedAndHeapAllocations : public IRMutator {
 
         if ((fixed_size_thread_allocation &&
              op->memory_type != MemoryType::Heap &&
-             op->memory_type != MemoryType::GPUShared) ||
+             op->memory_type != MemoryType::GPUShared &&
+             op->memory_type != MemoryType::GPUTexture) ||
             op->memory_type == MemoryType::Register ||
             op->memory_type == MemoryType::Stack) {
             // These allocations go in register or local memory
@@ -467,6 +468,7 @@ class ExtractSharedAndHeapAllocations : public IRMutator {
 
         user_assert(op->memory_type == MemoryType::Auto ||
                     op->memory_type == MemoryType::GPUShared ||
+                    op->memory_type == MemoryType::GPUTexture ||
                     op->memory_type == MemoryType::Heap)
             << "Allocation " << op->name << " must live in shared or heap memory, "
             << "but is scheduled to live in " << op->memory_type << " memory.\n";
@@ -1263,6 +1265,7 @@ class InjectThreadBarriers : public IRMutator {
             break;
         case MemoryType::Auto:
         case MemoryType::Heap:
+        case MemoryType::GPUTexture:
             debug(4) << "   memory type is heap or auto\n";
             device_stores.insert(op->name);
             break;
@@ -1286,6 +1289,7 @@ class InjectThreadBarriers : public IRMutator {
             break;
         case MemoryType::Auto:
         case MemoryType::Heap:
+        case MemoryType::GPUTexture:
             debug(4) << "   memory type is heap or auto\n";
             device_loads.insert(op->name);
             break;
diff --git a/src/Generator.cpp b/src/Generator.cpp
index 46a0b85d6301..daa8a067c6fb 100644
--- a/src/Generator.cpp
+++ b/src/Generator.cpp
@@ -747,7 +747,7 @@ std::string halide_type_to_c_type(const Type &t) {
 
 int generate_filter_main_inner(int argc, char **argv, std::ostream &cerr) {
     const char kUsage[] =
-        "gengen \n"
+        "gengen  \n"
         "  [-g GENERATOR_NAME] [-f FUNCTION_NAME] [-o OUTPUT_DIR] [-r RUNTIME_NAME] [-d 1|0]\n"
         "  [-e EMIT_OPTIONS] [-n FILE_BASE_NAME] [-p PLUGIN_NAME] [-s AUTOSCHEDULER_NAME]\n"
         "       target=target-string[,target-string...] [generator_arg=value [...]]\n"
diff --git a/src/Generator.h b/src/Generator.h
index 19a02cf0ebb5..2228dfaeac65 100644
--- a/src/Generator.h
+++ b/src/Generator.h
@@ -1785,6 +1785,7 @@ class GeneratorInput_Buffer : public GeneratorInputImpl<T, Func> {
     HALIDE_FORWARD_METHOD_CONST(ImageParam, dim)
     HALIDE_FORWARD_METHOD_CONST(ImageParam, host_alignment)
     HALIDE_FORWARD_METHOD(ImageParam, set_host_alignment)
+    HALIDE_FORWARD_METHOD(ImageParam, store_in)
     HALIDE_FORWARD_METHOD_CONST(ImageParam, dimensions)
     HALIDE_FORWARD_METHOD_CONST(ImageParam, left)
     HALIDE_FORWARD_METHOD_CONST(ImageParam, right)
@@ -2541,6 +2542,7 @@ class GeneratorOutput_Buffer : public GeneratorOutputImpl<T> {
     HALIDE_FORWARD_METHOD_CONST(OutputImageParam, dim)
     HALIDE_FORWARD_METHOD_CONST(OutputImageParam, host_alignment)
     HALIDE_FORWARD_METHOD(OutputImageParam, set_host_alignment)
+    HALIDE_FORWARD_METHOD(OutputImageParam, store_in)
     HALIDE_FORWARD_METHOD_CONST(OutputImageParam, dimensions)
     HALIDE_FORWARD_METHOD_CONST(OutputImageParam, left)
     HALIDE_FORWARD_METHOD_CONST(OutputImageParam, right)
diff --git a/src/IR.cpp b/src/IR.cpp
index 4505a6fe3fe7..bf44dc45b740 100644
--- a/src/IR.cpp
+++ b/src/IR.cpp
@@ -608,7 +608,9 @@ const char *const intrinsic_op_names[] = {
     "if_then_else",
     "if_then_else_mask",
     "image_load",
+    "image_load_texture",
     "image_store",
+    "image_store_texture",
     "lerp",
     "likely",
     "likely_if_innermost",
diff --git a/src/IR.h b/src/IR.h
index 40f5d6410c88..69c232968f2e 100644
--- a/src/IR.h
+++ b/src/IR.h
@@ -520,7 +520,9 @@ struct Call : public ExprNode<Call> {
         if_then_else,
         if_then_else_mask,
         image_load,
+        image_load_texture,
         image_store,
+        image_store_texture,
         lerp,
         likely,
         likely_if_innermost,
diff --git a/src/IRPrinter.cpp b/src/IRPrinter.cpp
index 953ba044883b..ccab2bffc59a 100644
--- a/src/IRPrinter.cpp
+++ b/src/IRPrinter.cpp
@@ -129,6 +129,9 @@ std::ostream &operator<<(std::ostream &out, const MemoryType &t) {
     case MemoryType::GPUShared:
         out << "GPUShared";
         break;
+    case MemoryType::GPUTexture:
+        out << "GPUTexture";
+        break;
     case MemoryType::LockedCache:
         out << "LockedCache";
         break;
diff --git a/src/InjectHostDevBufferCopies.cpp b/src/InjectHostDevBufferCopies.cpp
index 5827abfaf4bf..b745859fd9c3 100644
--- a/src/InjectHostDevBufferCopies.cpp
+++ b/src/InjectHostDevBufferCopies.cpp
@@ -52,20 +52,25 @@ class FindBufferUsage : public IRVisitor {
     }
 
     void visit(const Call *op) override {
-        if (op->is_intrinsic(Call::image_load)) {
+        if (op->is_intrinsic(Call::image_load) ||
+            op->is_intrinsic(Call::image_load_texture)) {
             internal_assert(!op->args.empty());
             if (is_buffer_var(op->args[1])) {
                 devices_touched.insert(current_device_api);
+                touched_as_texture = touched_as_texture || op->is_intrinsic(Call::image_load_texture);
             }
             for (size_t i = 0; i < op->args.size(); i++) {
                 if (i == 1) continue;
                 op->args[i].accept(this);
             }
-        } else if (op->is_intrinsic(Call::image_store)) {
+        } else if (op->is_intrinsic(Call::image_store) ||
+                   op->is_intrinsic(Call::image_store_texture)) {
             internal_assert(!op->args.empty());
             if (is_buffer_var(op->args[1])) {
                 devices_touched.insert(current_device_api);
                 devices_writing.insert(current_device_api);
+
+                touched_as_texture = touched_as_texture || op->is_intrinsic(Call::image_store_texture);
             }
             for (size_t i = 0; i < op->args.size(); i++) {
                 if (i == 1) continue;
@@ -126,6 +131,8 @@ class FindBufferUsage : public IRVisitor {
     // bits and device allocation messed with.
     std::set<DeviceAPI> devices_touched_by_extern;
 
+    bool touched_as_texture = false;
+
     FindBufferUsage(const std::string &buf, DeviceAPI d)
         : buffer(buf), current_device_api(d) {
     }
@@ -145,6 +152,8 @@ class InjectBufferCopiesForSingleBuffer : public IRMutator {
 
     bool is_external;
 
+    bool is_texture;
+
     enum FlagState {
         Unknown,
         False,
@@ -179,7 +188,7 @@ class InjectBufferCopiesForSingleBuffer : public IRMutator {
     }
 
     Stmt make_device_malloc(DeviceAPI target_device_api) {
-        Expr device_interface = make_device_interface_call(target_device_api);
+        Expr device_interface = make_device_interface_call(target_device_api, is_texture);
         Stmt device_malloc = call_extern_and_assert("halide_device_malloc",
                                                     {buffer_var(), device_interface});
         return device_malloc;
@@ -190,7 +199,7 @@ class InjectBufferCopiesForSingleBuffer : public IRMutator {
     }
 
     Stmt make_copy_to_device(DeviceAPI target_device_api) {
-        Expr device_interface = make_device_interface_call(target_device_api);
+        Expr device_interface = make_device_interface_call(target_device_api, is_texture);
         return call_extern_and_assert("halide_copy_to_device", {buffer_var(), device_interface});
     }
 
@@ -401,8 +410,8 @@ class InjectBufferCopiesForSingleBuffer : public IRMutator {
     }
 
 public:
-    InjectBufferCopiesForSingleBuffer(const std::string &b, bool e)
-        : buffer(b), is_external(e) {
+    InjectBufferCopiesForSingleBuffer(const std::string &b, bool e, bool t)
+        : buffer(b), is_external(e), is_texture(t) {
         if (is_external) {
             // The state of the buffer is totally unknown, which is
             // the default constructor for this->state
@@ -539,7 +548,6 @@ class InjectBufferCopies : public IRMutator {
                 body = Block::make(destructor, body);
 
                 // Then the device_and_host malloc
-                Expr device_interface = make_device_interface_call(device_api);
                 Stmt device_malloc = call_extern_and_assert("halide_device_and_host_malloc",
                                                             {buf, device_interface});
                 if (!is_one(condition)) {
@@ -561,11 +569,12 @@ class InjectBufferCopies : public IRMutator {
         Type type;
         vector<Expr> extents;
         Expr condition;
-        DeviceAPI device_api;
+        Expr device_interface;
 
     public:
-        InjectCombinedAllocation(string b, Type t, vector<Expr> e, Expr c, DeviceAPI d)
-            : buffer(std::move(b)), type(t), extents(std::move(e)), condition(std::move(c)), device_api(d) {
+        InjectCombinedAllocation(string b, Type t, vector<Expr> e, Expr c, Expr d)
+            : buffer(std::move(b)), type(t), extents(std::move(e)),
+              condition(std::move(c)), device_interface(std::move(d)) {
         }
     };
 
@@ -606,7 +615,7 @@ class InjectBufferCopies : public IRMutator {
 
         Stmt body = mutate(op->body);
 
-        InjectBufferCopiesForSingleBuffer injector(op->name, false);
+        InjectBufferCopiesForSingleBuffer injector(op->name, false, op->memory_type == MemoryType::GPUTexture);
         body = injector.mutate(body);
 
         string buffer_name = op->name + ".buffer";
@@ -634,8 +643,10 @@ class InjectBufferCopies : public IRMutator {
                 internal_assert(free_injecter.success);
             }
 
+            Expr device_interface = make_device_interface_call(touching_device, op->memory_type == MemoryType::GPUTexture);
+
             return InjectCombinedAllocation(op->name, op->type, op->extents,
-                                            op->condition, touching_device)
+                                            op->condition, device_interface)
                 .mutate(body);
         } else {
             // Only touched on host but passed to an extern stage, or
@@ -722,6 +733,10 @@ class InjectBufferCopiesForInputsAndOutputs : public IRMutator {
         void include(const Parameter &p) {
             if (p.defined()) {
                 result.insert(p.name());
+
+                if (p.memory_type() == MemoryType::GPUTexture) {
+                    result_textures.insert(p.name());
+                }
             }
         }
 
@@ -749,6 +764,7 @@ class InjectBufferCopiesForInputsAndOutputs : public IRMutator {
 
     public:
         set<string> result;
+        set<string> result_textures;
     };
 
 public:
@@ -760,7 +776,7 @@ class InjectBufferCopiesForInputsAndOutputs : public IRMutator {
             s.accept(&finder);
             Stmt new_stmt = s;
             for (const string &buf : finder.result) {
-                new_stmt = InjectBufferCopiesForSingleBuffer(buf, true).mutate(new_stmt);
+                new_stmt = InjectBufferCopiesForSingleBuffer(buf, true, finder.result_textures.count(buf)).mutate(new_stmt);
             }
             return new_stmt;
         } else {
diff --git a/src/OutputImageParam.cpp b/src/OutputImageParam.cpp
index 9702ef91b24e..a59ff13c43f1 100644
--- a/src/OutputImageParam.cpp
+++ b/src/OutputImageParam.cpp
@@ -102,4 +102,9 @@ OutputImageParam &OutputImageParam::set_estimates(const Region &estimates) {
     return *this;
 }
 
+OutputImageParam &OutputImageParam::store_in(MemoryType type) {
+    param.store_in(type);
+    return *this;
+}
+
 }  // namespace Halide
diff --git a/src/OutputImageParam.h b/src/OutputImageParam.h
index 4cb66ae40a70..a3aeb0ca5d7b 100644
--- a/src/OutputImageParam.h
+++ b/src/OutputImageParam.h
@@ -117,6 +117,10 @@ class OutputImageParam {
      * repeatedly, but slightly terser. The size of the estimates vector
      * must match the dimensionality of the ImageParam. */
     OutputImageParam &set_estimates(const Region &estimates);
+
+    /** Set the desired storage type for this parameter.  Only useful
+     * for MemoryType::GPUTexture at present */
+    OutputImageParam &store_in(MemoryType type);
 };
 
 }  // namespace Halide
diff --git a/src/Parameter.cpp b/src/Parameter.cpp
index 12d021348159..724e72653c3c 100644
--- a/src/Parameter.cpp
+++ b/src/Parameter.cpp
@@ -24,6 +24,7 @@ struct ParameterContents {
     std::vector<BufferConstraint> buffer_constraints;
     Expr scalar_min, scalar_max, scalar_estimate;
     const bool is_buffer;
+    MemoryType memory_type = MemoryType::Auto;
 
     ParameterContents(Type t, bool b, int d, const std::string &n)
         : type(t), dimensions(d), name(n), buffer(Buffer<>()), data(0),
@@ -349,5 +350,15 @@ void check_call_arg_types(const std::string &name, std::vector<Expr> *args, int
     }
 }
 
+void Parameter::store_in(MemoryType memory_type) {
+    check_is_buffer();
+    contents->memory_type = memory_type;
+}
+
+MemoryType Parameter::memory_type() const {
+    // check_is_buffer();
+    return contents->memory_type;
+}
+
 }  // namespace Internal
 }  // namespace Halide
diff --git a/src/Parameter.h b/src/Parameter.h
index 73e91d1060a1..95abe9bb89bd 100644
--- a/src/Parameter.h
+++ b/src/Parameter.h
@@ -18,6 +18,7 @@ template<typename T>
 class Buffer;
 struct Expr;
 struct Type;
+enum class MemoryType;
 
 namespace Internal {
 
@@ -157,6 +158,9 @@ class Parameter {
 
     /** Get the ArgumentEstimates appropriate for this Parameter. */
     ArgumentEstimates get_argument_estimates() const;
+
+    void store_in(MemoryType memory_type);
+    MemoryType memory_type() const;
 };
 
 /** Validate arguments to a call to a func, image or imageparam. */
diff --git a/src/StorageFlattening.cpp b/src/StorageFlattening.cpp
index 03d4f717cfd9..a13388a3747c 100644
--- a/src/StorageFlattening.cpp
+++ b/src/StorageFlattening.cpp
@@ -5,6 +5,7 @@
 #include "FuseGPUThreadLoops.h"
 #include "IRMutator.h"
 #include "IROperator.h"
+#include "IRPrinter.h"
 #include "Parameter.h"
 #include "Scope.h"
 
@@ -36,9 +37,11 @@ class FlattenDimensions : public IRMutator {
 private:
     const map<string, pair<Function, int>> &env;
     set<string> outputs;
+    set<string> textures;
     const Target &target;
     Scope<> realizations, shader_scope_realizations;
     bool in_shader = false;
+    bool in_gpu = false;
 
     Expr make_shape_var(string name, const string &field, size_t dim,
                         const Buffer<> &buf, const Parameter &param) {
@@ -111,6 +114,11 @@ class FlattenDimensions : public IRMutator {
             shader_scope_realizations.push(op->name);
         }
 
+        if (op->memory_type == MemoryType::GPUTexture) {
+            textures.insert(op->name);
+            debug(2) << "found texture " << op->name << "\n";
+        }
+
         Stmt body = mutate(op->body);
 
         // Compute the size
@@ -232,6 +240,13 @@ class FlattenDimensions : public IRMutator {
             }
         }
 
+        if (output_buf.defined()) {
+            debug(2) << "have output buf " << output_buf.name() << " " << output_buf.memory_type() << "\n";
+            if (output_buf.memory_type() == MemoryType::GPUTexture) {
+                textures.insert(op->name);
+            }
+        }
+
         Expr value = mutate(op->values[0]);
         if (in_shader && !shader_scope_realizations.contains(op->name)) {
             user_assert(op->args.size() == 3)
@@ -245,6 +260,21 @@ class FlattenDimensions : public IRMutator {
             Expr store = Call::make(value.type(), Call::image_store,
                                     args, Call::Intrinsic);
             return Evaluate::make(store);
+        } else if (in_gpu && textures.count(op->name)) {
+            debug(2) << " lower texture store to " << op->name << "\n";
+            Expr buffer_var =
+                Variable::make(type_of<halide_buffer_t *>(), op->name + ".buffer", output_buf);
+            vector<Expr> args(2);
+            args[0] = op->name;
+            args[1] = buffer_var;
+            for (size_t i = 0; i < op->args.size(); i++) {
+                Expr min = Variable::make(Int(32), op->name + ".min." + std::to_string(i));
+                args.push_back(op->args[i] - min);
+            }
+            args.push_back(value);
+            Expr store = Call::make(value.type(), Call::image_store_texture,
+                                    args, Call::Intrinsic);
+            return Evaluate::make(store);
         } else {
             Expr idx = mutate(flatten_args(op->name, op->args, Buffer<>(), output_buf));
             return Store::make(op->name, value, idx, output_buf, const_true(value.type().lanes()), ModulusRemainder());
@@ -255,9 +285,20 @@ class FlattenDimensions : public IRMutator {
         if (op->call_type == Call::Halide ||
             op->call_type == Call::Image) {
 
+            debug(2) << " load call to " << op->name << " " << textures.count(op->name) << "\n";
+            if (op->param.defined()) {
+                debug(2) << "     is param: "
+                         << " " << op->param.name() << " " << op->param.memory_type()
+                         << "\n";
+
+                if (op->param.memory_type() == MemoryType::GPUTexture) {
+                    textures.insert(op->name);
+                }
+            }
+
             internal_assert(op->value_index == 0);
 
-            if (in_shader && !shader_scope_realizations.contains(op->name)) {
+            if ((in_shader && !shader_scope_realizations.contains(op->name)) || (in_gpu && textures.count(op->name))) {
                 ReductionDomain rdom;
                 Expr buffer_var =
                     Variable::make(type_of<halide_buffer_t *>(), op->name + ".buffer",
@@ -276,13 +317,9 @@ class FlattenDimensions : public IRMutator {
                     args.push_back(mutate(op->args[i]) - min);
                     args.push_back(extent);
                 }
-                for (size_t i = op->args.size(); i < 3; i++) {
-                    args.emplace_back(0);
-                    args.emplace_back(1);
-                }
 
                 return Call::make(op->type,
-                                  Call::image_load,
+                                  textures.count(op->name) ? Call::image_load_texture : Call::image_load,
                                   args,
                                   Call::PureIntrinsic,
                                   FunctionPtr(),
@@ -362,13 +399,19 @@ class FlattenDimensions : public IRMutator {
 
     Stmt visit(const For *op) override {
         bool old_in_shader = in_shader;
+        bool old_in_gpu = in_gpu;
         if ((op->for_type == ForType::GPUBlock ||
              op->for_type == ForType::GPUThread) &&
             op->device_api == DeviceAPI::GLSL) {
             in_shader = true;
         }
+        if (op->for_type == ForType::GPUBlock ||
+            op->for_type == ForType::GPUThread) {
+            in_gpu = true;
+        }
         Stmt stmt = IRMutator::visit(op);
         in_shader = old_in_shader;
+        in_gpu = old_in_gpu;
         return stmt;
     }
 };
diff --git a/src/runtime/HalideRuntimeOpenCL.h b/src/runtime/HalideRuntimeOpenCL.h
index e3b27f84c800..9bc35def26c8 100644
--- a/src/runtime/HalideRuntimeOpenCL.h
+++ b/src/runtime/HalideRuntimeOpenCL.h
@@ -19,6 +19,7 @@ extern "C" {
 #define HALIDE_RUNTIME_OPENCL
 
 extern const struct halide_device_interface_t *halide_opencl_device_interface();
+extern const struct halide_device_interface_t *halide_opencl_image_device_interface();
 
 /** These are forward declared here to allow clients to override the
  *  Halide OpenCL runtime. Do not call them. */
@@ -90,6 +91,11 @@ extern const char *halide_opencl_get_build_options(void *user_context);
  * dirty bits are left unmodified. */
 extern int halide_opencl_wrap_cl_mem(void *user_context, struct halide_buffer_t *buf, uint64_t device_ptr);
 
+/** Same as halide_opencl_wrap_cl_mem but wraps a cl_mem created with
+ * clCreateImage
+ */
+extern int halide_opencl_image_wrap_cl_mem(void *user_context, struct halide_buffer_t *buf, uint64_t device_ptr);
+
 /** Disconnect a halide_buffer_t from the memory it was previously
  * wrapped around. Should only be called for a halide_buffer_t that
  * halide_opencl_wrap_device_ptr was previously called on. Frees any
diff --git a/src/runtime/opencl.cpp b/src/runtime/opencl.cpp
index b66b66cebd40..e9fd613f65c5 100644
--- a/src/runtime/opencl.cpp
+++ b/src/runtime/opencl.cpp
@@ -11,9 +11,8 @@ namespace Runtime {
 namespace Internal {
 namespace OpenCL {
 
-// Define the function pointers for the OpenCL API. OpenCL 1.2
-// currently disabled so we can work on build bots without it.
-//#define HAVE_OPENCL_12
+// Define the function pointers for the OpenCL API.
+#define HAVE_OPENCL_12
 #define CL_FN(ret, fn, args) WEAK ret(CL_API_CALL *fn) args;
 #include "cl_functions.h"
 
@@ -67,6 +66,7 @@ WEAK void load_libopencl(void *user_context) {
 }
 
 extern WEAK halide_device_interface_t opencl_device_interface;
+extern WEAK halide_device_interface_t opencl_image_device_interface;
 
 WEAK const char *get_opencl_error_name(cl_int err);
 WEAK int create_opencl_context(void *user_context, cl_context *ctx, cl_command_queue *q);
@@ -687,8 +687,11 @@ WEAK int halide_opencl_initialize_kernels(void *user_context, void **state_ptr,
         } else {
             debug(user_context) << (void *)program << "\n";
         }
-        (*state)->program = program;
 
+        // halide_print(user_context, "Source: \n");
+        // halide_print(user_context, src);
+
+        (*state)->program = program;
         debug(user_context) << "    clBuildProgram " << (void *)program
                             << " " << options.str() << "\n";
         err = clBuildProgram(program, 1, devices, options.str(), NULL, NULL);
@@ -696,7 +699,7 @@ WEAK int halide_opencl_initialize_kernels(void *user_context, void **state_ptr,
 
             {
                 // Allocate an appropriately sized buffer for the build log.
-                Printer<ErrorPrinter, 8192> p(user_context);
+                Printer<ErrorPrinter, 16384> p(user_context);
 
                 p << "CL: clBuildProgram failed: "
                   << get_opencl_error_name(err)
@@ -707,7 +710,7 @@ WEAK int halide_opencl_initialize_kernels(void *user_context, void **state_ptr,
                                           CL_PROGRAM_BUILD_LOG,
                                           p.capacity() - p.size() - 1, p.dst,
                                           NULL) != CL_SUCCESS) {
-                    p << "clGetProgramBuildInfo failed";
+                    p << "clGetProgramBuildInfo failed (Printer buffer too small?)";
                 }
             }
 
@@ -1198,7 +1201,8 @@ WEAK int halide_opencl_detach_cl_mem(void *user_context, halide_buffer_t *buf) {
     if (buf->device == NULL) {
         return 0;
     }
-    halide_assert(user_context, buf->device_interface == &opencl_device_interface);
+    halide_assert(user_context, buf->device_interface == &opencl_device_interface ||
+                                    buf->device_interface == &opencl_image_device_interface);
     free((device_handle *)buf->device);
     buf->device = 0;
     buf->device_interface->impl->release_module();
@@ -1210,7 +1214,8 @@ WEAK uintptr_t halide_opencl_get_cl_mem(void *user_context, halide_buffer_t *buf
     if (buf->device == NULL) {
         return 0;
     }
-    halide_assert(user_context, buf->device_interface == &opencl_device_interface);
+    halide_assert(user_context, buf->device_interface == &opencl_device_interface ||
+                                    buf->device_interface == &opencl_image_device_interface);
     return (uintptr_t)((device_handle *)buf->device)->mem;
 }
 
@@ -1488,3 +1493,386 @@ WEAK halide_device_interface_t opencl_device_interface = {
 }  // namespace Internal
 }  // namespace Runtime
 }  // namespace Halide
+
+extern "C" {
+
+WEAK int halide_opencl_image_device_malloc(void *user_context, halide_buffer_t *buf) {
+    debug(user_context)
+        << "CL: halide_opencl_image_device_malloc (user_context: " << user_context
+        << ", buf: " << buf << ")\n";
+
+    ClContext ctx(user_context);
+    if (ctx.error_code != CL_SUCCESS) {
+        return ctx.error_code;
+    }
+
+    size_t size = buf->size_in_bytes();
+    halide_assert(user_context, size != 0);
+    if (buf->device) {
+        halide_assert(user_context, validate_device_pointer(user_context, buf, size));
+        return 0;
+    }
+
+    for (int i = 0; i < buf->dimensions; i++) {
+        halide_assert(user_context, buf->dim[i].stride >= 0);
+    }
+
+    debug(user_context) << "    allocating " << *buf << "\n";
+
+#ifdef DEBUG_RUNTIME
+    uint64_t t_before = halide_current_time_ns(user_context);
+#endif
+
+    device_handle *dev_handle = (device_handle *)malloc(sizeof(device_handle));
+    if (dev_handle == NULL) {
+        return CL_OUT_OF_HOST_MEMORY;
+    }
+
+    cl_image_format format;
+    cl_image_desc desc;
+
+    struct halide_type_t type = buf->type;
+    if (type.code == halide_type_int) {
+        if (type.bits == 8) {
+            format.image_channel_data_type = CL_SIGNED_INT8;
+        } else if (type.bits == 16) {
+            format.image_channel_data_type = CL_SIGNED_INT16;
+        } else if (type.bits == 32) {
+            format.image_channel_data_type = CL_SIGNED_INT32;
+        } else {
+            halide_assert(user_context, false && "unhandled int bit width for image");
+        }
+    } else if (type.code == halide_type_uint) {
+        if (type.bits == 8) {
+            format.image_channel_data_type = CL_UNSIGNED_INT8;
+        } else if (type.bits == 16) {
+            format.image_channel_data_type = CL_UNSIGNED_INT16;
+        } else if (type.bits == 32) {
+            format.image_channel_data_type = CL_UNSIGNED_INT32;
+        } else {
+            halide_assert(user_context, false && "unhandled uint bit width for image");
+        }
+    } else if (type.code == halide_type_float) {
+        if (type.bits == 16) {
+            format.image_channel_data_type = CL_HALF_FLOAT;
+        } else if (type.bits == 32) {
+            format.image_channel_data_type = CL_FLOAT;
+        } else {
+            halide_assert(user_context, false && "unhandled float bit width for image");
+        }
+    } else {
+        halide_assert(user_context, false && "unhandled data type for image");
+    }
+
+    int last_dim_size = buf->dim[buf->dimensions - 1].extent;
+    format.image_channel_order = CL_R;
+
+    // if (buf->host == NULL) {
+    //     size_t size = buf->size_in_bytes();
+    //     debug(user_context) << "manually allocating buf->host";
+    //     buf->host = (uint8_t *)halide_malloc(user_context, size);
+    //     if (buf->host == NULL) {
+    //         return -1;
+    // debug(user_context) << *buf;
+    //     }
+    // }
+
+    debug(user_context) << "      format=(" << format.image_channel_data_type << ", " << format.image_channel_order << ")\n";
+
+    if (buf->dimensions == 1) {
+        desc.image_type = CL_MEM_OBJECT_IMAGE1D;
+    } else if (buf->dimensions == 2) {
+        desc.image_type = CL_MEM_OBJECT_IMAGE2D;
+    } else if (buf->dimensions == 3) {
+        desc.image_type = CL_MEM_OBJECT_IMAGE3D;
+    } else {
+        halide_assert(user_context, buf->dimensions >= 1 && buf->dimensions <= 3);
+    }
+    desc.image_width = buf->dim[0].extent;
+    desc.image_height = buf->dimensions >= 2 ? buf->dim[1].extent : 1;
+    desc.image_depth = buf->dimensions >= 3 ? buf->dim[1].extent : 1;
+    desc.image_array_size = 1;
+    // desc.image_row_pitch = buf->dimensions >= 2 ? buf->dim[1].stride * buf->type.bytes() : 0;
+    // desc.image_slice_pitch = buf->dimensions >= 3 ? buf->dim[2].stride * buf->type.bytes() : 0;
+    desc.image_row_pitch = 0;
+    halide_assert(user_context, buf->dimensions < 2 || buf->dim[1].stride == buf->dim[0].extent);
+    desc.image_slice_pitch = 0;
+    halide_assert(user_context, buf->dimensions < 3 || buf->dim[2].stride == buf->dim[0].extent * buf->dim[1].extent);
+    desc.num_mip_levels = 0;
+    desc.num_samples = 0;
+    desc.buffer = NULL;
+
+    debug(user_context) << "      desc=("
+                        << (int)desc.image_type << ", "
+                        << (int)desc.image_width << ", "
+                        << (int)desc.image_height << ", "
+                        << (int)desc.image_depth << ", "
+                        << (int)desc.image_array_size << ", "
+                        << (int)desc.image_row_pitch << ", "
+                        << (int)desc.image_slice_pitch << ", "
+                        << (void *)desc.buffer
+                        << ")\n";
+
+    cl_int err;
+    debug(user_context) << "    clCreateImage -> " << (int)size << " ";
+    cl_mem dev_ptr = clCreateImage(ctx.context, CL_MEM_READ_WRITE, &format, &desc, NULL, &err);
+    if (err != CL_SUCCESS || dev_ptr == 0) {
+        debug(user_context) << get_opencl_error_name(err) << "\n";
+        error(user_context) << "CL: clCreateImage failed: "
+                            << get_opencl_error_name(err);
+        free(dev_handle);
+        return err;
+    } else {
+        debug(user_context) << (void *)dev_ptr << " device_handle: " << dev_handle << "\n";
+    }
+
+    dev_handle->mem = dev_ptr;
+    dev_handle->offset = 0;
+    buf->device = (uint64_t)dev_handle;
+    buf->device_interface = &opencl_image_device_interface;
+    buf->device_interface->impl->use_module();
+
+    debug(user_context)
+        << "    Allocated device buffer " << (void *)buf->device
+        << " for buffer " << buf << "\n";
+
+    halide_assert(user_context, validate_device_pointer(user_context, buf, size));
+
+#ifdef DEBUG_RUNTIME
+    uint64_t t_after = halide_current_time_ns(user_context);
+    debug(user_context) << "    Time: " << (t_after - t_before) / 1.0e6 << " ms\n";
+#endif
+
+    return CL_SUCCESS;
+}
+
+WEAK int halide_opencl_image_buffer_copy(void *user_context, struct halide_buffer_t *src,
+                                         const struct halide_device_interface_t *dst_device_interface,
+                                         struct halide_buffer_t *dst) {
+    // We only handle copies to opencl or to host
+    debug(user_context)
+        << "CL: halide_opencl_image_buffer_copy (user_context: " << user_context
+        << ", src: " << src << ", dst: " << dst << ")\n";
+
+    halide_assert(user_context, dst_device_interface == NULL ||
+                                    dst_device_interface == &opencl_image_device_interface);
+
+    if ((src->device_dirty() || src->host == NULL) &&
+        src->device_interface != &opencl_image_device_interface) {
+        halide_assert(user_context, dst_device_interface == &opencl_image_device_interface);
+        // This is handled at the higher level.
+        return halide_error_code_incompatible_device_interface;
+    }
+
+    bool from_host = (src->device_interface != &opencl_image_device_interface) ||
+                     (src->device == 0) ||
+                     (src->host_dirty() && src->host != NULL);
+    bool to_host = !dst_device_interface;
+
+    halide_assert(user_context, from_host || src->device);
+    halide_assert(user_context, to_host || dst->device);
+
+    device_copy c = make_buffer_copy(src, from_host, dst, to_host);
+
+    int err = 0;
+    {
+        ClContext ctx(user_context);
+        if (ctx.error_code != CL_SUCCESS) {
+            return ctx.error_code;
+        }
+
+#ifdef DEBUG_RUNTIME
+        uint64_t t_before = halide_current_time_ns(user_context);
+        if (!from_host) {
+            halide_assert(user_context, validate_device_pointer(user_context, src));
+        }
+        if (!to_host) {
+            halide_assert(user_context, validate_device_pointer(user_context, dst));
+        }
+#endif
+
+        debug(user_context) << "    from " << (from_host ? "host" : "device")
+                            << " to " << (to_host ? "host" : "device") << ", "
+                            << (void *)c.src << " + " << 0
+                            << " -> " << (void *)c.dst << " + " << 0
+                            << ", " << c.chunk_size << " bytes\n";
+
+        halide_assert(user_context, c.chunk_size == src->size_in_bytes());
+        halide_assert(user_context, c.chunk_size == dst->size_in_bytes());
+        if (!from_host && to_host) {
+            int dim = dst->dimensions;
+            size_t offset[] = {0, 0, 0};
+            size_t region[] = {
+                static_cast<size_t>(dst->dim[0].extent),
+                dim >= 2 ? static_cast<size_t>(dst->dim[1].extent) : 1,
+                dim >= 3 ? static_cast<size_t>(dst->dim[2].extent) : 1};
+
+            // int row_pitch = dst->dimensions >= 2 ? dst->dim[1].stride * dst->type.bytes() : 0;
+            // int slice_pitch = dst->dimensions >= 3 ? dst->dim[2].stride * dst->type.bytes() : 0;
+            halide_assert(user_context, dst->dimensions < 2 || dst->dim[1].stride == dst->dim[0].extent);
+            halide_assert(user_context, dst->dimensions < 3 || dst->dim[2].stride == dst->dim[0].extent * dst->dim[1].extent);
+
+            err = clEnqueueReadImage(ctx.cmd_queue, ((device_handle *)c.src)->mem,
+                                     CL_FALSE, offset, region,
+                                     /* row_pitch */ 0, /* slice_pitch */ 0,
+                                     dst->host, 0, NULL, NULL);
+        } else if (from_host && !to_host) {
+            int dim = src->dimensions;
+            size_t offset[] = {0, 0, 0};
+            size_t region[] = {
+                static_cast<size_t>(src->dim[0].extent),
+                dim >= 2 ? static_cast<size_t>(src->dim[1].extent) : 1,
+                dim >= 3 ? static_cast<size_t>(src->dim[2].extent) : 1};
+            // int row_pitch = dim >= 2 ? src->dim[1].stride * src->type.bytes() : 0;
+            // int slice_pitch = dim >= 3 ? src->dim[2].stride * src->type.bytes() : 0;
+            halide_assert(user_context, src->dimensions < 2 || src->dim[1].stride == src->dim[0].extent);
+            halide_assert(user_context, src->dimensions < 3 || src->dim[2].stride == src->dim[0].extent * src->dim[1].extent);
+            err = clEnqueueWriteImage(ctx.cmd_queue, ((device_handle *)c.dst)->mem,
+                                      CL_FALSE, offset, region, /* row_pitch */ 0, /* slice_pitch */ 0, src->host,
+                                      0, NULL, NULL);
+        } else if (!from_host && !to_host) {
+            halide_assert(user_context, false && "image to image copies not implemented");
+            // err = clEnqueueCopyBuffer(ctx.cmd_queue, ((device_handle *)c.src)->mem, ((device_handle *)c.dst)->mem,
+            //                           src_idx + ((device_handle *)c.src)->offset, dst_idx + ((device_handle *)c.dst)->offset,
+            //                           c.chunk_size, 0, NULL, NULL);
+        }
+
+        if (err != CL_SUCCESS) {
+            debug(user_context) << get_opencl_error_name(err) << "\n";
+            error(user_context) << "CL: buffer transfer failed: "
+                                << get_opencl_error_name(err);
+            return err;
+        }
+
+        // The reads/writes above are all non-blocking, so empty the command
+        // queue before we proceed so that other host code won't write
+        // to the buffer while the above writes are still running.
+        clFinish(ctx.cmd_queue);
+
+#ifdef DEBUG_RUNTIME
+        uint64_t t_after = halide_current_time_ns(user_context);
+        debug(user_context) << "    Time: " << (t_after - t_before) / 1.0e6 << " ms\n";
+#endif
+    }
+
+    return err;
+}
+
+WEAK int halide_opencl_image_copy_to_device(void *user_context, halide_buffer_t *buf) {
+    return halide_opencl_image_buffer_copy(user_context, buf, &opencl_image_device_interface, buf);
+}
+
+WEAK int halide_opencl_image_copy_to_host(void *user_context, halide_buffer_t *buf) {
+    return halide_opencl_image_buffer_copy(user_context, buf, NULL, buf);
+}
+
+WEAK int halide_opencl_image_device_and_host_malloc(void *user_context, struct halide_buffer_t *buf) {
+    return halide_default_device_and_host_malloc(user_context, buf, &opencl_image_device_interface);
+}
+
+WEAK int halide_opencl_image_device_and_host_free(void *user_context, struct halide_buffer_t *buf) {
+    return halide_default_device_and_host_free(user_context, buf, &opencl_image_device_interface);
+}
+
+WEAK int halide_opencl_image_wrap_cl_mem(void *user_context, struct halide_buffer_t *buf, uint64_t mem) {
+    halide_assert(user_context, buf->device == 0);
+    if (buf->device != 0) {
+        return -2;
+    }
+    device_handle *dev_handle = (device_handle *)malloc(sizeof(device_handle));
+    if (dev_handle == NULL) {
+        return halide_error_code_out_of_memory;
+    }
+    dev_handle->mem = (cl_mem)mem;
+    dev_handle->offset = 0;
+    buf->device = (uint64_t)dev_handle;
+    buf->device_interface = &opencl_image_device_interface;
+    buf->device_interface->impl->use_module();
+#ifdef DEBUG_RUNTIME
+    if (!validate_device_pointer(user_context, buf)) {
+        free((device_handle *)buf->device);
+        buf->device = 0;
+        buf->device_interface->impl->release_module();
+        buf->device_interface = NULL;
+        return -3;
+    }
+#endif
+    return 0;
+}
+
+WEAK int halide_opencl_image_device_crop(void *user_context,
+                                         const struct halide_buffer_t *src,
+                                         struct halide_buffer_t *dst) {
+    halide_assert(user_context, false && "crop not supported on opencl image objects");
+    return -1;
+}
+
+WEAK int halide_opencl_image_device_slice(void *user_context,
+                                          const struct halide_buffer_t *src,
+                                          int slice_dim,
+                                          int slice_pos,
+                                          struct halide_buffer_t *dst) {
+    halide_assert(user_context, false && "slice not supported on opencl image objects");
+    return -1;
+}
+
+WEAK int halide_opencl_image_device_release_crop(void *user_context,
+                                                 struct halide_buffer_t *buf) {
+    halide_assert(user_context, false && "crop not supported on opencl image objects");
+    return -1;
+}
+}
+
+namespace Halide {
+namespace Runtime {
+namespace Internal {
+namespace OpenCL {
+
+WEAK halide_device_interface_impl_t opencl_image_device_interface_impl = {
+    halide_use_jit_module,
+    halide_release_jit_module,
+    halide_opencl_image_device_malloc,
+    halide_opencl_device_free,
+    halide_opencl_device_sync,
+    halide_opencl_device_release,
+    halide_opencl_image_copy_to_host,
+    halide_opencl_image_copy_to_device,
+    halide_opencl_image_device_and_host_malloc,
+    halide_opencl_image_device_and_host_free,
+    halide_opencl_image_buffer_copy,
+    halide_opencl_image_device_crop,
+    halide_opencl_image_device_slice,
+    halide_opencl_image_device_release_crop,
+    halide_opencl_image_wrap_cl_mem,
+    halide_opencl_detach_cl_mem,
+};
+
+WEAK halide_device_interface_t opencl_image_device_interface = {
+    halide_device_malloc,
+    halide_device_free,
+    halide_device_sync,
+    halide_device_release,
+    halide_copy_to_host,
+    halide_copy_to_device,
+    halide_device_and_host_malloc,
+    halide_device_and_host_free,
+    halide_buffer_copy,
+    halide_device_crop,
+    halide_device_slice,
+    halide_device_release_crop,
+    halide_device_wrap_native,
+    halide_device_detach_native,
+    NULL,
+    &opencl_image_device_interface_impl};
+
+}  // namespace OpenCL
+}  // namespace Internal
+}  // namespace Runtime
+}  // namespace Halide
+
+extern "C" {
+
+WEAK const struct halide_device_interface_t *halide_opencl_image_device_interface() {
+    return &opencl_image_device_interface;
+}
+}
\ No newline at end of file
diff --git a/src/runtime/runtime_api.cpp b/src/runtime/runtime_api.cpp
index 76d1b919f755..6071932000da 100644
--- a/src/runtime/runtime_api.cpp
+++ b/src/runtime/runtime_api.cpp
@@ -136,6 +136,8 @@ extern "C" __attribute__((used)) void *halide_runtime_api_functions[] = {
     (void *)&halide_opencl_get_device_type,
     (void *)&halide_opencl_get_platform_name,
     (void *)&halide_opencl_get_crop_offset,
+    (void *)&halide_opencl_image_device_interface,
+    (void *)&halide_opencl_image_wrap_cl_mem,
     (void *)&halide_opencl_initialize_kernels,
     (void *)&halide_opencl_run,
     (void *)&halide_opencl_set_build_options,
diff --git a/test/correctness/CMakeLists.txt b/test/correctness/CMakeLists.txt
index 17a4d4401aa0..85a3573cbdf0 100644
--- a/test/correctness/CMakeLists.txt
+++ b/test/correctness/CMakeLists.txt
@@ -147,6 +147,7 @@ tests(GROUPS correctness
       gpu_specialize.cpp
       gpu_store_in_register_with_no_lanes_loop.cpp
       gpu_sum_scan.cpp
+      gpu_texture.cpp
       gpu_thread_barrier.cpp
       gpu_transpose.cpp
       gpu_vectorized_shared_memory.cpp
diff --git a/test/correctness/gpu_texture.cpp b/test/correctness/gpu_texture.cpp
new file mode 100644
index 000000000000..3b1be49cbb5b
--- /dev/null
+++ b/test/correctness/gpu_texture.cpp
@@ -0,0 +1,145 @@
+#include "Halide.h"
+#include <stdio.h>
+
+using namespace Halide;
+using namespace Halide::Internal;
+
+int main(int argc, char **argv) {
+    // setenv("HL_JIT_TARGET", "host-opencl-debug", 1);
+
+    Target t = get_jit_target_from_environment();
+
+    if (!t.has_feature(halide_target_feature_opencl)) {
+        printf("[SKIP] No OpenCL target enabled.\n");
+        return 0;
+    }
+
+    // Check dynamic allocations per-block and per-thread into both
+    // shared and global
+    for (auto memory_type : {MemoryType::GPUTexture, MemoryType::Heap}) {
+        {
+            // 1D stores/loads
+            Buffer<int> input(100);
+            input.fill(10);
+            ImageParam param(Int(32), 1, "input");
+            param.set(input);
+            param.store_in(memory_type);  // check float stores
+
+            Func f("f"), g("g");
+            Var x("x"), xi("xi");
+            Var y("y");
+
+            f(x) = cast<float>(x);
+            g(x) = param(x) + cast<int>(f(2 * x));
+
+            g.gpu_tile(x, xi, 16);
+
+            f.compute_root().store_in(memory_type).gpu_blocks(x);  // store f as integer
+            g.output_buffer().store_in(memory_type);
+
+            Buffer<int> out = g.realize(100);
+            for (int x = 0; x < 100; x++) {
+                int correct = 2 * x + 10;
+                if (out(x) != correct) {
+                    printf("out[1D][%d](%d) = %d instead of %d\n", (int)memory_type, x, out(x), correct);
+                    return -1;
+                }
+            }
+        }
+        {
+            // 2D stores/loads
+            Buffer<int> input(10, 10);
+            input.fill(10);
+            ImageParam param(Int(32), 2);
+            param.set(input);
+            param.store_in(memory_type);  // check float stores
+
+            Func f("f"), g("g");
+            Var x("x"), xi("xi");
+            Var y("y");
+
+            f(x, y) = cast<float>(x + y);
+            g(x) = param(x, x) + cast<int>(f(2 * x, x));
+
+            g.gpu_tile(x, xi, 16, TailStrategy::GuardWithIf);
+
+            f.compute_root().store_in(memory_type).gpu_blocks(x, y);  // store f as integer
+            g.store_in(memory_type);
+
+            Buffer<int> out = g.realize(10);
+            for (int x = 0; x < 10; x++) {
+                int correct = 3 * x + 10;
+                if (out(x) != correct) {
+                    printf("out[2D][%d](%d) = %d instead of %d\n", (int)memory_type, x, out(x), correct);
+                    return -1;
+                }
+            }
+        }
+        {
+            // 3D stores/loads
+            Buffer<int> input(10, 10, 10);
+            input.fill(10);
+            ImageParam param(Int(32), 3);
+            param.set(input);
+            param.store_in(memory_type);  // check float stores
+
+            Func f("f"), g("g");
+            Var x("x"), xi("xi");
+            Var y("y"), z("z");
+
+            f(x, y, z) = cast<float>(x + y + z);
+            g(x) = param(x, x, x) + cast<int>(f(2 * x, x, x));
+
+            g.gpu_tile(x, xi, 16, TailStrategy::GuardWithIf);
+
+            f.compute_root().store_in(memory_type).gpu_blocks(x, y, z);  // store f as integer
+
+            g.store_in(memory_type);
+
+            Buffer<int> out = g.realize(10);
+            for (int x = 0; x < 10; x++) {
+                int correct = 4 * x + 10;
+                if (out(x) != correct) {
+                    printf("out[3D][%d](%d) = %d instead of %d\n", (int)memory_type, x, out(x), correct);
+                    return -1;
+                }
+            }
+        }
+        {
+            // 1D offset
+            Buffer<int> input(100);
+            input.set_min(5);
+            input.fill(10);
+            ImageParam param(Int(32), 1);
+            param.set(input);
+            param.store_in(memory_type);  // check float stores
+
+            Func f("f"), g("g");
+            Var x("x"), xi("xi");
+            Var y("y");
+
+            f(x) = cast<float>(x);
+            g(x) = param(x) + cast<int>(f(2 * x));
+
+            g.gpu_tile(x, xi, 16, TailStrategy::GuardWithIf);
+
+            f.compute_root().store_in(memory_type).gpu_blocks(x);  // store f as integer
+            g.store_in(memory_type);
+
+            Buffer<int> out(10);
+            out.set_min(10);
+            g.realize(out);
+            out.copy_to_host();
+            for (int x = 10; x < 20; x++) {
+                int correct = 2 * x + 10;
+                if (out(x) != correct) {
+                    printf("out[1D-shift][%d](%d) = %d instead of %d\n", (int)memory_type, x, out(x), correct);
+                    return -1;
+                }
+            }
+        }
+    }
+
+    printf("Success!\n");
+    return 0;
+}
diff --git a/test/generator/CMakeLists.txt b/test/generator/CMakeLists.txt
index ccb1a0abf96d..c49974a59d3b 100644
--- a/test/generator/CMakeLists.txt
+++ b/test/generator/CMakeLists.txt
@@ -263,6 +263,10 @@ halide_define_aot_test(gpu_object_lifetime FEATURES debug)
 # gpu_only_generator.cpp
 halide_define_aot_test(gpu_only)
 
+# gpu_texture_aottest.cpp
+# gpu_texture_generator.cpp
+halide_define_aot_test(gpu_texture)
+
 # image_from_array_aottest.cpp
 # image_from_array_generator.cpp
 halide_define_aot_test(image_from_array)
diff --git a/test/generator/gpu_texture_aottest.cpp b/test/generator/gpu_texture_aottest.cpp
new file mode 100644
index 000000000000..8ec8b306ac16
--- /dev/null
+++ b/test/generator/gpu_texture_aottest.cpp
@@ -0,0 +1,64 @@
+#include "HalideBuffer.h"
+#include "HalideRuntime.h"
+#include <assert.h>
+#include <math.h>
+#include <stdio.h>
+#if defined(TEST_OPENCL)
+#include "HalideRuntimeOpenCL.h"
+#endif
+
+#include "gpu_texture.h"
+using namespace Halide::Runtime;
+
+#if defined(TEST_OPENCL)
+
+#if !defined(HALIDE_RUNTIME_OPENCL)
+#error "TEST_OPENCL defined but HALIDE_RUNTIME_OPENCL not defined"
+#endif
+
+#endif
+
+int main(int argc, char **argv) {
+#if defined(TEST_OPENCL)
+    const int W = 32, H = 32;
+    Buffer<int> input(W, H);
+    for (int y = 0; y < input.height(); y++) {
+        for (int x = 0; x < input.width(); x++) {
+            input(x, y) = x + y;
+        }
+    }
+
+    // Explicitly copy data to the GPU.
+    input.set_host_dirty();
+
+    Buffer<int> output(W, H);
+
+    gpu_texture(input, output);
+
+    if (input.raw_buffer()->device_interface != halide_opencl_image_device_interface()) {
+        printf("Expected input to be copied to texture storage");
+        return -1;
+    }
+    if (output.raw_buffer()->device_interface != halide_opencl_image_device_interface()) {
+        printf("Expected output to be copied to texture storage");
+        return -1;
+    }
+
+    output.copy_to_host();
+
+    // Verify output.
+    for (int y = 0; y < H; y++) {
+        for (int x = 0; x < W; x++) {
+            if (input(x, y) * 2 != output(x, y)) {
+                printf("Error at %d, %d: %d != %d\n", x, y, input(x, y), output(x, y));
+                return -1;
+            }
+        }
+    }
+
+    printf("Success!\n");
+#else
+    printf("[SKIP] No OpenCL target enabled.\n");
+#endif
+    return 0;
+}
diff --git a/test/generator/gpu_texture_generator.cpp b/test/generator/gpu_texture_generator.cpp
new file mode 100644
index 000000000000..52c537d095c4
--- /dev/null
+++ b/test/generator/gpu_texture_generator.cpp
@@ -0,0 +1,30 @@
+#include "Halide.h"
+
+namespace {
+
+class GpuTexture : public Halide::Generator<GpuTexture> {
+public:
+    Input<Buffer<int32_t>> input{"input", 2};
+
+    Output<Buffer<int32_t>> output{"output", 2};
+
+    void generate() {
+        Var x("x"), y("y");
+
+        // Create a simple pipeline that scales pixel values by 2.
+        output(x, y) = input(x, y) * 2;
+
+        input.store_in(MemoryType::GPUTexture);
+        output.store_in(MemoryType::GPUTexture);
+
+        Target target = get_target();
+        if (target.has_gpu_feature()) {
+            Var xo, yo, xi, yi;
+            output.gpu_tile(x, y, xo, yo, xi, yi, 16, 16);
+        }
+    }
+};
+
+}  // namespace
+
+HALIDE_REGISTER_GENERATOR(GpuTexture, gpu_texture)