From 26d00762bbfb598852c641cc8aced737e672ed88 Mon Sep 17 00:00:00 2001 From: John Laxson Date: Wed, 23 Sep 2020 18:54:55 -0600 Subject: [PATCH 01/14] yes --- src/Closure.h | 5 +- src/CodeGen_OpenCL_Dev.cpp | 159 +++++++++++++- src/DeviceArgument.cpp | 23 +- src/DeviceArgument.h | 5 + src/DeviceInterface.cpp | 8 +- src/DeviceInterface.h | 2 +- src/Expr.h | 4 + src/FuseGPUThreadLoops.cpp | 6 +- src/Generator.h | 1 + src/IR.cpp | 2 + src/IR.h | 2 + src/IRPrinter.cpp | 3 + src/ImageParam.cpp | 4 + src/ImageParam.h | 4 + src/InjectHostDevBufferCopies.cpp | 41 ++-- src/Parameter.cpp | 11 + src/Parameter.h | 4 + src/StorageFlattening.cpp | 47 ++++- src/runtime/CMakeLists.txt | 1 + src/runtime/opencl.cpp | 338 +++++++++++++++++++++++++++++- src/runtime/opencl_image.cpp | 0 test/correctness/CMakeLists.txt | 1 + test/correctness/gpu_texture.cpp | 113 ++++++++++ 23 files changed, 741 insertions(+), 43 deletions(-) create mode 100644 src/runtime/opencl_image.cpp create mode 100644 test/correctness/gpu_texture.cpp 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..a9e2ca3b8d08 100644 --- a/src/CodeGen_OpenCL_Dev.cpp +++ b/src/CodeGen_OpenCL_Dev.cpp @@ -1,6 +1,7 @@ #include #include #include +#include #include "CSE.h" #include "CodeGen_Internal.h" @@ -242,6 +243,128 @@ 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(, , , , , + // , , ) + int dims = (op->args.size() - 2) / 2; + internal_assert(dims >= 1 && dims <= 3); + const StringImm *string_imm = op->args[0].as(); + if (!string_imm) { + internal_assert(op->args[0].as()); + string_imm = op->args[0].as()->value.as(); + } + 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()); + + string type_suffix; + if (op->type.is_int()) { + type_suffix = "i"; + } else if (op->type.is_uint()) { + type_suffix = "ui"; + } else if (op->type.is_float()) { + type_suffix = "f"; + } else { + internal_error << "Invalid type for read_image: " << op->type << "\n"; + } + + std::array coord; + for (int i = 0; i < dims; i++) { + coord[i] = print_expr(op->args[i*2 + 2]); + } + vector 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" << type_suffix << "(" << 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(, , , , , ) + const StringImm *string_imm = op->args[0].as(); + if (!string_imm) { + internal_assert(op->args[0].as()); + string_imm = op->args[0].as()->value.as(); + } + 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()); + + string type_suffix; + if (op->type.is_int()) { + type_suffix = "i"; + } else if (op->type.is_uint()) { + type_suffix = "ui"; + } else if (op->type.is_float()) { + type_suffix = "f"; + } else { + internal_error << "Invalid type for write_image: " << op->type << "\n"; + } + + std::array 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" << type_suffix << "(" << 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); } @@ -590,9 +713,11 @@ void CodeGen_OpenCL_Dev::CodeGen_OpenCL_C::visit(const Free *op) { return; } else { // Should have been freed internally - internal_assert(allocations.contains(op->name)); - allocations.pop(op->name); - stream << get_indent() << "#undef " << get_memory_space(op->name) << "\n"; + if (allocations.contains(op->name)) { + internal_assert(allocations.contains(op->name)); + allocations.pop(op->name); + stream << get_indent() << "#undef " << get_memory_space(op->name) << "\n"; + } } } @@ -765,11 +890,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 +938,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..19c245ceb3d4 100644 --- a/src/DeviceArgument.cpp +++ b/src/DeviceArgument.cpp @@ -15,15 +15,17 @@ std::vector HostClosure::arguments() { std::vector 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) << " "; + 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 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 @@ -48,17 +52,24 @@ void HostClosure::visit(const Call *op) { internal_assert(string_imm); + + 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..860e3274af52 100644 --- a/src/DeviceArgument.h +++ b/src/DeviceArgument.h @@ -36,6 +36,8 @@ struct DeviceArgument { */ bool is_buffer; + 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 +68,7 @@ struct DeviceArgument { DeviceArgument() : is_buffer(false), + is_texture(false), dimensions(0), size(0), packed_index(0), @@ -75,11 +78,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()); } @@ -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.h b/src/Generator.h index 19a02cf0ebb5..ee82f71a8cc6 100644 --- a/src/Generator.h +++ b/src/Generator.h @@ -1785,6 +1785,7 @@ class GeneratorInput_Buffer : public GeneratorInputImpl { 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) diff --git a/src/IR.cpp b/src/IR.cpp index 550c2c78e25f..925f0f1e2853 100644 --- a/src/IR.cpp +++ b/src/IR.cpp @@ -610,7 +610,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 998d4856c37a..194ff2407c60 100644 --- a/src/IR.h +++ b/src/IR.h @@ -520,7 +520,9 @@ struct Call : public ExprNode { 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/ImageParam.cpp b/src/ImageParam.cpp index cda49b501c4c..c3a8db3c604d 100644 --- a/src/ImageParam.cpp +++ b/src/ImageParam.cpp @@ -95,4 +95,8 @@ ImageParam &ImageParam::add_trace_tag(const std::string &trace_tag) { return *this; } +void ImageParam::store_in(MemoryType type) { + param.store_in(type); +} + } // namespace Halide diff --git a/src/ImageParam.h b/src/ImageParam.h index d4383bf4ed7f..f8b1ce7f0d02 100644 --- a/src/ImageParam.h +++ b/src/ImageParam.h @@ -32,6 +32,8 @@ class ImageParam : public OutputImageParam { /** Helper function to initialize the Func representation of this ImageParam. */ Func create_func() const; + MemoryType memory_type = MemoryType::Auto; + public: /** Construct a nullptr image parameter handle. */ ImageParam() = default; @@ -133,6 +135,8 @@ class ImageParam : public OutputImageParam { /** Add a trace tag to this ImageParam's Func. */ ImageParam &add_trace_tag(const std::string &trace_tag); + + void store_in(MemoryType type); }; } // namespace Halide diff --git a/src/InjectHostDevBufferCopies.cpp b/src/InjectHostDevBufferCopies.cpp index 5827abfaf4bf..eecd7d764b70 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 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,11 @@ class InjectBufferCopies : public IRMutator { Type type; vector extents; Expr condition; - DeviceAPI device_api; + Expr device_interface; public: - InjectCombinedAllocation(string b, Type t, vector 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 e, Expr c, Expr d) + : buffer(std::move(b)), type(t), extents(std::move(e)), condition(std::move(c)), device_interface(d) { } }; @@ -606,7 +614,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 +642,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 +732,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 +763,7 @@ class InjectBufferCopiesForInputsAndOutputs : public IRMutator { public: set result; + set result_textures; }; public: @@ -760,7 +775,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/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 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 *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 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..18b67a9838d4 100644 --- a/src/StorageFlattening.cpp +++ b/src/StorageFlattening.cpp @@ -36,9 +36,11 @@ class FlattenDimensions : public IRMutator { private: const map> &env; set outputs; + set textures; const 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 ¶m) { @@ -111,6 +113,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 @@ -245,6 +252,20 @@ 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(), op->name + ".buffer", output_buf); + vector args(2); + args[0] = op->name; + args[1] = buffer_var; + for (size_t i = 0; i < op->args.size(); i++) { + args.push_back(op->args[i]); + } + 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 +276,21 @@ 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() << " " + << "\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(), op->name + ".buffer", @@ -276,13 +309,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 +391,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/CMakeLists.txt b/src/runtime/CMakeLists.txt index 91a59716e63e..1bb6bc21a9f6 100644 --- a/src/runtime/CMakeLists.txt +++ b/src/runtime/CMakeLists.txt @@ -47,6 +47,7 @@ set(RUNTIME_CPP msan msan_stubs opencl + opencl_image opengl opengl_egl_context opengl_glx_context diff --git a/src/runtime/opencl.cpp b/src/runtime/opencl.cpp index b66b66cebd40..a26941b98166 100644 --- a/src/runtime/opencl.cpp +++ b/src/runtime/opencl.cpp @@ -13,7 +13,7 @@ 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 HAVE_OPENCL_12 #define CL_FN(ret, fn, args) WEAK ret(CL_API_CALL *fn) args; #include "cl_functions.h" @@ -67,6 +67,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); @@ -619,6 +620,8 @@ WEAK int halide_opencl_initialize_kernels(void *user_context, void **state_ptr, return ctx.error_code; } + // std::cout << src; + #ifdef DEBUG_RUNTIME uint64_t t_before = halide_current_time_ns(user_context); #endif @@ -687,8 +690,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 +702,7 @@ WEAK int halide_opencl_initialize_kernels(void *user_context, void **state_ptr, { // Allocate an appropriately sized buffer for the build log. - Printer p(user_context); + Printer p(user_context); p << "CL: clBuildProgram failed: " << get_opencl_error_name(err) @@ -707,7 +713,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?)"; } } @@ -1488,3 +1494,327 @@ 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 = 0; //buf->dim[1].stride * buf->type.bytes(); + desc.image_slice_pitch = 0; + desc.num_mip_levels = 0; + desc.num_samples = 0; + desc.buffer = NULL; + + debug(user_context) << " desc=(\n"; + debug(user_context) << " " << (int)desc.image_type << ",\n"; + + debug(user_context) << " " << (int)desc.image_width << ",\n"; + debug(user_context) << " " << (int)desc.image_height << ",\n"; + debug(user_context) << " " << (int)desc.image_depth << ",\n"; + + debug(user_context) << " " << (int)desc.image_array_size << ",\n"; + debug(user_context) << " " << (int)desc.image_row_pitch << ",\n"; + debug(user_context) << " " << (int)desc.image_slice_pitch << ",\n"; + debug(user_context) << " " << (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_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(dst->dim[0].extent), + dim >= 2 ? static_cast(dst->dim[1].extent) : 1, + dim >= 3 ? static_cast(dst->dim[1].extent) : 1}; + int pitch = dst->dimensions >= 2 ? dst->dim[1].stride * dst->type.bytes() : 0; + err = clEnqueueReadImage(ctx.cmd_queue, ((device_handle *)c.src)->mem, + CL_FALSE, offset, region, + 0, 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(src->dim[0].extent), + dim >= 2 ? static_cast(src->dim[1].extent) : 1, + dim >= 3 ? static_cast(src->dim[1].extent) : 1}; + int pitch = dim >= 2 ? src->dim[1].stride * src->type.bytes() : 0; + err = clEnqueueWriteImage(ctx.cmd_queue, ((device_handle *)c.dst)->mem, + CL_FALSE, offset, region, 0, 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); +} +} + +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, + nullptr, //halide_opencl_image_device_crop, + nullptr, //halide_opencl_image_device_slice, + nullptr, //halide_opencl_image_device_release_crop, + halide_opencl_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/opencl_image.cpp b/src/runtime/opencl_image.cpp new file mode 100644 index 000000000000..e69de29bb2d1 diff --git a/test/correctness/CMakeLists.txt b/test/correctness/CMakeLists.txt index 8c3c5d4638bf..026cfd45f965 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..cf8826750fce --- /dev/null +++ b/test/correctness/gpu_texture.cpp @@ -0,0 +1,113 @@ +#include "Halide.h" +#include + +using namespace Halide; +using namespace Halide::Internal; + +int main(int argc, char **argv) { + setenv("HL_JIT_TARGET", "host-opencl", 1); + setenv("HL_DEBUG_CODEGEN", "2", 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 input(100); + 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(x); + g(x) = param(x) + cast(f(2 * x)); + + g.gpu_tile(x, xi, 16); + + f.compute_root().store_in(memory_type).gpu_blocks(x); // store f as integer + g.store_in(memory_type); + + Buffer 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 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(x + y); + g(x) = param(x, x) + cast(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 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 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(x + y + z); + g(x) = param(x, x, x) + cast(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 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; + } + } + } + } + + printf("Success!\n"); + return 0; +} From 25d31da5e0c055a53718a5ca089ac80958894216 Mon Sep 17 00:00:00 2001 From: John Laxson Date: Wed, 23 Sep 2020 19:48:45 -0600 Subject: [PATCH 02/14] woof --- src/CodeGen_OpenCL_Dev.cpp | 240 +++++++++++++++--------------- src/DeviceArgument.cpp | 2 - src/InjectHostDevBufferCopies.cpp | 3 +- 3 files changed, 122 insertions(+), 123 deletions(-) diff --git a/src/CodeGen_OpenCL_Dev.cpp b/src/CodeGen_OpenCL_Dev.cpp index a9e2ca3b8d08..02938615e1ac 100644 --- a/src/CodeGen_OpenCL_Dev.cpp +++ b/src/CodeGen_OpenCL_Dev.cpp @@ -1,7 +1,7 @@ #include +#include #include #include -#include #include "CSE.h" #include "CodeGen_Internal.h" @@ -245,126 +245,126 @@ void CodeGen_OpenCL_Dev::CodeGen_OpenCL_C::visit(const Call *op) { } } else if (op->is_intrinsic(Call::image_load_texture)) { // image_load(, , , , , - // , , ) - int dims = (op->args.size() - 2) / 2; - internal_assert(dims >= 1 && dims <= 3); - const StringImm *string_imm = op->args[0].as(); - if (!string_imm) { - internal_assert(op->args[0].as()); - string_imm = op->args[0].as()->value.as(); - } - 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()); - - string type_suffix; - if (op->type.is_int()) { - type_suffix = "i"; - } else if (op->type.is_uint()) { - type_suffix = "ui"; - } else if (op->type.is_float()) { - type_suffix = "f"; - } else { - internal_error << "Invalid type for read_image: " << op->type << "\n"; - } - - std::array coord; - for (int i = 0; i < dims; i++) { - coord[i] = print_expr(op->args[i*2 + 2]); - } - vector 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" << type_suffix << "(" << 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(, , , , , ) - const StringImm *string_imm = op->args[0].as(); - if (!string_imm) { - internal_assert(op->args[0].as()); - string_imm = op->args[0].as()->value.as(); - } - 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()); - - string type_suffix; - if (op->type.is_int()) { - type_suffix = "i"; - } else if (op->type.is_uint()) { - type_suffix = "ui"; - } else if (op->type.is_float()) { - type_suffix = "f"; - } else { - internal_error << "Invalid type for write_image: " << op->type << "\n"; - } - - std::array 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" << type_suffix << "(" << 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"; + // , , ) + int dims = (op->args.size() - 2) / 2; + internal_assert(dims >= 1 && dims <= 3); + const StringImm *string_imm = op->args[0].as(); + if (!string_imm) { + internal_assert(op->args[0].as()); + string_imm = op->args[0].as()->value.as(); + } + 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()); + + string type_suffix; + if (op->type.is_int()) { + type_suffix = "i"; + } else if (op->type.is_uint()) { + type_suffix = "ui"; + } else if (op->type.is_float()) { + type_suffix = "f"; + } else { + internal_error << "Invalid type for read_image: " << op->type << "\n"; + } + + std::array coord; + for (int i = 0; i < dims; i++) { + coord[i] = print_expr(op->args[i * 2 + 2]); + } + vector 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" << type_suffix << "(" << 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(, , , , , ) + const StringImm *string_imm = op->args[0].as(); + if (!string_imm) { + internal_assert(op->args[0].as()); + string_imm = op->args[0].as()->value.as(); + } + 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()); + + string type_suffix; + if (op->type.is_int()) { + type_suffix = "i"; + } else if (op->type.is_uint()) { + type_suffix = "ui"; + } else if (op->type.is_float()) { + type_suffix = "f"; + } else { + internal_error << "Invalid type for write_image: " << op->type << "\n"; + } + + std::array 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" << type_suffix << "(" << 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(); - } + stream << write_image.str(); + } } else { CodeGen_C::visit(op); } diff --git a/src/DeviceArgument.cpp b/src/DeviceArgument.cpp index 19c245ceb3d4..faefbf908013 100644 --- a/src/DeviceArgument.cpp +++ b/src/DeviceArgument.cpp @@ -52,8 +52,6 @@ void HostClosure::visit(const Call *op) { internal_assert(string_imm); - - std::string bufname = string_imm->value; Buffer &ref = buffers[bufname]; ref.type = op->type; diff --git a/src/InjectHostDevBufferCopies.cpp b/src/InjectHostDevBufferCopies.cpp index eecd7d764b70..b745859fd9c3 100644 --- a/src/InjectHostDevBufferCopies.cpp +++ b/src/InjectHostDevBufferCopies.cpp @@ -573,7 +573,8 @@ class InjectBufferCopies : public IRMutator { public: InjectCombinedAllocation(string b, Type t, vector e, Expr c, Expr d) - : buffer(std::move(b)), type(t), extents(std::move(e)), condition(std::move(c)), device_interface(d) { + : buffer(std::move(b)), type(t), extents(std::move(e)), + condition(std::move(c)), device_interface(std::move(d)) { } }; From 76895bdacda70136c76a8d990dbe0a22b70a6e60 Mon Sep 17 00:00:00 2001 From: John Laxson Date: Wed, 23 Sep 2020 20:31:03 -0600 Subject: [PATCH 03/14] Handle mins --- src/StorageFlattening.cpp | 4 ++-- test/correctness/gpu_texture.cpp | 37 ++++++++++++++++++++++++++++++-- 2 files changed, 37 insertions(+), 4 deletions(-) diff --git a/src/StorageFlattening.cpp b/src/StorageFlattening.cpp index 18b67a9838d4..2145bd385705 100644 --- a/src/StorageFlattening.cpp +++ b/src/StorageFlattening.cpp @@ -260,7 +260,8 @@ class FlattenDimensions : public IRMutator { args[0] = op->name; args[1] = buffer_var; for (size_t i = 0; i < op->args.size(); i++) { - args.push_back(op->args[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, @@ -278,7 +279,6 @@ class FlattenDimensions : public IRMutator { debug(2) << " load call to " << op->name << " " << textures.count(op->name) << "\n"; if (op->param.defined()) { - debug(2) << " is param: " << " " << op->param.name() << " " << "\n"; diff --git a/test/correctness/gpu_texture.cpp b/test/correctness/gpu_texture.cpp index cf8826750fce..69bd96735ecf 100644 --- a/test/correctness/gpu_texture.cpp +++ b/test/correctness/gpu_texture.cpp @@ -5,8 +5,7 @@ using namespace Halide; using namespace Halide::Internal; int main(int argc, char **argv) { - setenv("HL_JIT_TARGET", "host-opencl", 1); - setenv("HL_DEBUG_CODEGEN", "2", 1); + // setenv("HL_JIT_TARGET", "host-opencl-debug", 1); Target t = get_jit_target_from_environment(); @@ -106,6 +105,40 @@ int main(int argc, char **argv) { } } } + { + // 1D offset + Buffer 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(x); + g(x) = param(x) + cast(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 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; + } + } + return 0; + } } printf("Success!\n"); From ef4e9b0e1ebab21a42e6668fe85175723b237b83 Mon Sep 17 00:00:00 2001 From: John Laxson Date: Wed, 23 Sep 2020 20:49:50 -0600 Subject: [PATCH 04/14] cleanup --- src/CodeGen_OpenCL_Dev.cpp | 42 ++++++++++++++------------------ src/DeviceArgument.h | 4 +++ src/runtime/CMakeLists.txt | 1 - src/runtime/opencl.cpp | 27 +++++++++----------- test/correctness/gpu_texture.cpp | 1 - 5 files changed, 33 insertions(+), 42 deletions(-) diff --git a/src/CodeGen_OpenCL_Dev.cpp b/src/CodeGen_OpenCL_Dev.cpp index 02938615e1ac..a21f2810de4f 100644 --- a/src/CodeGen_OpenCL_Dev.cpp +++ b/src/CodeGen_OpenCL_Dev.cpp @@ -175,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()) { @@ -258,17 +273,6 @@ void CodeGen_OpenCL_Dev::CodeGen_OpenCL_C::visit(const Call *op) { internal_assert(arg_type.lanes() <= 16); internal_assert(arg_type.lanes() == op->type.lanes()); - string type_suffix; - if (op->type.is_int()) { - type_suffix = "i"; - } else if (op->type.is_uint()) { - type_suffix = "ui"; - } else if (op->type.is_float()) { - type_suffix = "f"; - } else { - internal_error << "Invalid type for read_image: " << op->type << "\n"; - } - std::array coord; for (int i = 0; i < dims; i++) { coord[i] = print_expr(op->args[i * 2 + 2]); @@ -277,7 +281,7 @@ void CodeGen_OpenCL_Dev::CodeGen_OpenCL_C::visit(const Call *op) { // 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" << type_suffix << "(" << print_name(string_imm->value) << ", "; + 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: @@ -327,17 +331,6 @@ void CodeGen_OpenCL_Dev::CodeGen_OpenCL_C::visit(const Call *op) { Type value_type = op->args.back().type(); internal_assert(arg_type.lanes() == value_type.lanes()); - string type_suffix; - if (op->type.is_int()) { - type_suffix = "i"; - } else if (op->type.is_uint()) { - type_suffix = "ui"; - } else if (op->type.is_float()) { - type_suffix = "f"; - } else { - internal_error << "Invalid type for write_image: " << op->type << "\n"; - } - std::array coord; for (int i = 0; i < dims; i++) { coord[i] = print_expr(op->args[i + 2]); @@ -346,7 +339,8 @@ void CodeGen_OpenCL_Dev::CodeGen_OpenCL_C::visit(const Call *op) { // 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" << type_suffix << "(" << print_name(string_imm->value) << ", "; + 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: diff --git a/src/DeviceArgument.h b/src/DeviceArgument.h index 860e3274af52..8666650787d9 100644 --- a/src/DeviceArgument.h +++ b/src/DeviceArgument.h @@ -36,6 +36,10 @@ 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. diff --git a/src/runtime/CMakeLists.txt b/src/runtime/CMakeLists.txt index 1bb6bc21a9f6..91a59716e63e 100644 --- a/src/runtime/CMakeLists.txt +++ b/src/runtime/CMakeLists.txt @@ -47,7 +47,6 @@ set(RUNTIME_CPP msan msan_stubs opencl - opencl_image opengl opengl_egl_context opengl_glx_context diff --git a/src/runtime/opencl.cpp b/src/runtime/opencl.cpp index a26941b98166..87a4843930b3 100644 --- a/src/runtime/opencl.cpp +++ b/src/runtime/opencl.cpp @@ -11,8 +11,7 @@ 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 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" @@ -1593,22 +1592,16 @@ WEAK int halide_opencl_image_device_malloc(void *user_context, halide_buffer_t * 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 = 0; //buf->dim[1].stride * buf->type.bytes(); - desc.image_slice_pitch = 0; + 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.num_mip_levels = 0; desc.num_samples = 0; desc.buffer = NULL; debug(user_context) << " desc=(\n"; debug(user_context) << " " << (int)desc.image_type << ",\n"; - - debug(user_context) << " " << (int)desc.image_width << ",\n"; - debug(user_context) << " " << (int)desc.image_height << ",\n"; - debug(user_context) << " " << (int)desc.image_depth << ",\n"; - - debug(user_context) << " " << (int)desc.image_array_size << ",\n"; - debug(user_context) << " " << (int)desc.image_row_pitch << ",\n"; - debug(user_context) << " " << (int)desc.image_slice_pitch << ",\n"; + debug(user_context) << " " << (int)desc.image_width << ", " << (int)desc.image_height << ", " << (int)desc.image_depth << ",\n"; + debug(user_context) << " " << (int)desc.image_array_size << ", " << (int)desc.image_row_pitch << ", " << (int)desc.image_slice_pitch << ",\n"; debug(user_context) << " " << (void *)desc.buffer << ")\n"; cl_int err; @@ -1704,10 +1697,11 @@ WEAK int halide_opencl_image_buffer_copy(void *user_context, struct halide_buffe static_cast(dst->dim[0].extent), dim >= 2 ? static_cast(dst->dim[1].extent) : 1, dim >= 3 ? static_cast(dst->dim[1].extent) : 1}; - int pitch = dst->dimensions >= 2 ? dst->dim[1].stride * dst->type.bytes() : 0; + 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; err = clEnqueueReadImage(ctx.cmd_queue, ((device_handle *)c.src)->mem, CL_FALSE, offset, region, - 0, 0, + row_pitch, slice_pitch, dst->host, 0, NULL, NULL); } else if (from_host && !to_host) { int dim = src->dimensions; @@ -1716,9 +1710,10 @@ WEAK int halide_opencl_image_buffer_copy(void *user_context, struct halide_buffe static_cast(src->dim[0].extent), dim >= 2 ? static_cast(src->dim[1].extent) : 1, dim >= 3 ? static_cast(src->dim[1].extent) : 1}; - int pitch = dim >= 2 ? src->dim[1].stride * src->type.bytes() : 0; + 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; err = clEnqueueWriteImage(ctx.cmd_queue, ((device_handle *)c.dst)->mem, - CL_FALSE, offset, region, 0, 0, src->host, + CL_FALSE, offset, region, row_pitch, slice_pitch, src->host, 0, NULL, NULL); } else if (!from_host && !to_host) { halide_assert(user_context, false && "image to image copies not implemented"); diff --git a/test/correctness/gpu_texture.cpp b/test/correctness/gpu_texture.cpp index 69bd96735ecf..9a8fb5805568 100644 --- a/test/correctness/gpu_texture.cpp +++ b/test/correctness/gpu_texture.cpp @@ -137,7 +137,6 @@ int main(int argc, char **argv) { return -1; } } - return 0; } } From e7e8ac178f7afb2ff5ba9d7a649451c86c8b3281 Mon Sep 17 00:00:00 2001 From: John Laxson Date: Wed, 23 Sep 2020 20:56:30 -0600 Subject: [PATCH 05/14] runtime helpers --- src/runtime/HalideRuntimeOpenCL.h | 5 +++++ src/runtime/opencl.cpp | 32 +++++++++++++++++++++++++++++-- src/runtime/runtime_api.cpp | 1 + 3 files changed, 36 insertions(+), 2 deletions(-) diff --git a/src/runtime/HalideRuntimeOpenCL.h b/src/runtime/HalideRuntimeOpenCL.h index e3b27f84c800..18d6c03545cf 100644 --- a/src/runtime/HalideRuntimeOpenCL.h +++ b/src/runtime/HalideRuntimeOpenCL.h @@ -90,6 +90,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 87a4843930b3..bb409f2afc1d 100644 --- a/src/runtime/opencl.cpp +++ b/src/runtime/opencl.cpp @@ -1203,7 +1203,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(); @@ -1215,7 +1216,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; } @@ -1758,6 +1760,32 @@ WEAK int halide_opencl_image_device_and_host_malloc(void *user_context, struct h 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; +} } namespace Halide { diff --git a/src/runtime/runtime_api.cpp b/src/runtime/runtime_api.cpp index 76d1b919f755..e320692cd6cc 100644 --- a/src/runtime/runtime_api.cpp +++ b/src/runtime/runtime_api.cpp @@ -136,6 +136,7 @@ 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_wrap_cl_mem, (void *)&halide_opencl_initialize_kernels, (void *)&halide_opencl_run, (void *)&halide_opencl_set_build_options, From a54d8224ea64590a458e3e7343e0a57d87c9ffc2 Mon Sep 17 00:00:00 2001 From: John Laxson Date: Wed, 23 Sep 2020 20:57:44 -0600 Subject: [PATCH 06/14] bye --- src/runtime/opencl_image.cpp | 0 1 file changed, 0 insertions(+), 0 deletions(-) delete mode 100644 src/runtime/opencl_image.cpp diff --git a/src/runtime/opencl_image.cpp b/src/runtime/opencl_image.cpp deleted file mode 100644 index e69de29bb2d1..000000000000 From 33d589de24e4c1ab4d8d75d97ab36dfee2fbd5fa Mon Sep 17 00:00:00 2001 From: John Laxson Date: Wed, 23 Sep 2020 21:07:03 -0600 Subject: [PATCH 07/14] add remaining device interface methods --- src/CodeGen_OpenCL_Dev.cpp | 8 +++----- src/runtime/opencl.cpp | 36 ++++++++++++++++++++++++++++-------- 2 files changed, 31 insertions(+), 13 deletions(-) diff --git a/src/CodeGen_OpenCL_Dev.cpp b/src/CodeGen_OpenCL_Dev.cpp index a21f2810de4f..0557bc0a0909 100644 --- a/src/CodeGen_OpenCL_Dev.cpp +++ b/src/CodeGen_OpenCL_Dev.cpp @@ -707,11 +707,9 @@ void CodeGen_OpenCL_Dev::CodeGen_OpenCL_C::visit(const Free *op) { return; } else { // Should have been freed internally - if (allocations.contains(op->name)) { - internal_assert(allocations.contains(op->name)); - allocations.pop(op->name); - stream << get_indent() << "#undef " << get_memory_space(op->name) << "\n"; - } + internal_assert(allocations.contains(op->name)); + allocations.pop(op->name); + stream << get_indent() << "#undef " << get_memory_space(op->name) << "\n"; } } diff --git a/src/runtime/opencl.cpp b/src/runtime/opencl.cpp index bb409f2afc1d..4116795c6657 100644 --- a/src/runtime/opencl.cpp +++ b/src/runtime/opencl.cpp @@ -619,8 +619,6 @@ WEAK int halide_opencl_initialize_kernels(void *user_context, void **state_ptr, return ctx.error_code; } - // std::cout << src; - #ifdef DEBUG_RUNTIME uint64_t t_before = halide_current_time_ns(user_context); #endif @@ -1204,7 +1202,7 @@ WEAK int halide_opencl_detach_cl_mem(void *user_context, halide_buffer_t *buf) { return 0; } halide_assert(user_context, buf->device_interface == &opencl_device_interface || - buf->device_interface == &opencl_image_device_interface); + buf->device_interface == &opencl_image_device_interface); free((device_handle *)buf->device); buf->device = 0; buf->device_interface->impl->release_module(); @@ -1217,7 +1215,7 @@ WEAK uintptr_t halide_opencl_get_cl_mem(void *user_context, halide_buffer_t *buf return 0; } halide_assert(user_context, buf->device_interface == &opencl_device_interface || - buf->device_interface == &opencl_image_device_interface); + buf->device_interface == &opencl_image_device_interface); return (uintptr_t)((device_handle *)buf->device)->mem; } @@ -1786,6 +1784,28 @@ WEAK int halide_opencl_image_wrap_cl_mem(void *user_context, struct halide_buffe #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 { @@ -1805,10 +1825,10 @@ WEAK halide_device_interface_impl_t opencl_image_device_interface_impl = { halide_opencl_image_device_and_host_malloc, halide_opencl_image_device_and_host_free, halide_opencl_image_buffer_copy, - nullptr, //halide_opencl_image_device_crop, - nullptr, //halide_opencl_image_device_slice, - nullptr, //halide_opencl_image_device_release_crop, - halide_opencl_wrap_cl_mem, + 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, }; From 63397d4d8ee0e9688c12d6d8c062bab4a977eada Mon Sep 17 00:00:00 2001 From: John Laxson Date: Wed, 23 Sep 2020 23:12:27 -0600 Subject: [PATCH 08/14] cleanup --- src/runtime/opencl.cpp | 45 ++++++++++++++++++++++++++++-------------- 1 file changed, 30 insertions(+), 15 deletions(-) diff --git a/src/runtime/opencl.cpp b/src/runtime/opencl.cpp index 4116795c6657..f18efd3cd09c 100644 --- a/src/runtime/opencl.cpp +++ b/src/runtime/opencl.cpp @@ -1592,17 +1592,26 @@ WEAK int halide_opencl_image_device_malloc(void *user_context, halide_buffer_t * 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 = 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=(\n"; - debug(user_context) << " " << (int)desc.image_type << ",\n"; - debug(user_context) << " " << (int)desc.image_width << ", " << (int)desc.image_height << ", " << (int)desc.image_depth << ",\n"; - debug(user_context) << " " << (int)desc.image_array_size << ", " << (int)desc.image_row_pitch << ", " << (int)desc.image_slice_pitch << ",\n"; - debug(user_context) << " " << (void *)desc.buffer << ")\n"; + 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 << " "; @@ -1696,12 +1705,16 @@ WEAK int halide_opencl_image_buffer_copy(void *user_context, struct halide_buffe size_t region[] = { static_cast(dst->dim[0].extent), dim >= 2 ? static_cast(dst->dim[1].extent) : 1, - dim >= 3 ? static_cast(dst->dim[1].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; + dim >= 3 ? static_cast(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, slice_pitch, + /* row_pitch */ 0, /* slice_pitch */ 0, dst->host, 0, NULL, NULL); } else if (from_host && !to_host) { int dim = src->dimensions; @@ -1709,11 +1722,13 @@ WEAK int halide_opencl_image_buffer_copy(void *user_context, struct halide_buffe size_t region[] = { static_cast(src->dim[0].extent), dim >= 2 ? static_cast(src->dim[1].extent) : 1, - dim >= 3 ? static_cast(src->dim[1].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; + dim >= 3 ? static_cast(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, slice_pitch, src->host, + 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"); From f5d0427be4e0a80c3597fd57f12e0f6eacf0b792 Mon Sep 17 00:00:00 2001 From: John Laxson Date: Thu, 24 Sep 2020 11:39:38 -0600 Subject: [PATCH 09/14] move store_in to OutputImageParam --- src/Generator.h | 1 + src/ImageParam.cpp | 4 ---- src/ImageParam.h | 4 ---- src/OutputImageParam.cpp | 5 +++++ src/OutputImageParam.h | 4 ++++ 5 files changed, 10 insertions(+), 8 deletions(-) diff --git a/src/Generator.h b/src/Generator.h index ee82f71a8cc6..2228dfaeac65 100644 --- a/src/Generator.h +++ b/src/Generator.h @@ -2542,6 +2542,7 @@ class GeneratorOutput_Buffer : public GeneratorOutputImpl { 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/ImageParam.cpp b/src/ImageParam.cpp index c3a8db3c604d..cda49b501c4c 100644 --- a/src/ImageParam.cpp +++ b/src/ImageParam.cpp @@ -95,8 +95,4 @@ ImageParam &ImageParam::add_trace_tag(const std::string &trace_tag) { return *this; } -void ImageParam::store_in(MemoryType type) { - param.store_in(type); -} - } // namespace Halide diff --git a/src/ImageParam.h b/src/ImageParam.h index f8b1ce7f0d02..d4383bf4ed7f 100644 --- a/src/ImageParam.h +++ b/src/ImageParam.h @@ -32,8 +32,6 @@ class ImageParam : public OutputImageParam { /** Helper function to initialize the Func representation of this ImageParam. */ Func create_func() const; - MemoryType memory_type = MemoryType::Auto; - public: /** Construct a nullptr image parameter handle. */ ImageParam() = default; @@ -135,8 +133,6 @@ class ImageParam : public OutputImageParam { /** Add a trace tag to this ImageParam's Func. */ ImageParam &add_trace_tag(const std::string &trace_tag); - - void store_in(MemoryType type); }; } // namespace Halide 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 From 41abcf2364a8d8d425d573f306713cf5bc85f2d1 Mon Sep 17 00:00:00 2001 From: John Laxson Date: Thu, 24 Sep 2020 12:01:44 -0600 Subject: [PATCH 10/14] Revert "move store_in to OutputImageParam" This reverts commit f5d0427be4e0a80c3597fd57f12e0f6eacf0b792. --- src/Generator.h | 1 - src/ImageParam.cpp | 4 ++++ src/ImageParam.h | 4 ++++ src/OutputImageParam.cpp | 5 ----- src/OutputImageParam.h | 4 ---- 5 files changed, 8 insertions(+), 10 deletions(-) diff --git a/src/Generator.h b/src/Generator.h index 2228dfaeac65..ee82f71a8cc6 100644 --- a/src/Generator.h +++ b/src/Generator.h @@ -2542,7 +2542,6 @@ class GeneratorOutput_Buffer : public GeneratorOutputImpl { 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/ImageParam.cpp b/src/ImageParam.cpp index cda49b501c4c..c3a8db3c604d 100644 --- a/src/ImageParam.cpp +++ b/src/ImageParam.cpp @@ -95,4 +95,8 @@ ImageParam &ImageParam::add_trace_tag(const std::string &trace_tag) { return *this; } +void ImageParam::store_in(MemoryType type) { + param.store_in(type); +} + } // namespace Halide diff --git a/src/ImageParam.h b/src/ImageParam.h index d4383bf4ed7f..f8b1ce7f0d02 100644 --- a/src/ImageParam.h +++ b/src/ImageParam.h @@ -32,6 +32,8 @@ class ImageParam : public OutputImageParam { /** Helper function to initialize the Func representation of this ImageParam. */ Func create_func() const; + MemoryType memory_type = MemoryType::Auto; + public: /** Construct a nullptr image parameter handle. */ ImageParam() = default; @@ -133,6 +135,8 @@ class ImageParam : public OutputImageParam { /** Add a trace tag to this ImageParam's Func. */ ImageParam &add_trace_tag(const std::string &trace_tag); + + void store_in(MemoryType type); }; } // namespace Halide diff --git a/src/OutputImageParam.cpp b/src/OutputImageParam.cpp index a59ff13c43f1..9702ef91b24e 100644 --- a/src/OutputImageParam.cpp +++ b/src/OutputImageParam.cpp @@ -102,9 +102,4 @@ 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 a3aeb0ca5d7b..4cb66ae40a70 100644 --- a/src/OutputImageParam.h +++ b/src/OutputImageParam.h @@ -117,10 +117,6 @@ 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 From 50301087c0219b0bacc2267a636131375a4efc52 Mon Sep 17 00:00:00 2001 From: John Laxson Date: Thu, 24 Sep 2020 11:39:38 -0600 Subject: [PATCH 11/14] move store_in to OutputImageParam --- src/Generator.h | 1 + src/ImageParam.cpp | 4 ---- src/ImageParam.h | 4 ---- src/OutputImageParam.cpp | 5 +++++ src/OutputImageParam.h | 4 ++++ 5 files changed, 10 insertions(+), 8 deletions(-) diff --git a/src/Generator.h b/src/Generator.h index ee82f71a8cc6..2228dfaeac65 100644 --- a/src/Generator.h +++ b/src/Generator.h @@ -2542,6 +2542,7 @@ class GeneratorOutput_Buffer : public GeneratorOutputImpl { 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/ImageParam.cpp b/src/ImageParam.cpp index c3a8db3c604d..cda49b501c4c 100644 --- a/src/ImageParam.cpp +++ b/src/ImageParam.cpp @@ -95,8 +95,4 @@ ImageParam &ImageParam::add_trace_tag(const std::string &trace_tag) { return *this; } -void ImageParam::store_in(MemoryType type) { - param.store_in(type); -} - } // namespace Halide diff --git a/src/ImageParam.h b/src/ImageParam.h index f8b1ce7f0d02..d4383bf4ed7f 100644 --- a/src/ImageParam.h +++ b/src/ImageParam.h @@ -32,8 +32,6 @@ class ImageParam : public OutputImageParam { /** Helper function to initialize the Func representation of this ImageParam. */ Func create_func() const; - MemoryType memory_type = MemoryType::Auto; - public: /** Construct a nullptr image parameter handle. */ ImageParam() = default; @@ -135,8 +133,6 @@ class ImageParam : public OutputImageParam { /** Add a trace tag to this ImageParam's Func. */ ImageParam &add_trace_tag(const std::string &trace_tag); - - void store_in(MemoryType type); }; } // namespace Halide 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 From 8382eff78f73b0686e629656d4f8d8318764c9dc Mon Sep 17 00:00:00 2001 From: John Laxson Date: Fri, 25 Sep 2020 10:12:30 -0600 Subject: [PATCH 12/14] Output support and generator test --- python_bindings/src/PyEnums.cpp | 5 +- python_bindings/src/PyImageParam.cpp | 1 + src/StorageFlattening.cpp | 10 +++- src/runtime/HalideRuntimeOpenCL.h | 1 + src/runtime/opencl.cpp | 2 +- src/runtime/runtime_api.cpp | 1 + test/correctness/gpu_texture.cpp | 5 +- test/generator/CMakeLists.txt | 4 ++ test/generator/gpu_texture_aottest.cpp | 64 ++++++++++++++++++++++++ test/generator/gpu_texture_generator.cpp | 30 +++++++++++ 10 files changed, 118 insertions(+), 5 deletions(-) create mode 100644 test/generator/gpu_texture_aottest.cpp create mode 100644 test/generator/gpu_texture_generator.cpp diff --git a/python_bindings/src/PyEnums.cpp b/python_bindings/src/PyEnums.cpp index 11b1cceec591..49ff5c80d6ca 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_(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/StorageFlattening.cpp b/src/StorageFlattening.cpp index 2145bd385705..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" @@ -239,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) @@ -280,7 +288,7 @@ class FlattenDimensions : public IRMutator { 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.name() << " " << op->param.memory_type() << "\n"; if (op->param.memory_type() == MemoryType::GPUTexture) { diff --git a/src/runtime/HalideRuntimeOpenCL.h b/src/runtime/HalideRuntimeOpenCL.h index 18d6c03545cf..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. */ diff --git a/src/runtime/opencl.cpp b/src/runtime/opencl.cpp index f18efd3cd09c..e9fd613f65c5 100644 --- a/src/runtime/opencl.cpp +++ b/src/runtime/opencl.cpp @@ -1658,7 +1658,7 @@ WEAK int halide_opencl_image_buffer_copy(void *user_context, struct halide_buffe dst_device_interface == &opencl_image_device_interface); if ((src->device_dirty() || src->host == NULL) && - src->device_interface != &opencl_device_interface) { + 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; diff --git a/src/runtime/runtime_api.cpp b/src/runtime/runtime_api.cpp index e320692cd6cc..6071932000da 100644 --- a/src/runtime/runtime_api.cpp +++ b/src/runtime/runtime_api.cpp @@ -136,6 +136,7 @@ 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, diff --git a/test/correctness/gpu_texture.cpp b/test/correctness/gpu_texture.cpp index 9a8fb5805568..c41339f29c6f 100644 --- a/test/correctness/gpu_texture.cpp +++ b/test/correctness/gpu_texture.cpp @@ -21,7 +21,7 @@ int main(int argc, char **argv) { // 1D stores/loads Buffer input(100); input.fill(10); - ImageParam param(Int(32), 1); + ImageParam param(Int(32), 1, "input"); param.set(input); param.store_in(memory_type); // check float stores @@ -35,7 +35,7 @@ int main(int argc, char **argv) { g.gpu_tile(x, xi, 16); f.compute_root().store_in(memory_type).gpu_blocks(x); // store f as integer - g.store_in(memory_type); + g.output_buffer().store_in(memory_type); Buffer out = g.realize(100); for (int x = 0; x < 100; x++) { @@ -45,6 +45,7 @@ int main(int argc, char **argv) { return -1; } } + return -0; } { // 2D stores/loads 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 +#include +#include +#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 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 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 { +public: + Input> input{"input", 2}; + + Output> 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) From b36aab3dddafc5e19260a7e7ed2c13f0494db9da Mon Sep 17 00:00:00 2001 From: John Laxson Date: Sat, 26 Sep 2020 08:28:18 -0600 Subject: [PATCH 13/14] fix bypassed test --- test/correctness/gpu_texture.cpp | 1 - 1 file changed, 1 deletion(-) diff --git a/test/correctness/gpu_texture.cpp b/test/correctness/gpu_texture.cpp index c41339f29c6f..3b1be49cbb5b 100644 --- a/test/correctness/gpu_texture.cpp +++ b/test/correctness/gpu_texture.cpp @@ -45,7 +45,6 @@ int main(int argc, char **argv) { return -1; } } - return -0; } { // 2D stores/loads From 3a1c03fe803c7e9c83f1c94b76945a0ebcf282d4 Mon Sep 17 00:00:00 2001 From: Steven Johnson Date: Tue, 29 Sep 2020 10:15:23 -0700 Subject: [PATCH 14/14] tickle buildbots --- src/Generator.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) 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"