From 26d00762bbfb598852c641cc8aced737e672ed88 Mon Sep 17 00:00:00 2001 From: John Laxson Date: Wed, 23 Sep 2020 18:54:55 -0600 Subject: [PATCH 01/29] 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/29] 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/29] 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/29] 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/29] 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/29] 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/29] 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/29] 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/29] 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/29] 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/29] 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/29] 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/29] 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 511082c7b424fb256600256d53af33741c164d1d Mon Sep 17 00:00:00 2001 From: John Laxson Date: Wed, 30 Sep 2020 11:03:23 -0600 Subject: [PATCH 14/29] build fixes --- src/runtime/opencl.cpp | 18 ++++-------------- test/correctness/gpu_texture.cpp | 2 -- 2 files changed, 4 insertions(+), 16 deletions(-) diff --git a/src/runtime/opencl.cpp b/src/runtime/opencl.cpp index e9fd613f65c5..efe40bba7047 100644 --- a/src/runtime/opencl.cpp +++ b/src/runtime/opencl.cpp @@ -1563,20 +1563,8 @@ WEAK int halide_opencl_image_device_malloc(void *user_context, halide_buffer_t * } 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) { @@ -1803,8 +1791,10 @@ WEAK int halide_opencl_image_wrap_cl_mem(void *user_context, struct halide_buffe 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; + for (int dim = 0; dim < src->dimensions; dim++) { + halide_assert(user_context, src->dim[dim] == dst->dim[dim] && "crop not supported on opencl image objects"); + } + return 0; } WEAK int halide_opencl_image_device_slice(void *user_context, diff --git a/test/correctness/gpu_texture.cpp b/test/correctness/gpu_texture.cpp index 3b1be49cbb5b..97af4097a0fe 100644 --- a/test/correctness/gpu_texture.cpp +++ b/test/correctness/gpu_texture.cpp @@ -5,8 +5,6 @@ using namespace Halide; using namespace Halide::Internal; int main(int argc, char **argv) { - // setenv("HL_JIT_TARGET", "host-opencl-debug", 1); - Target t = get_jit_target_from_environment(); if (!t.has_feature(halide_target_feature_opencl)) { From dbad5b10d49edfb0e8acd6d65c8b77b0bddc9125 Mon Sep 17 00:00:00 2001 From: John Laxson Date: Wed, 30 Sep 2020 11:04:52 -0600 Subject: [PATCH 15/29] comment --- test/correctness/gpu_texture.cpp | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/test/correctness/gpu_texture.cpp b/test/correctness/gpu_texture.cpp index 97af4097a0fe..613f6f78cf45 100644 --- a/test/correctness/gpu_texture.cpp +++ b/test/correctness/gpu_texture.cpp @@ -12,8 +12,7 @@ int main(int argc, char **argv) { return 0; } - // Check dynamic allocations per-block and per-thread into both - // shared and global + // Check dynamic allocations into Heap and Texture memory for (auto memory_type : {MemoryType::GPUTexture, MemoryType::Heap}) { { // 1D stores/loads From 979b05aad41c7e2200f0833a0a293b2a1c20677b Mon Sep 17 00:00:00 2001 From: John Laxson Date: Wed, 30 Sep 2020 12:22:51 -0600 Subject: [PATCH 16/29] assert me not --- src/runtime/opencl.cpp | 74 +++++++++++++++++++++++++++++++++--------- 1 file changed, 58 insertions(+), 16 deletions(-) diff --git a/src/runtime/opencl.cpp b/src/runtime/opencl.cpp index efe40bba7047..79a12f1c55c1 100644 --- a/src/runtime/opencl.cpp +++ b/src/runtime/opencl.cpp @@ -1567,6 +1567,19 @@ WEAK int halide_opencl_image_device_malloc(void *user_context, halide_buffer_t * debug(user_context) << " format=(" << format.image_channel_data_type << ", " << format.image_channel_order << ")\n"; + if (buf->dim[0].stride != 1) { + error(user_context) << "image buffer must be dense on inner dimension"; + return halide_error_code_device_malloc_failed; + } + if (buf->dimensions >= 2 && buf->dim[1].stride != buf->dim[0].extent) { + error(user_context) << "image buffer must be dense on inner dimension"; + return halide_error_code_device_malloc_failed; + } + if (buf->dimensions >= 3 && buf->dim[2].stride != buf->dim[0].extent * buf->dim[1].extent) { + error(user_context) << "image buffer must be dense on inner dimension"; + return halide_error_code_device_malloc_failed; + } + if (buf->dimensions == 1) { desc.image_type = CL_MEM_OBJECT_IMAGE1D; } else if (buf->dimensions == 2) { @@ -1574,7 +1587,8 @@ WEAK int halide_opencl_image_device_malloc(void *user_context, halide_buffer_t * } else if (buf->dimensions == 3) { desc.image_type = CL_MEM_OBJECT_IMAGE3D; } else { - halide_assert(user_context, buf->dimensions >= 1 && buf->dimensions <= 3); + error(user_context) << "image buffer must have 1-3 dimensions"; + return halide_error_code_device_malloc_failed; } desc.image_width = buf->dim[0].extent; desc.image_height = buf->dimensions >= 2 ? buf->dim[1].extent : 1; @@ -1583,9 +1597,7 @@ WEAK int halide_opencl_image_device_malloc(void *user_context, halide_buffer_t * // 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; @@ -1685,8 +1697,10 @@ WEAK int halide_opencl_image_buffer_copy(void *user_context, struct halide_buffe << " -> " << (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 (src->size_in_bytes() != dst->size_in_bytes() || c.chunk_size != src->size_in_bytes()) { + error(user_context) << "image buffer copies must be for whole buffer"; + return halide_error_code_device_buffer_copy_failed; + } if (!from_host && to_host) { int dim = dst->dimensions; size_t offset[] = {0, 0, 0}; @@ -1697,9 +1711,14 @@ WEAK int halide_opencl_image_buffer_copy(void *user_context, struct halide_buffe // 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); - + if (dst->dimensions >= 2 && dst->dim[1].stride != dst->dim[0].extent) { + error(user_context) << "image buffer copies must be dense on inner dimension"; + return halide_error_code_device_buffer_copy_failed; + } + if (dst->dimensions >= 3 && dst->dim[2].stride != dst->dim[0].extent * dst->dim[1].extent) { + error(user_context) << "image buffer copies must be dense on inner dimension"; + return halide_error_code_device_buffer_copy_failed; + } err = clEnqueueReadImage(ctx.cmd_queue, ((device_handle *)c.src)->mem, CL_FALSE, offset, region, /* row_pitch */ 0, /* slice_pitch */ 0, @@ -1713,13 +1732,20 @@ WEAK int halide_opencl_image_buffer_copy(void *user_context, struct halide_buffe 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); + if (src->dimensions >= 2 && src->dim[1].stride != src->dim[0].extent) { + error(user_context) << "image buffer copies must be dense on inner dimension"; + return halide_error_code_device_buffer_copy_failed; + } + if (src->dimensions >= 3 && src->dim[2].stride != src->dim[0].extent * src->dim[1].extent) { + error(user_context) << "image buffer copies must be dense on inner dimension"; + return halide_error_code_device_buffer_copy_failed; + } err = clEnqueueWriteImage(ctx.cmd_queue, ((device_handle *)c.dst)->mem, CL_FALSE, offset, region, /* row_pitch */ 0, /* slice_pitch */ 0, src->host, 0, NULL, NULL); } else if (!from_host && !to_host) { - halide_assert(user_context, false && "image to image copies not implemented"); + error(user_context) << "image to image copies not implemented"; + return halide_error_code_device_buffer_copy_failed; // 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); @@ -1771,6 +1797,19 @@ WEAK int halide_opencl_image_wrap_cl_mem(void *user_context, struct halide_buffe if (dev_handle == NULL) { return halide_error_code_out_of_memory; } + + cl_int mem_type = 0; + cl_int result = clGetMemObjectInfo((cl_mem)mem, CL_MEM_TYPE, sizeof(mem_type), &mem_type, NULL); + if (result != CL_SUCCESS || (mem_type != CL_MEM_OBJECT_IMAGE1D && + mem_type != CL_MEM_OBJECT_IMAGE2D && + mem_type != CL_MEM_OBJECT_IMAGE3D)) { + error(user_context) << "CL: Bad device pointer passed to halide_opencl_image_wrap_cl_mem: " << (void *)mem + << ": clGetMemObjectInfo returned " + << get_opencl_error_name(result) + << " with type " << mem_type; + return halide_error_code_device_wrap_native_failed; + } + dev_handle->mem = (cl_mem)mem; dev_handle->offset = 0; buf->device = (uint64_t)dev_handle; @@ -1792,7 +1831,10 @@ WEAK int halide_opencl_image_device_crop(void *user_context, const struct halide_buffer_t *src, struct halide_buffer_t *dst) { for (int dim = 0; dim < src->dimensions; dim++) { - halide_assert(user_context, src->dim[dim] == dst->dim[dim] && "crop not supported on opencl image objects"); + if (src->dim[dim] != dst->dim[dim]) { + error(user_context) << "crop not supported on opencl image objects"; + return halide_error_code_device_crop_unsupported; + } } return 0; } @@ -1802,14 +1844,14 @@ WEAK int halide_opencl_image_device_slice(void *user_context, 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; + error(user_context) << "slice not supported on opencl image objects"; + return halide_error_code_device_crop_unsupported; } 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; + error(user_context) << "crop not supported on opencl image objects"; + return halide_error_code_device_crop_unsupported; } } From 4141309ccb4f12636b46d3aedd6c2c811aef38a1 Mon Sep 17 00:00:00 2001 From: John Laxson Date: Wed, 30 Sep 2020 14:22:53 -0600 Subject: [PATCH 17/29] more assert/error cleanup --- src/runtime/opencl.cpp | 25 ++++++++----------------- 1 file changed, 8 insertions(+), 17 deletions(-) diff --git a/src/runtime/opencl.cpp b/src/runtime/opencl.cpp index 79a12f1c55c1..221178ef21a9 100644 --- a/src/runtime/opencl.cpp +++ b/src/runtime/opencl.cpp @@ -1532,6 +1532,7 @@ WEAK int halide_opencl_image_device_malloc(void *user_context, halide_buffer_t * cl_image_desc desc; struct halide_type_t type = buf->type; + format.image_channel_data_type = -1; if (type.code == halide_type_int) { if (type.bits == 8) { format.image_channel_data_type = CL_SIGNED_INT8; @@ -1539,8 +1540,6 @@ WEAK int halide_opencl_image_device_malloc(void *user_context, halide_buffer_t * 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) { @@ -1549,33 +1548,25 @@ WEAK int halide_opencl_image_device_malloc(void *user_context, halide_buffer_t * 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"); + } + if (format.image_channel_data_type == -1) { + error(user_context) << "Unhandled datatype for opencl texture object: " << type; + return halide_error_code_device_malloc_failed; } format.image_channel_order = CL_R; debug(user_context) << " format=(" << format.image_channel_data_type << ", " << format.image_channel_order << ")\n"; - if (buf->dim[0].stride != 1) { - error(user_context) << "image buffer must be dense on inner dimension"; - return halide_error_code_device_malloc_failed; - } - if (buf->dimensions >= 2 && buf->dim[1].stride != buf->dim[0].extent) { - error(user_context) << "image buffer must be dense on inner dimension"; - return halide_error_code_device_malloc_failed; - } - if (buf->dimensions >= 3 && buf->dim[2].stride != buf->dim[0].extent * buf->dim[1].extent) { + if (buf->dim[0].stride != 1 || + (buf->dimensions >= 2 && buf->dim[1].stride != buf->dim[0].extent) || + (buf->dimensions >= 3 && buf->dim[2].stride != buf->dim[0].extent * buf->dim[1].extent)) { error(user_context) << "image buffer must be dense on inner dimension"; return halide_error_code_device_malloc_failed; } From c0f09a8941ba0e91ee22adb955da8dbcfc7cb83f Mon Sep 17 00:00:00 2001 From: John Laxson Date: Thu, 1 Oct 2020 19:24:25 -0600 Subject: [PATCH 18/29] intrinsic rename, cleanup --- src/Closure.h | 4 ++-- src/CodeGen_OpenCL_Dev.cpp | 6 +++--- src/DeviceArgument.cpp | 22 ++++++++++----------- src/DeviceArgument.h | 10 +++++----- src/DeviceInterface.cpp | 4 ++-- src/DeviceInterface.h | 3 ++- src/IR.cpp | 2 -- src/IR.h | 2 -- src/InjectHostDevBufferCopies.cpp | 32 ++++++++++++++----------------- src/StorageFlattening.cpp | 6 ++---- src/runtime/opencl.cpp | 13 +------------ 11 files changed, 41 insertions(+), 63 deletions(-) diff --git a/src/Closure.h b/src/Closure.h index 858709f4205b..d4252b4829d9 100644 --- a/src/Closure.h +++ b/src/Closure.h @@ -56,13 +56,13 @@ class Closure : public IRVisitor { bool write; /** The buffer is a texture */ - bool texture; + MemoryType memory_type; /** The size of the buffer if known, otherwise zero. */ size_t size; Buffer() - : dimensions(0), read(false), write(false), texture(false), size(0) { + : dimensions(0), read(false), write(false), memory_type(MemoryType::Auto), size(0) { } }; diff --git a/src/CodeGen_OpenCL_Dev.cpp b/src/CodeGen_OpenCL_Dev.cpp index 0557bc0a0909..1e529902aea8 100644 --- a/src/CodeGen_OpenCL_Dev.cpp +++ b/src/CodeGen_OpenCL_Dev.cpp @@ -258,7 +258,7 @@ 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)) { + } else if (op->is_intrinsic(Call::image_load)) { // image_load(, , , , , // , , ) int dims = (op->args.size() - 2) / 2; @@ -316,7 +316,7 @@ void CodeGen_OpenCL_Dev::CodeGen_OpenCL_C::visit(const Call *op) { // Widen to the correct type print_assignment(op->type, "convert_" + print_type(op->type) + "(" + id + ")"); } - } else if (op->is_intrinsic(Call::image_store_texture)) { + } else if (op->is_intrinsic(Call::image_store)) { // image_store(, , , , , ) const StringImm *string_imm = op->args[0].as(); if (!string_imm) { @@ -882,7 +882,7 @@ 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) { - if (args[i].is_texture) { + if (args[i].memory_type == MemoryType::GPUTexture) { int dims = args[i].dimensions; internal_assert(dims >= 1 && dims <= 3) << "dims = " << dims << "\n"; if (args[i].read && args[i].write) { diff --git a/src/DeviceArgument.cpp b/src/DeviceArgument.cpp index faefbf908013..3a1d6cb5bb43 100644 --- a/src/DeviceArgument.cpp +++ b/src/DeviceArgument.cpp @@ -15,17 +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, false, v.second, 0); + res.emplace_back(v.first, false, MemoryType::Auto, 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) << " "; + if (b.second.memory_type == MemoryType::GPUTexture) debug(2) << " "; debug(2) << " dims=" << (int)b.second.dimensions; debug(2) << "\n"; - DeviceArgument arg(b.first, true, b.second.texture, b.second.type, b.second.dimensions, b.second.size); + DeviceArgument arg(b.first, true, b.second.memory_type, b.second.type, b.second.dimensions, b.second.size); arg.read = b.second.read; arg.write = b.second.write; res.push_back(arg); @@ -36,10 +36,8 @@ 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_texture)) { + op->is_intrinsic(Call::image_store)) { // The argument to the call is either a StringImm or a broadcasted // StringImm if this is part of a vectorized expression @@ -55,17 +53,17 @@ void HostClosure::visit(const Call *op) { std::string bufname = string_imm->value; Buffer &ref = buffers[bufname]; ref.type = op->type; - ref.texture = op->is_intrinsic(Call::image_load_texture) || - op->is_intrinsic(Call::image_store_texture); + ref.memory_type = op->is_intrinsic(Call::image_load) || + op->is_intrinsic(Call::image_store) ? + MemoryType::GPUTexture : + MemoryType::Auto; 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::image_load)) { 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_texture)) { + op->is_intrinsic(Call::image_store)) { ref.write = true; ref.dimensions = op->args.size() - 3; } diff --git a/src/DeviceArgument.h b/src/DeviceArgument.h index 8666650787d9..d1e8ed4cb77c 100644 --- a/src/DeviceArgument.h +++ b/src/DeviceArgument.h @@ -36,11 +36,11 @@ struct DeviceArgument { */ bool is_buffer; - /** If is_buffer == true and is_texture == true, this argument should be + /** If is_buffer == true and memory_type == GPUTexture, this argument should be * passed and accessed through texture sampler operations instead of * directly as a memory array */ - bool is_texture; + MemoryType memory_type; /** 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) */ @@ -72,7 +72,7 @@ struct DeviceArgument { DeviceArgument() : is_buffer(false), - is_texture(false), + memory_type(MemoryType::Auto), dimensions(0), size(0), packed_index(0), @@ -82,13 +82,13 @@ struct DeviceArgument { DeviceArgument(const std::string &_name, bool _is_buffer, - bool _is_texture, + MemoryType _mem, Type _type, uint8_t _dimensions, size_t _size = 0) : name(_name), is_buffer(_is_buffer), - is_texture(_is_texture), + memory_type(_mem), dimensions(_dimensions), type(_type), size(_size), diff --git a/src/DeviceInterface.cpp b/src/DeviceInterface.cpp index 643836843ba0..ac25f1f43515 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, bool texture) { +Expr make_device_interface_call(DeviceAPI device_api, MemoryType memory_type) { if (device_api == DeviceAPI::Host) { return make_zero(type_of()); } @@ -173,7 +173,7 @@ Expr make_device_interface_call(DeviceAPI device_api, bool texture) { interface_name = "halide_cuda_device_interface"; break; case DeviceAPI::OpenCL: - if (texture) { + if (memory_type == MemoryType::GPUTexture) { interface_name = "halide_opencl_image_device_interface"; } else { interface_name = "halide_opencl_device_interface"; diff --git a/src/DeviceInterface.h b/src/DeviceInterface.h index fb1c12028f6d..32b8230736e5 100644 --- a/src/DeviceInterface.h +++ b/src/DeviceInterface.h @@ -6,6 +6,7 @@ */ #include "Target.h" +#include "Expr.h" namespace Halide { @@ -37,7 +38,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, bool texture = false); +Expr make_device_interface_call(DeviceAPI device_api, MemoryType memory_type = MemoryType::Auto); } // namespace Internal } // namespace Halide diff --git a/src/IR.cpp b/src/IR.cpp index bf44dc45b740..4505a6fe3fe7 100644 --- a/src/IR.cpp +++ b/src/IR.cpp @@ -608,9 +608,7 @@ 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 69c232968f2e..40f5d6410c88 100644 --- a/src/IR.h +++ b/src/IR.h @@ -520,9 +520,7 @@ 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/InjectHostDevBufferCopies.cpp b/src/InjectHostDevBufferCopies.cpp index b745859fd9c3..c85cc55c55c8 100644 --- a/src/InjectHostDevBufferCopies.cpp +++ b/src/InjectHostDevBufferCopies.cpp @@ -53,24 +53,23 @@ class FindBufferUsage : public IRVisitor { void visit(const Call *op) override { if (op->is_intrinsic(Call::image_load) || - op->is_intrinsic(Call::image_load_texture)) { + op->is_intrinsic(Call::image_load)) { 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); + touched_as_texture = touched_as_texture || op->is_intrinsic(Call::image_load); } 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) || - op->is_intrinsic(Call::image_store_texture)) { + } else if (op->is_intrinsic(Call::image_store)) { 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); + touched_as_texture = touched_as_texture || op->is_intrinsic(Call::image_store); } for (size_t i = 0; i < op->args.size(); i++) { if (i == 1) continue; @@ -152,7 +151,7 @@ class InjectBufferCopiesForSingleBuffer : public IRMutator { bool is_external; - bool is_texture; + MemoryType memory_type; enum FlagState { Unknown, @@ -188,7 +187,7 @@ class InjectBufferCopiesForSingleBuffer : public IRMutator { } Stmt make_device_malloc(DeviceAPI target_device_api) { - Expr device_interface = make_device_interface_call(target_device_api, is_texture); + Expr device_interface = make_device_interface_call(target_device_api, memory_type); Stmt device_malloc = call_extern_and_assert("halide_device_malloc", {buffer_var(), device_interface}); return device_malloc; @@ -199,7 +198,7 @@ class InjectBufferCopiesForSingleBuffer : public IRMutator { } Stmt make_copy_to_device(DeviceAPI target_device_api) { - Expr device_interface = make_device_interface_call(target_device_api, is_texture); + Expr device_interface = make_device_interface_call(target_device_api, memory_type); return call_extern_and_assert("halide_copy_to_device", {buffer_var(), device_interface}); } @@ -410,8 +409,8 @@ class InjectBufferCopiesForSingleBuffer : public IRMutator { } public: - InjectBufferCopiesForSingleBuffer(const std::string &b, bool e, bool t) - : buffer(b), is_external(e), is_texture(t) { + InjectBufferCopiesForSingleBuffer(const std::string &b, bool e, MemoryType m) + : buffer(b), is_external(e), memory_type(m) { if (is_external) { // The state of the buffer is totally unknown, which is // the default constructor for this->state @@ -615,7 +614,7 @@ class InjectBufferCopies : public IRMutator { Stmt body = mutate(op->body); - InjectBufferCopiesForSingleBuffer injector(op->name, false, op->memory_type == MemoryType::GPUTexture); + InjectBufferCopiesForSingleBuffer injector(op->name, false, op->memory_type); body = injector.mutate(body); string buffer_name = op->name + ".buffer"; @@ -643,7 +642,7 @@ class InjectBufferCopies : public IRMutator { internal_assert(free_injecter.success); } - Expr device_interface = make_device_interface_call(touching_device, op->memory_type == MemoryType::GPUTexture); + Expr device_interface = make_device_interface_call(touching_device, op->memory_type); return InjectCombinedAllocation(op->name, op->type, op->extents, op->condition, device_interface) @@ -733,10 +732,7 @@ 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()); - } + result_storage[p.name()] = p.memory_type(); } } @@ -764,7 +760,7 @@ class InjectBufferCopiesForInputsAndOutputs : public IRMutator { public: set result; - set result_textures; + std::map result_storage; }; public: @@ -776,7 +772,7 @@ class InjectBufferCopiesForInputsAndOutputs : public IRMutator { s.accept(&finder); Stmt new_stmt = s; for (const string &buf : finder.result) { - new_stmt = InjectBufferCopiesForSingleBuffer(buf, true, finder.result_textures.count(buf)).mutate(new_stmt); + new_stmt = InjectBufferCopiesForSingleBuffer(buf, true, finder.result_storage.at(buf)).mutate(new_stmt); } return new_stmt; } else { diff --git a/src/StorageFlattening.cpp b/src/StorageFlattening.cpp index a13388a3747c..e7cc413bfd69 100644 --- a/src/StorageFlattening.cpp +++ b/src/StorageFlattening.cpp @@ -241,7 +241,6 @@ 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); } @@ -261,7 +260,6 @@ class FlattenDimensions : public IRMutator { 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); @@ -272,7 +270,7 @@ class FlattenDimensions : public IRMutator { args.push_back(op->args[i] - min); } args.push_back(value); - Expr store = Call::make(value.type(), Call::image_store_texture, + Expr store = Call::make(value.type(), Call::image_store, args, Call::Intrinsic); return Evaluate::make(store); } else { @@ -319,7 +317,7 @@ class FlattenDimensions : public IRMutator { } return Call::make(op->type, - textures.count(op->name) ? Call::image_load_texture : Call::image_load, + Call::image_load, args, Call::PureIntrinsic, FunctionPtr(), diff --git a/src/runtime/opencl.cpp b/src/runtime/opencl.cpp index 221178ef21a9..e946b39a67b0 100644 --- a/src/runtime/opencl.cpp +++ b/src/runtime/opencl.cpp @@ -688,9 +688,6 @@ WEAK int halide_opencl_initialize_kernels(void *user_context, void **state_ptr, debug(user_context) << (void *)program << "\n"; } - // halide_print(user_context, "Source: \n"); - // halide_print(user_context, src); - (*state)->program = program; debug(user_context) << " clBuildProgram " << (void *)program << " " << options.str() << "\n"; @@ -1585,8 +1582,6 @@ 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 = 0; desc.image_slice_pitch = 0; desc.num_mip_levels = 0; @@ -1700,8 +1695,6 @@ WEAK int halide_opencl_image_buffer_copy(void *user_context, struct halide_buffe dim >= 2 ? static_cast(dst->dim[1].extent) : 1, 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; if (dst->dimensions >= 2 && dst->dim[1].stride != dst->dim[0].extent) { error(user_context) << "image buffer copies must be dense on inner dimension"; return halide_error_code_device_buffer_copy_failed; @@ -1721,8 +1714,7 @@ 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[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; + if (src->dimensions >= 2 && src->dim[1].stride != src->dim[0].extent) { error(user_context) << "image buffer copies must be dense on inner dimension"; return halide_error_code_device_buffer_copy_failed; @@ -1737,9 +1729,6 @@ WEAK int halide_opencl_image_buffer_copy(void *user_context, struct halide_buffe } else if (!from_host && !to_host) { error(user_context) << "image to image copies not implemented"; return halide_error_code_device_buffer_copy_failed; - // 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) { From b8348e03fd71c73b7d19ce4b0323af854d60e2ba Mon Sep 17 00:00:00 2001 From: John Laxson Date: Thu, 1 Oct 2020 19:27:20 -0600 Subject: [PATCH 19/29] format --- src/DeviceInterface.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/DeviceInterface.h b/src/DeviceInterface.h index 32b8230736e5..121cead51e20 100644 --- a/src/DeviceInterface.h +++ b/src/DeviceInterface.h @@ -5,8 +5,8 @@ * Methods for managing device allocations when jitting */ -#include "Target.h" #include "Expr.h" +#include "Target.h" namespace Halide { From 27362670a97106d2d43df4a1e28242a0d8f9e58f Mon Sep 17 00:00:00 2001 From: Steven Johnson Date: Fri, 2 Oct 2020 11:08:08 -0700 Subject: [PATCH 20/29] Fix signed/unsigned comparison issue --- src/runtime/opencl.cpp | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/src/runtime/opencl.cpp b/src/runtime/opencl.cpp index e946b39a67b0..c0e81340cd3a 100644 --- a/src/runtime/opencl.cpp +++ b/src/runtime/opencl.cpp @@ -1529,7 +1529,8 @@ WEAK int halide_opencl_image_device_malloc(void *user_context, halide_buffer_t * cl_image_desc desc; struct halide_type_t type = buf->type; - format.image_channel_data_type = -1; + const cl_channel_type CL_INVALID = 0xffff; + format.image_channel_data_type = CL_INVALID; if (type.code == halide_type_int) { if (type.bits == 8) { format.image_channel_data_type = CL_SIGNED_INT8; @@ -1553,7 +1554,7 @@ WEAK int halide_opencl_image_device_malloc(void *user_context, halide_buffer_t * format.image_channel_data_type = CL_FLOAT; } } - if (format.image_channel_data_type == -1) { + if (format.image_channel_data_type == CL_INVALID) { error(user_context) << "Unhandled datatype for opencl texture object: " << type; return halide_error_code_device_malloc_failed; } From a4137b63df7a6434aa868e58f4dcfc0def185a7e Mon Sep 17 00:00:00 2001 From: John Laxson Date: Mon, 5 Oct 2020 20:29:47 -0500 Subject: [PATCH 21/29] Set storage for Buffer<> nodes --- src/InjectHostDevBufferCopies.cpp | 1 + 1 file changed, 1 insertion(+) diff --git a/src/InjectHostDevBufferCopies.cpp b/src/InjectHostDevBufferCopies.cpp index 9e55116faf4e..6e221b6f9377 100644 --- a/src/InjectHostDevBufferCopies.cpp +++ b/src/InjectHostDevBufferCopies.cpp @@ -739,6 +739,7 @@ class InjectBufferCopiesForInputsAndOutputs : public IRMutator { void include(const Buffer<> &b) { if (b.defined()) { result.insert(b.name()); + result_storage[b.name()] = MemoryType::Auto; } } From dec6954ccf26795fb59485ba6239dec9313038fa Mon Sep 17 00:00:00 2001 From: John Laxson Date: Mon, 5 Oct 2020 20:41:02 -0500 Subject: [PATCH 22/29] redundancies have been made redundant --- src/InjectHostDevBufferCopies.cpp | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/src/InjectHostDevBufferCopies.cpp b/src/InjectHostDevBufferCopies.cpp index 6e221b6f9377..db1afd228232 100644 --- a/src/InjectHostDevBufferCopies.cpp +++ b/src/InjectHostDevBufferCopies.cpp @@ -52,8 +52,7 @@ class FindBufferUsage : public IRVisitor { } void visit(const Call *op) override { - if (op->is_intrinsic(Call::image_load) || - op->is_intrinsic(Call::image_load)) { + if (op->is_intrinsic(Call::image_load)) { internal_assert(!op->args.empty()); if (is_buffer_var(op->args[1])) { devices_touched.insert(current_device_api); From b905086bdf7bc1e710d4e28acd604e40c5966dd4 Mon Sep 17 00:00:00 2001 From: John Laxson Date: Thu, 8 Oct 2020 15:09:49 -0500 Subject: [PATCH 23/29] touched_as_texture is dead --- src/InjectHostDevBufferCopies.cpp | 5 ----- 1 file changed, 5 deletions(-) diff --git a/src/InjectHostDevBufferCopies.cpp b/src/InjectHostDevBufferCopies.cpp index b4ac06dcc058..3ddb5afeb87b 100644 --- a/src/InjectHostDevBufferCopies.cpp +++ b/src/InjectHostDevBufferCopies.cpp @@ -56,7 +56,6 @@ class FindBufferUsage : public IRVisitor { 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); } for (size_t i = 0; i < op->args.size(); i++) { if (i == 1) { @@ -69,8 +68,6 @@ class FindBufferUsage : public IRVisitor { 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); } for (size_t i = 0; i < op->args.size(); i++) { if (i == 1) { @@ -133,8 +130,6 @@ 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) { } From 64d1506056fc9d49c99732e9d3bd61fa107a2403 Mon Sep 17 00:00:00 2001 From: John Laxson Date: Thu, 8 Oct 2020 15:09:59 -0500 Subject: [PATCH 24/29] Guard texture with OpenCL --- test/generator/gpu_texture_generator.cpp | 6 ++++-- 1 file changed, 4 insertions(+), 2 deletions(-) diff --git a/test/generator/gpu_texture_generator.cpp b/test/generator/gpu_texture_generator.cpp index 52c537d095c4..4e9300535c3d 100644 --- a/test/generator/gpu_texture_generator.cpp +++ b/test/generator/gpu_texture_generator.cpp @@ -14,8 +14,10 @@ class GpuTexture : public Halide::Generator { // 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); + if(get_target().has_feature(Target::OpenCL)) { + input.store_in(MemoryType::GPUTexture); + output.store_in(MemoryType::GPUTexture); + } Target target = get_target(); if (target.has_gpu_feature()) { From 502e6854cc0d5c3c37dbc1f072fdfbb750c4a5ed Mon Sep 17 00:00:00 2001 From: John Laxson Date: Thu, 8 Oct 2020 15:18:30 -0500 Subject: [PATCH 25/29] lint --- test/generator/gpu_texture_generator.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/test/generator/gpu_texture_generator.cpp b/test/generator/gpu_texture_generator.cpp index 4e9300535c3d..6b4d4658dcb4 100644 --- a/test/generator/gpu_texture_generator.cpp +++ b/test/generator/gpu_texture_generator.cpp @@ -14,7 +14,7 @@ class GpuTexture : public Halide::Generator { // Create a simple pipeline that scales pixel values by 2. output(x, y) = input(x, y) * 2; - if(get_target().has_feature(Target::OpenCL)) { + if (get_target().has_feature(Target::OpenCL)) { input.store_in(MemoryType::GPUTexture); output.store_in(MemoryType::GPUTexture); } From 537d54c44c95621abd4728a57291032a7b5b12a4 Mon Sep 17 00:00:00 2001 From: John Laxson Date: Mon, 12 Oct 2020 23:06:52 -0600 Subject: [PATCH 26/29] Check OpenCL version on tests --- src/runtime/cl_functions.h | 17 ++++---- src/runtime/opencl.cpp | 60 +++++++++++++++++++++++--- test/correctness/gpu_texture.cpp | 10 +++++ test/generator/gpu_texture_aottest.cpp | 9 ++++ 4 files changed, 83 insertions(+), 13 deletions(-) diff --git a/src/runtime/cl_functions.h b/src/runtime/cl_functions.h index 9258e99c35ed..0e170890569c 100644 --- a/src/runtime/cl_functions.h +++ b/src/runtime/cl_functions.h @@ -9,6 +9,10 @@ #define CL_FN(ret, fn, args) #endif +#ifndef CL_12_FN +#define CL_12_FN(ret, fn, args) CL_FN(ret, fn, args) +#endif + /* Platform API */ CL_FN(cl_int, clGetPlatformIDs, (cl_uint /* num_entries */, @@ -37,20 +41,18 @@ CL_FN(cl_int, void * /* param_value */, size_t * /* param_value_size_ret */)); -#ifdef HAVE_OPENCL_12 -CL_FN(cl_int, +CL_12_FN(cl_int, clCreateSubDevices, (cl_device_id /* in_device */, const cl_device_partition_property * /* properties */, cl_uint /* num_devices */, cl_device_id * /* out_devices */, cl_uint * /* num_devices_ret */)); -CL_FN(cl_int, +CL_12_FN(cl_int, clRetainDevice, (cl_device_id /* device */)); -CL_FN(cl_int, +CL_12_FN(cl_int, clReleaseDevice, (cl_device_id /* device */)); -#endif /* Context APIs */ CL_FN(cl_context, @@ -116,15 +118,13 @@ CL_FN(cl_mem, const void * /* buffer_create_info */, cl_int * /* errcode_ret */)); -#ifdef HAVE_OPENCL_12 -CL_FN(cl_mem, +CL_12_FN(cl_mem, clCreateImage, (cl_context /* context */, cl_mem_flags /* flags */, const cl_image_format * /* image_format */, const cl_image_desc * /* image_desc */, void * /* host_ptr */, cl_int * /* errcode_ret */)); -#endif CL_FN(cl_int, clRetainMemObject, (cl_mem /* memobj */)); @@ -352,5 +352,6 @@ CL_FN(cl_int, cl_event * /* event */)); #undef CL_FN +#undef CL_12_FN // clang-format on diff --git a/src/runtime/opencl.cpp b/src/runtime/opencl.cpp index c0e81340cd3a..3fa6f648d93e 100644 --- a/src/runtime/opencl.cpp +++ b/src/runtime/opencl.cpp @@ -48,9 +48,9 @@ extern "C" WEAK void *halide_opencl_get_symbol(void *user_context, const char *n } template -ALWAYS_INLINE T get_cl_symbol(void *user_context, const char *name) { +ALWAYS_INLINE T get_cl_symbol(void *user_context, const char *name, bool req) { T s = (T)halide_opencl_get_symbol(user_context, name); - if (!s) { + if (!s && req) { error(user_context) << "OpenCL API not found: " << name << "\n"; } return s; @@ -61,7 +61,8 @@ WEAK void load_libopencl(void *user_context) { debug(user_context) << " load_libopencl (user_context: " << user_context << ")\n"; halide_assert(user_context, clCreateContext == NULL); -#define CL_FN(ret, fn, args) fn = get_cl_symbol(user_context, #fn); +#define CL_FN(ret, fn, args) fn = get_cl_symbol(user_context, #fn, true); +#define CL_12_FN(ret, fn, args) fn = get_cl_symbol(user_context, #fn, false); #include "cl_functions.h" } @@ -607,6 +608,55 @@ WEAK int halide_opencl_device_free(void *user_context, halide_buffer_t *buf) { return 0; } +WEAK int halide_opencl_compute_capability(void *user_context, int *major, int *minor) { + if (!lib_opencl) { + // If OpenCL can't be found, we want to return 0, 0 and it's not + // considered an error. So we should be very careful about + // looking for OpenCL without tripping any errors in the rest + // of this runtime. + void *sym = halide_opencl_get_symbol(user_context, "clCreateContext"); + if (!sym) { + *major = *minor = 0; + return 0; + } + } + + { + ClContext ctx(user_context); + if (ctx.error_code != 0) { + return ctx.error_code; + } + + cl_int err; + + cl_device_id devices[1]; + err = clGetContextInfo(ctx.context, CL_CONTEXT_DEVICES, sizeof(devices), devices, NULL); + if (err != CL_SUCCESS) { + error(user_context) << "CL: clGetContextInfo failed: " + << get_opencl_error_name(err); + return err; + } + + char device_version[256] = ""; + err = clGetDeviceInfo(devices[0], CL_DEVICE_VERSION, sizeof(device_version), device_version, NULL); + if (err != CL_SUCCESS) { + error(user_context) << "CL: clGetDeviceInfo failed: " + << get_opencl_error_name(err); + return err; + } + + // This should always be of the format "OpenCL X.Y" per the spec + if (strlen(device_version) < 10) { + return -1; + } + + *major = device_version[7] - '0'; + *minor = device_version[9] - '0'; + } + + return 0; +} + WEAK int halide_opencl_initialize_kernels(void *user_context, void **state_ptr, const char *src, int size) { debug(user_context) << "CL: halide_opencl_init_kernels (user_context: " << user_context @@ -1483,7 +1533,7 @@ WEAK halide_device_interface_t opencl_device_interface = { halide_device_release_crop, halide_device_wrap_native, halide_device_detach_native, - NULL, + halide_opencl_compute_capability, &opencl_device_interface_impl}; } // namespace OpenCL @@ -1875,7 +1925,7 @@ WEAK halide_device_interface_t opencl_image_device_interface = { halide_device_release_crop, halide_device_wrap_native, halide_device_detach_native, - NULL, + halide_opencl_compute_capability, &opencl_image_device_interface_impl}; } // namespace OpenCL diff --git a/test/correctness/gpu_texture.cpp b/test/correctness/gpu_texture.cpp index 613f6f78cf45..1f48073c68b7 100644 --- a/test/correctness/gpu_texture.cpp +++ b/test/correctness/gpu_texture.cpp @@ -1,4 +1,5 @@ #include "Halide.h" +#include "HalideRuntimeOpenCL.h" #include using namespace Halide; @@ -12,6 +13,15 @@ int main(int argc, char **argv) { return 0; } + const auto *interface = get_device_interface_for_device_api(DeviceAPI::OpenCL); + assert(interface->compute_capability != nullptr); + int major, minor; + int err = interface->compute_capability(nullptr, &major, &minor); + if (major == 1 && minor < 2) { + printf("[SKIP] OpenCL %d.%d is less than required 1.2.\n", major, minor); + return 0; + } + // Check dynamic allocations into Heap and Texture memory for (auto memory_type : {MemoryType::GPUTexture, MemoryType::Heap}) { { diff --git a/test/generator/gpu_texture_aottest.cpp b/test/generator/gpu_texture_aottest.cpp index 8ec8b306ac16..494ed854f5fd 100644 --- a/test/generator/gpu_texture_aottest.cpp +++ b/test/generator/gpu_texture_aottest.cpp @@ -20,6 +20,15 @@ using namespace Halide::Runtime; int main(int argc, char **argv) { #if defined(TEST_OPENCL) + const auto *interface = halide_opencl_device_interface(); + assert(interface->compute_capability != nullptr); + int major, minor; + int err = interface->compute_capability(nullptr, &major, &minor); + if (major == 1 && minor < 2) { + printf("[SKIP] OpenCl %d.%d is less than required 1.2.\n", major, minor); + return 0; + } + const int W = 32, H = 32; Buffer input(W, H); for (int y = 0; y < input.height(); y++) { From 55d1a88c3ca548c1438b23233a82f3450dda6da9 Mon Sep 17 00:00:00 2001 From: Steven Johnson Date: Wed, 14 Oct 2020 08:50:24 -0700 Subject: [PATCH 27/29] Avoid unused-variable error --- test/correctness/gpu_texture.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/test/correctness/gpu_texture.cpp b/test/correctness/gpu_texture.cpp index 1f48073c68b7..62ae5feb77a2 100644 --- a/test/correctness/gpu_texture.cpp +++ b/test/correctness/gpu_texture.cpp @@ -17,7 +17,7 @@ int main(int argc, char **argv) { assert(interface->compute_capability != nullptr); int major, minor; int err = interface->compute_capability(nullptr, &major, &minor); - if (major == 1 && minor < 2) { + if (err != 0 || (major == 1 && minor < 2)) { printf("[SKIP] OpenCL %d.%d is less than required 1.2.\n", major, minor); return 0; } From f80455fb61596b510ebf69952219fdeed91a424f Mon Sep 17 00:00:00 2001 From: Steven Johnson Date: Wed, 14 Oct 2020 12:55:36 -0700 Subject: [PATCH 28/29] Avoid another unused-variable error --- test/generator/gpu_texture_aottest.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/test/generator/gpu_texture_aottest.cpp b/test/generator/gpu_texture_aottest.cpp index 494ed854f5fd..9a958d309a3f 100644 --- a/test/generator/gpu_texture_aottest.cpp +++ b/test/generator/gpu_texture_aottest.cpp @@ -24,7 +24,7 @@ int main(int argc, char **argv) { assert(interface->compute_capability != nullptr); int major, minor; int err = interface->compute_capability(nullptr, &major, &minor); - if (major == 1 && minor < 2) { + if (err != 0 || (major == 1 && minor < 2)) { printf("[SKIP] OpenCl %d.%d is less than required 1.2.\n", major, minor); return 0; } From 5b958b74aa7ae1909f9e95ce3a877cb65744113f Mon Sep 17 00:00:00 2001 From: Steven Johnson Date: Fri, 16 Oct 2020 11:39:11 -0700 Subject: [PATCH 29/29] Skip generator_aotcpp_gpu_texture in tests --- Makefile | 3 +++ 1 file changed, 3 insertions(+) diff --git a/Makefile b/Makefile index 92bbfd2bbed3..f8564139625d 100644 --- a/Makefile +++ b/Makefile @@ -1100,6 +1100,9 @@ GENERATOR_AOTCPP_TESTS := $(filter-out generator_aotcpp_multitarget,$(GENERATOR_ # remove AOT-CPP tests that don't (yet) work for C++ backend # (each tagged with the *known* blocking issue(s)) +# https://github.com/halide/Halide/issues/2084 (only if opencl enabled) +GENERATOR_AOTCPP_TESTS := $(filter-out generator_aotcpp_gpu_texture,$(GENERATOR_AOTCPP_TESTS)) + # https://github.com/halide/Halide/issues/2084 (only if opencl enabled) GENERATOR_AOTCPP_TESTS := $(filter-out generator_aotcpp_acquire_release,$(GENERATOR_AOTCPP_TESTS))