The graphics API is moving towards next generation.
Vulkan/Metal on the native and WebGPU on the web.
Due to the limited programming model, we cannot get the best compute performance in WebGL.
Now that the mainline already have both WebGPU and vulkan support, this PR phases out WebGL.
tvm_option(USE_CUDA "Build with CUDA" OFF)
tvm_option(USE_OPENCL "Build with OpenCL" OFF)
tvm_option(USE_VULKAN "Build with Vulkan" OFF)
-tvm_option(USE_OPENGL "Build with OpenGL" OFF)
tvm_option(USE_METAL "Build with Metal" OFF)
tvm_option(USE_ROCM "Build with ROCM" OFF)
tvm_option(ROCM_PATH "The path to rocm" /opt/rocm)
include(cmake/modules/CUDA.cmake)
include(cmake/modules/Hexagon.cmake)
include(cmake/modules/OpenCL.cmake)
-include(cmake/modules/OpenGL.cmake)
include(cmake/modules/OpenMP.cmake)
include(cmake/modules/Vulkan.cmake)
include(cmake/modules/Metal.cmake)
return "vpi";
case kDLROCM:
return "rocm";
- case kOpenGL:
- return "opengl";
case kDLExtDev:
return "ext_dev";
case kDLWebGPU:
*/
TVM_DLL Stage& double_buffer(); // NOLINT(*)
/*!
- * \brief Schedule for OpenGL fragment shader.
- * \return reference to self.
- */
- Stage& opengl(); // NOLINT(*)
- /*!
* \brief whether the stage has been scheduled.
* \return whether the stage has been scheduled.
*/
std::string scope;
/*! \brief Whether this is an output stage */
bool is_output{false};
- /*! \brief Whether this is an OpenGL stage */
- bool is_opengl{false};
/*! \brief Whether apply double buffer optimization to this stage */
bool double_buffer{false};
/*!
v->Visit("attach_stage", &attach_stage);
v->Visit("scope", &scope);
v->Visit("is_output", &is_output);
- v->Visit("is_opengl", &is_opengl);
v->Visit("double_buffer", &double_buffer);
v->Visit("group", &group);
v->Visit("num_child_stages", &num_child_stages);
static constexpr const char* shift_right = "shift_right";
static constexpr const char* popcount = "popcount";
static constexpr const char* likely = "likely";
- static constexpr const char* glsl_texture_store = "glsl_texture_store";
static constexpr const char* prefetch = "prefetch";
static constexpr const char* isnan = "isnan";
static constexpr const char* isfinite = "isfinite";
constexpr const char* pipeline_stage_scope = "pipeline_stage_scope";
/*! \brief pipeline execution scope, implies the scope can be pipelined. */
constexpr const char* pipeline_exec_scope = "pipeline_exec_scope";
-/*!
- * \brief Mark that this stage is an OpenGL shader. Since OpenGL shader only
- * allows writing out to one element of the output texture, the Provide node
- * gets translated to a special Call::glsl_texture_store statement instead of a
- * Store statement.
- */
-constexpr const char* opengl_stage_scope = "opengl_stage_scope";
/*!
* \brief Mark that it is in the device scope.
f_name = "device_api.gpu";
} else if (target == "cl" || target == "opencl" || target == "sdaccel") {
f_name = "device_api.opencl";
- } else if (target == "gl" || target == "opengl") {
- f_name = "device_api.opengl";
} else if (target == "mtl" || target == "metal") {
f_name = "device_api.metal";
} else if (target == "vulkan") {
+++ /dev/null
-/*
- * Licensed to the Apache Software Foundation (ASF) under one
- * or more contributor license agreements. See the NOTICE file
- * distributed with this work for additional information
- * regarding copyright ownership. The ASF licenses this file
- * to you under the Apache License, Version 2.0 (the
- * "License"); you may not use this file except in compliance
- * with the License. You may obtain a copy of the License at
- *
- * http://www.apache.org/licenses/LICENSE-2.0
- *
- * Unless required by applicable law or agreed to in writing,
- * software distributed under the License is distributed on an
- * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
- * KIND, either express or implied. See the License for the
- * specific language governing permissions and limitations
- * under the License.
- */
-
-/*!
- * Optional module when build opencl is switched to off
- */
-#include "../../runtime/opengl/opengl_module.h"
-#include "../source/codegen_source_base.h"
-
-namespace tvm {
-namespace runtime {
-
-Module OpenGLModuleCreate(std::unordered_map<std::string, OpenGLShader> shaders, std::string fmt,
- std::unordered_map<std::string, FunctionInfo> fmap) {
- LOG(WARNING) << "OpenGL runtime not enabled, return a source module...";
- auto data = ToJSON(shaders);
- return codegen::DeviceSourceModuleCreate(data, "gl", fmap, "opengl");
-}
-
-} // namespace runtime
-} // namespace tvm
+++ /dev/null
-/*
- * Licensed to the Apache Software Foundation (ASF) under one
- * or more contributor license agreements. See the NOTICE file
- * distributed with this work for additional information
- * regarding copyright ownership. The ASF licenses this file
- * to you under the Apache License, Version 2.0 (the
- * "License"); you may not use this file except in compliance
- * with the License. You may obtain a copy of the License at
- *
- * http://www.apache.org/licenses/LICENSE-2.0
- *
- * Unless required by applicable law or agreed to in writing,
- * software distributed under the License is distributed on an
- * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
- * KIND, either express or implied. See the License for the
- * specific language governing permissions and limitations
- * under the License.
- */
-
-/*!
- * \file codegen_opengl.cc
- *
- * We are targeting OpenGL 3.3. The reason of not targeting a recent version
- * of OpenGL is to have better compatibility of WebGL 2.
- */
-#include "codegen_opengl.h"
-
-#include <string>
-#include <unordered_map>
-#include <utility>
-#include <vector>
-
-#include "../../runtime/thread_storage_scope.h"
-#include "../build_common.h"
-
-namespace tvm {
-namespace codegen {
-
-CodeGenOpenGL::CodeGenOpenGL() : output_(nullptr), output_iter_var_(nullptr) {}
-
-void CodeGenOpenGL::InitFuncState(const PrimFunc& f) {
- CodeGenC::InitFuncState(f);
- output_ = nullptr;
- inputs_.clear();
- output_iter_var_ = nullptr;
- thread_extent_var_ = "";
- this->decl_stream.str("");
- this->stream.str("");
-}
-
-void CodeGenOpenGL::AddFunction(const PrimFunc& f) {
- // clear previous generated state.
- this->InitFuncState(f);
-
- this->decl_stream << "#version 300 es\n";
- this->decl_stream << "precision highp float;\n";
-
- // skip the first underscore, so SSA variable starts from _1
- GetUniqueName("_");
-
- // Allocate argument names. Store in `var_idmap_`.
- for (auto arg : f->params) {
- auto arg_name = GetUniqueName(arg.get()->name_hint);
- var_idmap_[arg.get()] = arg_name;
-
- if (auto* ptr = arg->type_annotation.as<PointerTypeNode>()) {
- if (auto* prim = ptr->element_type.as<PrimTypeNode>()) {
- RegisterHandleType(arg.get(), prim->dtype);
- }
- }
- }
-
- thread_extent_var_ = GetUniqueName("thread_extent");
- this->decl_stream << "uniform int " << thread_extent_var_ << ";\n";
-
- this->stream << "void main() {\n";
-
- int func_scope = this->BeginScope();
- this->PrintStmt(f->body);
- this->EndScope(func_scope);
-
- this->PrintIndent();
- this->stream << "}\n\n";
-
- // Declare arguments.
- for (auto arg : f->params) {
- if (this->inputs_.find(arg.get()) != this->inputs_.cend()) {
- // Declare input texture.
- // Format:
- // - Float: "uniform sampler2D {name};"
- // - Int: "uniform isampler2D {name};"
- // - UInt: "uniform usampler2D {name};"
-
- auto arg_name = GetVarID(arg.get());
-
- auto type_it = this->handle_data_type_.find(arg.get());
- CHECK(type_it != this->handle_data_type_.cend()) << "Cannot find type.";
- DLDataType type = type_it->second;
- CHECK_EQ(type.lanes, 1) << "Vector type not supported.";
-
- switch (type.code) {
- case kDLInt:
- this->decl_stream << "uniform isampler2D " << arg_name << ";\n";
- break;
- case kDLUInt:
- this->decl_stream << "uniform usampler2D " << arg_name << ";\n";
- break;
- case kDLFloat:
- this->decl_stream << "uniform sampler2D " << arg_name << ";\n";
- break;
- default:
- LOG(FATAL) << "Unsupported type code.";
- }
-
- } else if (this->output_ == arg.get()) {
- // Declare output texture.
- // Format: "out {type} {name};"
-
- auto arg_name = GetVarID(arg.get());
-
- auto type_it = this->handle_data_type_.find(arg.get());
- CHECK(type_it != this->handle_data_type_.cend()) << "Cannot find type.";
- auto type = type_it->second;
-
- this->decl_stream << "out ";
- PrintType(type, this->decl_stream);
- this->decl_stream << " " << arg_name << ";\n";
-
- } else {
- // Declare uniform value.
- // Format: "uniform {type} {name};"
-
- auto arg_name = GetVarID(arg.get());
- auto type = arg.get()->dtype;
-
- this->decl_stream << "uniform ";
- PrintType(type, this->decl_stream);
- this->decl_stream << " " << arg_name << ";\n";
- }
- }
-
- std::vector<std::string> arg_names;
- std::vector<runtime::OpenGLArgKind> arg_kinds;
- for (auto arg : f->params) {
- std::string name = GetVarID(arg.get());
-
- runtime::OpenGLArgKind kind;
- if (inputs_.find(arg.get()) != inputs_.cend()) {
- kind = runtime::OpenGLArgKind::kInputTexture;
- } else if (output_ == arg.get()) {
- kind = runtime::OpenGLArgKind::kOutputTexture;
- } else {
- kind = runtime::OpenGLArgKind::kUniform;
- }
-
- arg_names.push_back(name);
- arg_kinds.push_back(kind);
- }
-
- auto global_symbol = f->GetAttr<String>(tvm::attr::kGlobalSymbol);
- CHECK(global_symbol.defined())
- << "CodeGenOpenGL: Expect PrimFunc to have the global_symbol attribute";
-
- shaders_[static_cast<std::string>(global_symbol.value())] =
- runtime::OpenGLShader(this->decl_stream.str() + this->stream.str(), std::move(arg_names),
- std::move(arg_kinds), this->thread_extent_var_);
-}
-
-std::unordered_map<std::string, runtime::OpenGLShader> CodeGenOpenGL::Finish() { return shaders_; }
-
-void CodeGenOpenGL::BindThreadIndex(const IterVar& iv) {
- CHECK_EQ(iv->thread_tag, "threadIdx.x") << "Must be threadIdx.x";
- CHECK(var_idmap_.find(iv->var.get()) == var_idmap_.end()) << "Only support one thread iter var";
- CHECK(output_iter_var_ == nullptr) << "Only support one thread iter var";
-
- var_idmap_[iv->var.get()] = iv->thread_tag;
- output_iter_var_ = iv->var.get();
-
- // Declare threadIdx local variable.
- this->PrintIndent();
- this->stream << "ivec2 threadIdx = ivec2(" << runtime::kTextureRowSize
- << " * int(gl_FragCoord.y) + int(gl_FragCoord.x), 0);\n";
-
- // Return directly if threadIdx.x >= thread_extent.
- this->PrintIndent();
- this->stream << "if (threadIdx.x >= " << thread_extent_var_ << ") {\n";
- this->PrintIndent();
- this->stream << " return;\n";
- this->PrintIndent();
- this->stream << "}\n";
-}
-
-void CodeGenOpenGL::VisitStmt_(const StoreNode* op) {
- LOG(FATAL) << "Store statement not supported in OpenGL."
- << " Texture store should be a Call statement.";
-}
-
-// texelFetch(tex, ivec2(idx & kTextureRowMask, idx >> kTextureRowBits), 0).r
-std::string CodeGenOpenGL::TexelFetch(const VarNode* buffer, PrimExpr index) {
- std::ostringstream os;
- os << "texelFetch(" << GetVarID(buffer) << ", ivec2(int(";
- PrintExpr(index, os);
- os << ") & " << runtime::kTextureRowMask << ", int(";
- PrintExpr(index, os);
- os << ") >> " << runtime::kTextureRowBits << "), 0).r";
- return os.str();
-}
-
-// Print a reference expression to a buffer.
-// Format: texelFetch(buffer, index, 0).r
-std::string CodeGenOpenGL::GetBufferRef(DataType t, const VarNode* buffer, PrimExpr index) {
- CHECK_EQ(t.lanes(), 1) << "Vector type not supported.";
- CHECK(HandleTypeMatch(buffer, t)) << "Type mismatch not supported.";
-
- if (buffer == this->output_) {
- // This is the output texture.
- return GetVarID(buffer);
- } else {
- // This is an input texture.
- this->inputs_.insert(buffer);
- return TexelFetch(buffer, index);
- }
-}
-
-void CodeGenOpenGL::PrintType(DataType t, std::ostream& os) {
- switch (t.code()) {
- case kDLInt:
- CHECK_EQ(t.bits(), 32) << "Only support 32-bit int.";
- os << "int";
- break;
- case kDLUInt:
- CHECK_EQ(t.bits(), 32) << "Only support 32-bit uint.";
- os << "uint";
- break;
- case kDLFloat:
- CHECK_EQ(t.bits(), 32) << "Only support 32-bit float.";
- os << "float";
- break;
- default:
- LOG(FATAL) << "Unsupported type code.";
- }
-}
-
-// Codegen for immediate values
-
-void CodeGenOpenGL::VisitExpr_(const IntImmNode* op, std::ostream& os) {
- CHECK_EQ(op->dtype, DataType::Int(32)) << "GLSL 3.0 only supports 32-bit ints.";
- CodeGenC::VisitExpr_(op, os);
-}
-
-void CodeGenOpenGL::VisitExpr_(const FloatImmNode* op, std::ostream& os) {
- CHECK_EQ(op->dtype, DataType::Float(32)) << "GLSL 3.0 only supports 32-bit floats.";
- CodeGenC::VisitExpr_(op, os);
-}
-
-void CodeGenOpenGL::VisitExpr_(const StringImmNode*, std::ostream& os) {
- LOG(FATAL) << "GLSL 3.0 doesn't support strings.";
-}
-
-void CodeGenOpenGL::VisitStmt_(const EvaluateNode* op) {
- auto call = op->value.as<CallNode>();
- if (call == nullptr || call->name != CallNode::glsl_texture_store) {
- // Fallback to normal logic.
- CodeGenC::VisitStmt_(op);
- }
-
- CHECK_EQ(call->args.size(), 2);
- auto buffer = call->args[0].as<VarNode>();
- auto value = call->args[1];
-
- // Doesn't support store to vector.
- auto type = value.dtype();
- CHECK_EQ(type.lanes(), 1) << "Vectorized store not implemented, type = " << type;
-
- CHECK(inputs_.find(buffer) == inputs_.cend())
- << "Texture has been read from before. Must not store to it.";
- if (output_ == nullptr) {
- output_ = buffer; // Record that this texture is the output.
- } else {
- CHECK(output_ == buffer) << "GLSL can only write to 1 texture.";
- }
-
- this->PrintIndent();
- this->stream << GetVarID(buffer) << " = " << PrintExpr(value) << ";\n";
-}
-
-runtime::Module BuildOpenGL(IRModule mod, std::string target) {
- bool output_ssa = false;
- CodeGenOpenGL cg;
- cg.Init(output_ssa);
-
- for (auto kv : mod->functions) {
- CHECK(kv.second->IsInstance<PrimFuncNode>()) << "CodeGenOpenGL: Can only take PrimFunc";
- auto f = Downcast<PrimFunc>(kv.second);
- auto calling_conv = f->GetAttr<Integer>(tvm::attr::kCallingConv);
- CHECK(calling_conv == CallingConv::kDeviceKernelLaunch)
- << "CodeGenOpenGL: expect calling_conv equals CallingConv::kDeviceKernelLaunch";
- cg.AddFunction(f);
- }
-
- auto shaders = cg.Finish();
- return OpenGLModuleCreate(shaders, "gl", ExtractFuncInfo(mod));
-}
-
-TVM_REGISTER_GLOBAL("target.build.opengl").set_body_typed(BuildOpenGL);
-
-} // namespace codegen
-} // namespace tvm
+++ /dev/null
-/*
- * Licensed to the Apache Software Foundation (ASF) under one
- * or more contributor license agreements. See the NOTICE file
- * distributed with this work for additional information
- * regarding copyright ownership. The ASF licenses this file
- * to you under the Apache License, Version 2.0 (the
- * "License"); you may not use this file except in compliance
- * with the License. You may obtain a copy of the License at
- *
- * http://www.apache.org/licenses/LICENSE-2.0
- *
- * Unless required by applicable law or agreed to in writing,
- * software distributed under the License is distributed on an
- * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
- * KIND, either express or implied. See the License for the
- * specific language governing permissions and limitations
- * under the License.
- */
-
-/*!
- * \file codegen_opengl.h
- * \brief Generate OpenGL device code.
- */
-#ifndef TVM_TARGET_SOURCE_CODEGEN_OPENGL_H_
-#define TVM_TARGET_SOURCE_CODEGEN_OPENGL_H_
-
-#include <tvm/target/codegen.h>
-
-#include <string>
-#include <unordered_map>
-#include <unordered_set>
-
-#include "../../runtime/opengl/opengl_module.h"
-#include "codegen_c.h"
-
-namespace tvm {
-namespace codegen {
-
-class CodeGenOpenGL final : public CodeGenC {
- public:
- CodeGenOpenGL();
- std::unordered_map<std::string, runtime::OpenGLShader> Finish();
-
- void AddFunction(const PrimFunc& f);
- void InitFuncState(const PrimFunc& f) final;
- void BindThreadIndex(const IterVar& iv) final;
- void VisitStmt_(const StoreNode* op) final;
- std::string TexelFetch(const VarNode* buffer, PrimExpr index);
- std::string GetBufferRef(DataType t, const VarNode* buffer, PrimExpr index) final;
- void PrintType(DataType t, std::ostream& os) final; // NOLINT(*)
-
- // Codegen for immediate values
- void VisitExpr_(const IntImmNode* op, std::ostream& os) final; // NOLINT(*)
- void VisitExpr_(const FloatImmNode* op, std::ostream& os) final; // NOLINT(*)
- void VisitExpr_(const StringImmNode* op, std::ostream& os) final; // NOLINT(*)
-
- // Match glsl_texture_store Call.
- void VisitStmt_(const EvaluateNode* op) final; // NOLINT(*)
-
- private:
- const VarNode* output_{nullptr};
- std::unordered_set<const VarNode*> inputs_;
- const VarNode* output_iter_var_{nullptr};
- std::unordered_map<std::string, runtime::OpenGLShader> shaders_;
- std::string thread_extent_var_;
-};
-
-} // namespace codegen
-} // namespace tvm
-
-#endif // TVM_TARGET_SOURCE_CODEGEN_OPENGL_H_
+++ /dev/null
-/*
- * Licensed to the Apache Software Foundation (ASF) under one
- * or more contributor license agreements. See the NOTICE file
- * distributed with this work for additional information
- * regarding copyright ownership. The ASF licenses this file
- * to you under the Apache License, Version 2.0 (the
- * "License"); you may not use this file except in compliance
- * with the License. You may obtain a copy of the License at
- *
- * http://www.apache.org/licenses/LICENSE-2.0
- *
- * Unless required by applicable law or agreed to in writing,
- * software distributed under the License is distributed on an
- * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
- * KIND, either express or implied. See the License for the
- * specific language governing permissions and limitations
- * under the License.
- */
-
-/*!
- * \file intrin_rule_opencl.cc
- * \brief OpenCL intrinsic rules.
- */
-#include "../intrin_rule.h"
-
-namespace tvm {
-namespace codegen {
-namespace intrin {
-
-TVM_REGISTER_GLOBAL("tvm.intrin.rule.opengl.floor").set_body(DispatchExtern<Direct>);
-
-TVM_REGISTER_GLOBAL("tvm.intrin.rule.opengl.ceil").set_body(DispatchExtern<Direct>);
-
-TVM_REGISTER_GLOBAL("tvm.intrin.rule.opengl.exp").set_body(DispatchExtern<Direct>);
-
-TVM_REGISTER_GLOBAL("tvm.intrin.rule.opengl.exp2").set_body(DispatchExtern<Direct>);
-
-TVM_REGISTER_GLOBAL("tvm.intrin.rule.opengl.exp10").set_body(DispatchExtern<Direct>);
-
-TVM_REGISTER_GLOBAL("tvm.intrin.rule.opengl.log").set_body(DispatchExtern<Direct>);
-
-TVM_REGISTER_GLOBAL("tvm.intrin.rule.opengl.log2").set_body(DispatchExtern<Direct>);
-
-TVM_REGISTER_GLOBAL("tvm.intrin.rule.opengl.log10").set_body(DispatchExtern<Direct>);
-
-TVM_REGISTER_GLOBAL("tvm.intrin.rule.opengl.tanh").set_body(DispatchExtern<Direct>);
-
-TVM_REGISTER_GLOBAL("tvm.intrin.rule.opengl.sqrt").set_body(DispatchExtern<Direct>);
-
-TVM_REGISTER_GLOBAL("tvm.intrin.rule.opengl.pow").set_body(DispatchExtern<Direct>);
-
-TVM_REGISTER_GLOBAL("tvm.intrin.rule.opengl.popcount").set_body(DispatchExtern<Direct>);
-
-TVM_REGISTER_GLOBAL("tvm.intrin.rule.opengl.sin").set_body(DispatchExtern<Direct>);
-
-TVM_REGISTER_GLOBAL("tvm.intrin.rule.opengl.sinh").set_body(DispatchExtern<Direct>);
-
-TVM_REGISTER_GLOBAL("tvm.intrin.rule.opengl.cos").set_body(DispatchExtern<Direct>);
-
-TVM_REGISTER_GLOBAL("tvm.intrin.rule.opengl.cosh").set_body(DispatchExtern<Direct>);
-
-} // namespace intrin
-} // namespace codegen
-} // namespace tvm
* \brief Construct a Target node from the given name and options.
* \param target_name The major target name. Should be one of
* {"aocl", "aocl_sw_emu", "c", "cuda", "ext_dev", "hexagon", "hybrid", "llvm",
- * "metal", "nvptx", "opencl", "opengl", "rocm", "sdaccel", "stackvm", "vulkan"}
+ * "metal", "nvptx", "opencl", "rocm", "sdaccel", "stackvm", "vulkan"}
* \param options Additional options appended to the target
* \return The constructed Target
*/
t->device_type = kDLAOCL;
t->keys_array.push_back("aocl");
t->keys_array.push_back("hls");
- } else if (target_name == "opengl") {
- t->device_type = kOpenGL;
- t->keys_array.push_back("opengl");
} else if (target_name == "stackvm") {
t->device_type = kDLCPU;
} else if (target_name == "ext_dev") {
return *this;
}
-Stage& Stage::opengl() {
- CHECK(!is_scheduled()) << "Must be a fresh schedule";
- StageNode* self = operator->();
-
- auto all_iter_vars = self->all_iter_vars; // curr version of all_iter_vars
- CHECK(!all_iter_vars.empty()) << "At least one iter var";
-
- // Fuse all data parallel dimensions to 1.
- IterVar fused = all_iter_vars[0];
- for (size_t i = 1; i != all_iter_vars.size(); ++i) {
- auto iter_var = all_iter_vars[i];
- switch (iter_var->iter_type) {
- case IterVarType::kDataPar: {
- fuse(fused, all_iter_vars[i], &fused);
- break;
- }
- case IterVarType::kThreadIndex: {
- LOG(ERROR) << "A fresh schedule shouldn't have thread index iter var";
- break;
- }
- case IterVarType::kCommReduce:
- case IterVarType::kOrdered:
- case IterVarType::kOpaque: {
- break;
- }
- default: {
- LOG(ERROR) << "Invalid iter var type " << IterVarType2String(iter_var->iter_type);
- break;
- }
- }
- }
-
- // Bind the only dimension to threadIdx.x.
- bind(fused, thread_axis(Range(nullptr), "threadIdx.x"));
-
- // Mark this stage as OpenGL.
- (*this)->is_opengl = true;
-
- return *this;
-}
-
Stage CopyStage(const Stage& s) {
ObjectPtr<StageNode> n = make_object<StageNode>(*s.operator->());
return Stage(n);
TVM_REGISTER_GLOBAL("te.StageDoubleBuffer").set_body_method(&Stage::double_buffer);
-TVM_REGISTER_GLOBAL("te.StageOpenGL").set_body_method(&Stage::opengl);
-
TVM_REGISTER_GLOBAL("te.ScheduleNormalize").set_body_method(&Schedule::normalize);
TVM_REGISTER_GLOBAL("te.ScheduleCreateGroup").set_body_method(&Schedule::create_group);
pipeline =
AttrStmtNode::make(s->op, tir::attr::realize_scope, StringImmNode::make(s->scope), pipeline);
- if (s->is_opengl) {
- pipeline =
- AttrStmtNode::make(s->op, tir::attr::opengl_stage_scope, StringImmNode::make(""), pipeline);
- }
return pipeline;
}
vinfo[dim].align_factor = tuple->args[1].as<IntImmNode>()->value;
vinfo[dim].align_offset = tuple->args[2].as<IntImmNode>()->value;
return this->VisitStmt(op->body);
- } else if (op->attr_key == attr::opengl_stage_scope) {
- is_opengl_ = true;
}
return StmtExprMutator::VisitStmt_(op);
}
const BufferEntry& e = it->second;
CHECK(!e.released) << "Read a buffer that is already out of scope";
- if (is_opengl_) {
- return EvaluateNode::make(CallNode::make(DataType(), CallNode::glsl_texture_store,
- {e.buffer->data, op->value}, CallNode::Intrinsic));
- } else {
- Stmt body = e.buffer.vstore(e.RelIndex(op->indices), op->value);
- if (create_bound_attributes_ && ShapeIsValid(e.buffer->shape)) {
- shape_collector_.push_back(std::make_pair(e.buffer->data, e.buffer->shape));
- }
- // To create bound attribute collector should has at least one item.
- if (create_bound_attributes_ && shape_collector_.size()) {
- for (size_t i = 0; i < shape_collector_.size(); ++i) {
- body = AttrStmtNode::make(shape_collector_[i].first, tir::attr::buffer_bound,
- MakeBound(e.buffer->dtype, shape_collector_[i].second), body);
- }
+
+ Stmt body = e.buffer.vstore(e.RelIndex(op->indices), op->value);
+ if (create_bound_attributes_ && ShapeIsValid(e.buffer->shape)) {
+ shape_collector_.push_back(std::make_pair(e.buffer->data, e.buffer->shape));
+ }
+ // To create bound attribute collector should has at least one item.
+ if (create_bound_attributes_ && shape_collector_.size()) {
+ for (size_t i = 0; i < shape_collector_.size(); ++i) {
+ body = AttrStmtNode::make(shape_collector_[i].first, tir::attr::buffer_bound,
+ MakeBound(e.buffer->dtype, shape_collector_[i].second), body);
}
- return body;
}
+ return body;
}
Stmt VisitStmt_(const BufferRealizeNode* op) final {
IRVisitorWithAnalyzer* bound_analyzer_;
// The size of cacheline
int cache_line_size_;
- // The current stage is an OpenGL shader.
- bool is_opengl_{false};
// Whether to mark load/store with theirs bounds.
bool create_bound_attributes_{false};
};
projects / platform / upstream / tvm.git / commitdiff