diff --git a/.clang-format b/.clang-format index cceda7f..c2129ec 100644 --- a/.clang-format +++ b/.clang-format @@ -1,2 +1,13 @@ -BasedOnStyle: Microsoft -IndentWidth: 4 \ No newline at end of file +# BasedOnStyle: Microsoft +# IndentWidth: 4 + +# http://clang.llvm.org/docs/ClangFormatStyleOptions.html +BasedOnStyle: Chromium + +ColumnLimit: 100 + +# Use 4 space indents +IndentWidth: 4 +ObjCBlockIndentWidth: 4 +AccessModifierOffset: -2 +InsertBraces: true \ No newline at end of file diff --git a/.gitignore b/.gitignore index 7ad2d18..ea62771 100644 --- a/.gitignore +++ b/.gitignore @@ -26,3 +26,4 @@ __pycache__ # Build folder build/ +.cache/ \ No newline at end of file diff --git a/.gitmodules b/.gitmodules new file mode 100644 index 0000000..1e50d5b --- /dev/null +++ b/.gitmodules @@ -0,0 +1,3 @@ +[submodule "dawn"] + path = dawn + url = https://dawn.googlesource.com/dawn diff --git a/.vscode/settings.json b/.vscode/settings.json index 8152b4d..64c11c0 100644 --- a/.vscode/settings.json +++ b/.vscode/settings.json @@ -1,5 +1,8 @@ { "editor.semanticTokenColorCustomizations": { "enabled": true, // enable for all themes - } + }, + "clang-format.executable": "/usr/bin/clang-format", + "editor.inlayHints.enabled": "off", + "git.ignoreLimitWarning": true } \ No newline at end of file diff --git a/CMakeLists.txt b/CMakeLists.txt index 7cd86a5..6f92131 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -12,7 +12,20 @@ add_definitions(${LLVM_DEFINITIONS}) list(APPEND CMAKE_MODULE_PATH "${LLVM_CMAKE_DIR}") include(AddLLVM) -set(CMAKE_CXX_STANDARD 17) +set(CMAKE_CXX_STANDARD 20) +set(CMAKE_EXPORT_COMPILE_COMMANDS ON) + +# Disable Tint tests and command line tools BEFORE add_subdirectory + +set(DAWN_FETCH_DEPENDENCIES ON) + +set(DAWN_BUILD_SAMPLES OFF CACHE BOOL "" FORCE) +set(DAWN_BUILD_TESTS OFF CACHE BOOL "" FORCE) +set(DAWN_BUILD_NODE_BINDINGS OFF CACHE BOOL "" FORCE) + +set(TINT_BUILD_BENCHMARKS OFF CACHE BOOL "" FORCE) +set(TINT_BUILD_TESTS OFF CACHE BOOL "" FORCE) +set(TINT_BUILD_AS_OTHER_OS OFF CACHE BOOL "" FORCE) add_subdirectory(dawn) @@ -21,4 +34,4 @@ add_subdirectory(tools) if( LLVM_INCLUDE_TESTS ) add_subdirectory(unittests) -endif() +endif() \ No newline at end of file diff --git a/README.md b/README.md index 8c328f2..c0e41db 100644 --- a/README.md +++ b/README.md @@ -4,7 +4,11 @@ A very naive WGSL backend for LLVM intended to be used to port CUDA programs to ## Building -1. Tested with LLVM 19.0.0 +1. If you are cloning this repo for the first time, then clone this repo recursively to get `dawn`'s repo as well + +```bash +git clone https://github.com/grx6741/llvm-wgsl.git --recursive +``` 2. Install LLVM 19.0.0 from source or use a package manager. @@ -15,6 +19,9 @@ sudo apt install llvm-19-dev 3. Compile with `cmake`. ```bash +# Note, may not work with other compilers like gcc +export CC=clang +export CXX=clang++ cmake -B build -G Ninja # only tested with Ninja ``` diff --git a/lib/Target/WGSLBackend/CMakeLists.txt b/lib/Target/WGSLBackend/CMakeLists.txt index c420f22..c5926b7 100644 --- a/lib/Target/WGSLBackend/CMakeLists.txt +++ b/lib/Target/WGSLBackend/CMakeLists.txt @@ -5,7 +5,6 @@ add_subdirectory(TargetInfo) add_llvm_target(WGSLBackendCodeGen WGSLBackend.cpp WGSLTargetMachine.cpp - WGSL.cpp LINK_COMPONENTS Analysis @@ -26,4 +25,7 @@ add_llvm_target(WGSLBackendCodeGen WGSLBackend ) - target_link_libraries(LLVMWGSLBackendCodeGen PRIVATE tint_api) +target_link_libraries(LLVMWGSLBackendCodeGen + PRIVATE + tint_api +) diff --git a/lib/Target/WGSLBackend/WGSL.cpp b/lib/Target/WGSLBackend/WGSL.cpp deleted file mode 100644 index 91c2c4c..0000000 --- a/lib/Target/WGSLBackend/WGSL.cpp +++ /dev/null @@ -1,49 +0,0 @@ -#include "WGSL.h" -#include - -#include "llvm/Support/ErrorHandling.h" - -#define NOT_SUPPORTED(x) llvm::report_fatal_error(x " is not supported") - - -WGSLType::WGSLType(std::string name, PrimitiveType Maintype, uint32_t numElements, PrimitiveType UnderlyingType) - : m_Name{name}, m_Type{Maintype}, m_NumElements{numElements}, m_UnderlyingType{UnderlyingType} -{ - if (m_Type == PrimitiveType::Tarray) { - if (!isValidArrayElementType(m_UnderlyingType)) - llvm::report_fatal_error("Invalid array element type"); - } -} - -void WGSLType::print(llvm::raw_string_ostream &os, bool type) const -{ - os << m_Name << " "; - assert(!isAbstractType(m_Type) && "Trying to print abstract type"); - - if (!type) - return; - - switch (m_Type) - { - case PrimitiveType::Ti32: - case PrimitiveType::Tf32: - case PrimitiveType::Tu32: - case PrimitiveType::Tbool: - os << WGSL::PrimitiveTypeNames.at(m_Type); - break; - case PrimitiveType::Tarray: - os << ": array<" << WGSL::PrimitiveTypeNames.at(m_UnderlyingType) << ", " << m_NumElements << ">"; - break; - case PrimitiveType::Tpointer: - NOT_SUPPORTED("pointer type"); - break; - - case PrimitiveType::Tvoid: - case PrimitiveType::Ti64: - case PrimitiveType::Tf64: - break; - case PrimitiveType::Tnone: - llvm::report_fatal_error("Attempted to print an abstract type (Tnone)"); - break; - } -} diff --git a/lib/Target/WGSLBackend/WGSL.h b/lib/Target/WGSLBackend/WGSL.h deleted file mode 100644 index 41c99d8..0000000 --- a/lib/Target/WGSLBackend/WGSL.h +++ /dev/null @@ -1,59 +0,0 @@ -#ifndef WGSL -#define WGSL - -#include -#include - -namespace WGSL -{ - enum class PrimitiveType - { - Ti32, - Ti64, // abstract-int - Tf32, - Tf64, // abstract-float - Tu32, - Tbool, - Tvoid, - Tarray, - Tpointer, - Tnone - }; - - inline std::unordered_map PrimitiveTypeNames = { - {PrimitiveType::Ti32, "i32"}, - {PrimitiveType::Ti64, "i64"}, - {PrimitiveType::Tf32, "f32"}, - {PrimitiveType::Tf64, "f64"}, - {PrimitiveType::Tu32, "u32"}, - {PrimitiveType::Tbool, "bool"}, - {PrimitiveType::Tarray, "array"}, - {PrimitiveType::Tpointer, "pointer"}, - {PrimitiveType::Tnone, "none"}, - }; -} - -inline bool isAbstractType(const PrimitiveType type) { return type == PrimitiveType::Ti64 || type == PrimitiveType::Tf64; } -inline bool isValidArrayElementType(const PrimitiveType type) { return type == PrimitiveType::Ti32 || type == PrimitiveType::Tu32 || type == PrimitiveType::Tf32 || type == PrimitiveType::Tbool; } - -void printPrimitiveType(const PrimitiveType type, llvm::raw_string_ostream& os); - -class WGSLType -{ - private: - PrimitiveType m_Type{PrimitiveType::Tnone}; - std::string m_Name{""}; - - uint32_t m_NumElements{0}; // if array, then length - - // If this is an array, then this is the type of the elements - // If this i a pointer, then this is the type of the pointed to object - PrimitiveType m_UnderlyingType{PrimitiveType::Tnone}; - - public: - WGSLType(std::string name, PrimitiveType Maintype, uint32_t numElements = 0, PrimitiveType UnderlyingType = PrimitiveType::Tnone); - - void print(llvm::raw_string_ostream &os, bool type = true) const; -}; - -#endif // WGSL diff --git a/lib/Target/WGSLBackend/WGSLBackend.cpp b/lib/Target/WGSLBackend/WGSLBackend.cpp index 25e628c..c396ca5 100644 --- a/lib/Target/WGSLBackend/WGSLBackend.cpp +++ b/lib/Target/WGSLBackend/WGSLBackend.cpp @@ -1,14 +1,13 @@ #include "WGSLBackend.h" #include "WGSLTargetMachine.h" -#include "WGSL.h" +#include "llvm/Demangle/Demangle.h" #include "llvm/MC/TargetRegistry.h" +#include "llvm/Support/raw_ostream.h" -namespace llvm_wgsl -{ +namespace llvm_wgsl { -extern "C" void LLVMInitializeWGSLBackendTarget() -{ +extern "C" void LLVMInitializeWGSLBackendTarget() { // Register the target. RegisterTargetMachine X(TheWGSLBackendTarget); } @@ -16,149 +15,112 @@ extern "C" void LLVMInitializeWGSLBackendTarget() // public char WGSLWriter::ID = 0; -bool WGSLWriter::doInitialization(Module &M) -{ +bool WGSLWriter::doInitialization(llvm::Module& M) { + // TODO : do GPU arch check and kernel function check in + // runOnFunction method per function instead checkGPUArchitecture(M); getAllKernelFunctions(M); return false; } -bool WGSLWriter::doFinalization(Module &M) -{ +bool WGSLWriter::doFinalization(llvm::Module& M) { return true; } -bool WGSLWriter::runOnFunction(Function &F) -{ - if (!kernelFunctions.count(&F)) - return false; +bool WGSLWriter::runOnFunction(llvm::Function& F) { + /* - StringRef kernelName = F.getName(); - Type *returnType = F.getReturnType(); - - WGSL::PrimitiveType returnPrimitiveType = WGSL::PrimitiveType::Tvoid; - if (returnType->isIntegerTy(32)) - returnPrimitiveType = WGSL::PrimitiveType::Ti32; - else if (returnType->isFloatTy()) - returnPrimitiveType = WGSL::PrimitiveType::Tf32; - else if (returnType->isVoidTy()) - returnPrimitiveType = WGSL::PrimitiveType::Tvoid; - else if (returnType->isPointerTy()) { - returnPrimitiveType = WGSL::PrimitiveType::Tpointer; - } - else { - llvm::report_fatal_error("Unsupported return type in WGSL"); - errs() << "Unsupported return type in WGSL : " << returnType->getTypeID() << "\n"; + if (!m_KernelFunctions.count(&F)) { + return false; } - std::string out; - llvm::raw_string_ostream os(out); - - os << "fn " << kernelName << "("; - - for (auto argIter = F.arg_begin(); argIter != F.arg_end(); ++argIter) - { - const Argument &arg = *argIter; - Type *argType = arg.getType(); - - WGSL::PrimitiveType paramType = WGSL::PrimitiveType::Tnone; - if (argType->isIntegerTy(32)) { - paramType = WGSL::PrimitiveType::Ti32; - } else if (argType->isFloatTy()) { - paramType = WGSL::PrimitiveType::Tf32; - } else if (argType->isIntegerTy(1)) { - paramType = WGSL::PrimitiveType::Tbool; - } else if (argType->isPointerTy()) { - paramType = WGSL::PrimitiveType::Tpointer; - } else { - llvm::report_fatal_error("Unsupported argument type in WGSL"); - } - - if (argIter != F.arg_begin()) { - os << ", "; - } - - os << arg.getName() << ": " << WGSL::PrimitiveTypeNames.at(paramType); - } + std::string functionDefinitionstr; + llvm::raw_string_ostream functionDefinitionStream(functionDefinitionstr); + + llvm::StringRef kernelName = F.getName(); + std::string kernelNameStr = kernelName.str(); + + char* demangledName = itaniumDemangle(kernelName); - os << ")"; + llvm::outs() << "Demangled kernel name: " << (demangledName ? demangledName : "null") << "\n"; - if (returnPrimitiveType != WGSL::PrimitiveType::Tvoid) - os << " -> " << WGSL::PrimitiveTypeNames.at(returnPrimitiveType); + m_Demangler.partialDemangle(kernelNameStr.c_str()); - os << " {}"; + char *demangled_name = m_Demangler.getFunctionName(nullptr, nullptr); + assert(demangled_name && "Demangled name should not be null"); - llvm::outs() << os.str() << "\n"; + functionDefinitionStream << "fn " << demangled_name << "("; + + free(demangled_name); + + char *parameters = m_Demangler.getFunctionParameters(nullptr, nullptr); + assert(parameters && "Demangled parameters should not be null"); + free(parameters); + functionDefinitionstr.clear(); + + */ return false; } // private -void WGSLWriter::checkGPUArchitecture(const Module &M) -{ +void WGSLWriter::checkGPUArchitecture(const llvm::Module& M) { // Only supports sm_50 for now - for (auto &F : M) - { - if (F.isDeclaration()) - continue; // Skip declarations + for (auto& F : M) { + if (F.isDeclaration()) { + continue; // Skip declarations + } llvm::AttributeList attrs = F.getAttributes(); - if (attrs.hasFnAttr("target-cpu")) - { + if (attrs.hasFnAttr("target-cpu")) { llvm::StringRef cpu = attrs.getFnAttr("target-cpu").getValueAsString(); - if (cpu != "sm_50") - { + if (cpu != "sm_50") { llvm::outs() << "We Only support sm_50 for now, but found: " << cpu << "\n"; report_fatal_error("Unsupported GPU architecture"); } - return; + return; } } } -void WGSLWriter::getAllKernelFunctions(const llvm::Module &M) -{ - llvm::NamedMDNode *annotations = M.getNamedMetadata("nvvm.annotations"); - if (!annotations) - { +void WGSLWriter::getAllKernelFunctions(const llvm::Module& M) { + llvm::NamedMDNode* annotations = M.getNamedMetadata("nvvm.annotations"); + if (!annotations) { llvm::outs() << "No WGSL kernel functions found in the module.\n"; report_fatal_error("No WGSL kernel functions found"); } - for (const auto *MD : annotations->operands()) - { - if (MD->getNumOperands() < 2) + for (const auto* MD : annotations->operands()) { + if (MD->getNumOperands() < 2) { continue; + } - auto *fnVal = llvm::dyn_cast(MD->getOperand(0).get()); - auto *strVal = llvm::dyn_cast(MD->getOperand(1).get()); + auto* fnVal = llvm::dyn_cast(MD->getOperand(0).get()); + auto* strVal = llvm::dyn_cast(MD->getOperand(1).get()); - if (!fnVal || !strVal) + if (!fnVal || !strVal) { continue; + } - if (strVal->getString() == "kernel") - { - if (auto *F = llvm::dyn_cast(fnVal->getValue())) - { - kernelFunctions.insert(F); + if (strVal->getString() == "kernel") { + if (auto* F = llvm::dyn_cast(fnVal->getValue())) { + m_KernelFunctions.insert(F); llvm::outs() << "Found WGSL kernel function: " << F->getName() << "\n"; } } } - if (kernelFunctions.empty()) - { + if (m_KernelFunctions.empty()) { llvm::outs() << "No WGSL kernel functions found in the module.\n"; llvm_unreachable("No WGSL kernel functions found"); } - if (kernelFunctions.size() > 1) - { + if (m_KernelFunctions.size() > 1) { llvm::outs() << "Multiple WGSL kernel functions found, only one is supported.\n"; llvm_unreachable("Multiple WGSL kernel functions found"); } } -} // namespace llvm_wgsl +} // namespace llvm_wgsl \ No newline at end of file diff --git a/lib/Target/WGSLBackend/WGSLBackend.h b/lib/Target/WGSLBackend/WGSLBackend.h index 3319090..a6272d2 100644 --- a/lib/Target/WGSLBackend/WGSLBackend.h +++ b/lib/Target/WGSLBackend/WGSLBackend.h @@ -1,13 +1,11 @@ -#include #include -#include #include "IDMap.h" #include "llvm/ADT/StringExtras.h" #include "llvm/Analysis/LoopInfo.h" -#include "llvm/CodeGen/IntrinsicLowering.h" #include "llvm/CodeGen/Passes.h" +#include "llvm/Demangle/Demangle.h" #include "llvm/IR/Attributes.h" #include "llvm/IR/DataLayout.h" #include "llvm/IR/InstVisitor.h" @@ -22,60 +20,161 @@ #include "llvm/Transforms/Scalar.h" #include "src/tint/api/tint.h" - -namespace llvm_wgsl -{ - -using namespace llvm; - -class WGSLMCAsmInfo : public MCAsmInfo -{ +#include "src/tint/lang/core/address_space.h" +#include "src/tint/lang/core/ir/builder.h" +#include "src/tint/lang/core/ir/convert.h" +#include "src/tint/lang/core/ir/module.h" +#include "src/tint/lang/core/ir/validator.h" +#include "src/tint/lang/core/number.h" +#include "src/tint/lang/core/type/array_count.h" +#include "src/tint/lang/core/type/i32.h" +#include "src/tint/lang/core/type/void.h" +#include "src/tint/lang/wgsl/writer/ir_to_program/program_options.h" +#include "src/tint/lang/wgsl/writer/options.h" +#include "src/tint/lang/wgsl/writer/writer.h" +/* +#include "src/tint/utils/result/result.h" + */ + +namespace llvm_wgsl { + +class WGSLMCAsmInfo : public MCAsmInfo { public: - WGSLMCAsmInfo() - { - PrivateGlobalPrefix = ""; - } + WGSLMCAsmInfo() { PrivateGlobalPrefix = ""; } }; -using FunctionInfoVariant = std::variant; - -// WGSLWriter - This class is the main chunk of code that converts an LLVM module to a WGSL compute shader. -class WGSLWriter : public FunctionPass, public InstVisitor -{ - private: - std::unordered_set kernelFunctions; - raw_ostream &FileOut; - const Instruction *CurInstr = nullptr; - +// WGSLWriter - This class is the main chunk of code that converts an LLVM module to a WGSL compute +// shader. +class WGSLWriter : public llvm::FunctionPass, public llvm::InstVisitor { public: static char ID; - explicit WGSLWriter(raw_ostream &o) : FunctionPass(ID), FileOut(o) - { - tint::Initialize(); - - tint::Shutdown(); + inline explicit WGSLWriter(llvm::raw_ostream& o) : llvm::FunctionPass(ID), m_FileOut(o) { + using namespace tint; + using namespace tint::core; + using namespace tint::core::ir; + using namespace tint::core::type; + /* + */ + + tint::Initialize(); + + // type::Manager types; + ir::Module mod; + ir::Builder builder{mod}; + + auto* types = &mod.Types(); + + auto* i32_ty = types->Get(); + auto* u32_ty = types->Get(); + auto* f32_ty = types->Get(); + + auto* runtime_count = types->Get(); + + auto* a_type = types->Get( + core::AddressSpace::kStorage, + types->Get(u32_ty, runtime_count, 0u, 0u, 0u, 0u), + core::Access::kRead); + auto* a = builder.Var("a", a_type); + a->SetBindingPoint(0u, 0u); + mod.root_block->Append(a); + + auto* b_type = types->Get( + core::AddressSpace::kStorage, + types->Get(f32_ty, runtime_count, 0u, 0u, 0u, 0u), + core::Access::kRead); + auto* b = builder.Var("b", b_type); + b->SetBindingPoint(0u, 1u); + mod.root_block->Append(b); + + auto* c_type = types->Get( + core::AddressSpace::kStorage, + types->Get(i32_ty, runtime_count, 0u, 0u, 0u, 0u), + core::Access::kReadWrite); + auto* c = builder.Var("c", c_type); + c->SetBindingPoint(0u, 2u); + mod.root_block->Append(c); + + auto* entry_ret_type = types->Get(); + auto* entry = builder.ComputeFunction("addKernel"); + + auto* id_param = builder.FunctionParam("id", types->Get(u32_ty, 3u)); + + auto entry_param_id_attr = IOAttributes(); + entry_param_id_attr.builtin = core::BuiltinValue::kGlobalInvocationId; + id_param->SetAttributes(entry_param_id_attr); + + entry->AppendParam(id_param); + + mod.functions.Clear(); + mod.functions.Push(entry); + + builder.Append(entry->Block(), [&] { + // let i = id.x; + auto* id = entry->Params()[0]; + auto* var_6 = builder.Access(u32_ty, id, builder.Constant(u32(0))); + auto* var_i = builder.Let(var_6); + mod.SetName(var_i, "i"); + + // c[i] + auto* var_8 = builder.Access( + types->ptr(), c, var_i); + + // i32(a[i]) + auto* var_9 = builder.Access( + types->ptr(), a, var_i); + auto* var_10 = builder.Load(var_9); + auto* var_11 = builder.Convert(i32_ty, var_10); + + // i32(b[i]) + auto* var_12 = builder.Access( + types->ptr(), b, var_i); + auto* var_13 = builder.Load(var_12); + auto* var_14 = builder.Convert(i32_ty, var_13); + + // c[i] = i32(a[i]) + i32(b[i]) + auto* var_15 = builder.Add(i32_ty, var_11, var_14); + builder.Store(var_8, var_15); + + builder.Return(entry); + }); + + if (auto res = Validate(mod); res != Success) { + llvm::errs() << res.Failure().reason << "\n"; + return; + } + + auto wgsl_out = wgsl::writer::WgslFromIR(mod, wgsl::writer::ProgramOptions{}); + if (wgsl_out != Success) { + llvm::errs() << "WGSL generation failed: " << wgsl_out.Failure().reason << "\n"; + return; + } + + llvm::outs() << wgsl_out->wgsl << "\n"; + + tint::Shutdown(); } - virtual inline StringRef getPassName() const - { - return "WGSL backend"; - } + virtual inline llvm::StringRef getPassName() const { return "WGSL backend"; } - inline void getAnalysisUsage(AnalysisUsage &AU) const - { - AU.addRequired(); + inline void getAnalysisUsage(llvm::AnalysisUsage& AU) const { + AU.addRequired(); AU.setPreservesCFG(); } - virtual bool doInitialization(Module &M); - virtual bool doFinalization(Module &M); - virtual bool runOnFunction(Function &F); + virtual bool doInitialization(llvm::Module& M); + virtual bool doFinalization(llvm::Module& M); + virtual bool runOnFunction(llvm::Function& F); private: + std::unordered_set m_KernelFunctions; + llvm::raw_ostream& m_FileOut; + const llvm::Instruction* m_CurInstr = nullptr; + llvm::ItaniumPartialDemangler m_Demangler; + // On Initialization - void checkGPUArchitecture(const Module &M); - void getAllKernelFunctions(const Module &M); + void checkGPUArchitecture(const llvm::Module& M); + void getAllKernelFunctions(const llvm::Module& M); }; -} // namespace llvm_wgsl +} // namespace llvm_wgsl diff --git a/test/cu_ll_tests/vector_add_device_expected.wgsl b/test/cu_ll_tests/vector_add_device_expected.wgsl new file mode 100644 index 0000000..459f8de --- /dev/null +++ b/test/cu_ll_tests/vector_add_device_expected.wgsl @@ -0,0 +1,11 @@ +@group(0) @binding(0) var a: array; +@group(0) @binding(1) var b: array; +@group(0) @binding(2) var c: array; + +@compute @workgroup_size(1) +fn addKernel( + @builtin(global_invocation_id) id: vec3 +) { + let i = id.x; + c[i] = i32(a[i]) + i32(b[i]); +} \ No newline at end of file diff --git a/test/cu_ll_tests/vector_add_device_tint_ir.wgsl b/test/cu_ll_tests/vector_add_device_tint_ir.wgsl new file mode 100644 index 0000000..cb845cc --- /dev/null +++ b/test/cu_ll_tests/vector_add_device_tint_ir.wgsl @@ -0,0 +1,23 @@ +$B1: { # root + %a:ptr, read> = var @binding_point(0, 0) + %b:ptr, read> = var @binding_point(0, 1) + %c:ptr, read_write> = var @binding_point(0, 2) +} + +%addKernel = @compute @workgroup_size(1, 1, 1) +func(%id:vec3 [@global_invocation_id]):void { + $B2: { + %6:u32 = access %id, 0u + %i:u32 = let %6 + %8:ptr = access %c, %i + %9:ptr = access %a, %i + %10:u32 = load %9 + %11:i32 = convert %10 + %12:ptr = access %b, %i + %13:f32 = load %12 + %14:i32 = convert %13 + %15:i32 = add %11, %14 + store %8, %15 + ret + } +} \ No newline at end of file diff --git a/tools/llvm-wgsl/CMakeLists.txt b/tools/llvm-wgsl/CMakeLists.txt index ef95628..e5da5b4 100644 --- a/tools/llvm-wgsl/CMakeLists.txt +++ b/tools/llvm-wgsl/CMakeLists.txt @@ -4,4 +4,6 @@ add_executable(llvm-wgsl llvm-wgsl.cpp) llvm_map_components_to_libnames(llvm_libs WGSLBackendCodeGen WGSLBackendInfo) -target_link_libraries(llvm-wgsl LLVM ${llvm_libs}) \ No newline at end of file +target_link_libraries(llvm-wgsl LLVM + ${llvm_libs} + ) \ No newline at end of file