From f9b77bc717417c892c955be019aa271d715c2929 Mon Sep 17 00:00:00 2001 From: Kiniaxx Date: Thu, 16 Jul 2026 11:25:26 +0200 Subject: [PATCH] Added AMD GPU backend support. --- .gitignore | 4 + init_env.sh | 54 +++++ src/include/OSL/llvm_util.h | 8 + src/liboslexec/CMakeLists.txt | 36 ++- src/liboslexec/backendllvm.cpp | 299 ++++++++++++++++++++++- src/liboslexec/backendllvm.h | 3 +- src/liboslexec/llvm_instance.cpp | 20 +- src/liboslexec/llvm_util.cpp | 133 ++++++++++ src/liboslexec/oslexec_pvt.h | 55 +++++ src/liboslexec/shadingsys.cpp | 300 ++++++++++++++++++++++- src/oslc/CMakeLists.txt | 2 +- src/oslc/oslcmain.cpp | 34 ++- src/shaders/matte.oso | 18 ++ src/shaders/unnamed_group_1.ll | 33 +++ src/shaders/unnamed_group_1_O1.ll | 33 +++ src/testrender/CMakeLists.txt | 15 +- src/testrender/cuda/rend_lib.cu | 2 + src/testrender/gpu_raytracer.h | 22 ++ src/testrender/hip_raytracer.cpp | 308 ++++++++++++++++++++++++ src/testrender/hip_raytracer.h | 26 ++ src/testshade/CMakeLists.txt | 20 +- src/testshade/simplerend.cpp | 6 +- src/testshade/testshade.cpp | 73 +++++- testsuite/render-cornell/Untitled.ipynb | 6 + 24 files changed, 1490 insertions(+), 20 deletions(-) create mode 100644 init_env.sh create mode 100644 src/shaders/matte.oso create mode 100644 src/shaders/unnamed_group_1.ll create mode 100644 src/shaders/unnamed_group_1_O1.ll create mode 100644 src/testrender/gpu_raytracer.h create mode 100644 src/testrender/hip_raytracer.cpp create mode 100644 src/testrender/hip_raytracer.h create mode 100644 testsuite/render-cornell/Untitled.ipynb diff --git a/.gitignore b/.gitignore index c3c04174c4..3f2ed6e64e 100644 --- a/.gitignore +++ b/.gitignore @@ -11,3 +11,7 @@ build .vscode .envrc .DS_Store +build/ +.ipynb_checkpoints/ +ext/ +rocm.gpg diff --git a/init_env.sh b/init_env.sh new file mode 100644 index 0000000000..78faa1aca8 --- /dev/null +++ b/init_env.sh @@ -0,0 +1,54 @@ +#!/bin/bash + +echo ">>> 1. Odblokowanie systemu i czyszczenie błędów (Opcja Nuklearna)..." +# Usuwamy problematyczne skrypty postinst, które blokowały dpkg brakiem py3compile +sudo rm -f /var/lib/dpkg/info/python3-yaml.postinst +sudo rm -f /var/lib/dpkg/info/python3-pygments.postinst +sudo rm -f /var/lib/dpkg/info/libboost-mpi-python1.83.0.postinst + +# Wymuszamy naprawę systemu i usunięcie zepsutych pakietów +sudo dpkg --configure -a +sudo apt-get install -f -y + +echo ">>> 2. Instalacja bezpiecznych zależności (minimalny zestaw)..." +sudo apt-get update +sudo apt-get install -y \ + build-essential \ + cmake \ + git \ + wget \ + pkg-config \ + clang \ + llvm-dev \ + libboost-dev \ + libboost-filesystem-dev \ + libboost-thread-dev \ + libboost-system-dev \ + libopenexr-dev \ + libopenimageio-dev \ + libpugixml-dev \ + libtbb-dev \ + zlib1g-dev \ + python3-minimal + +echo ">>> 3. Konfiguracja AMD ROCm..." +if [ ! -f /etc/apt/keyrings/rocm.gpg ]; then + sudo mkdir -p /etc/apt/keyrings + wget -qO - https://repo.radeon.com/rocm/rocm.gpg.key | gpg --dearmor | sudo tee /etc/apt/keyrings/rocm.gpg > /dev/null +fi + +echo "deb [arch=amd64 signed-by=/etc/apt/keyrings/rocm.gpg] https://repo.radeon.com/rocm/apt/latest/ noble main" | sudo tee /etc/apt/sources.list.d/rocm.list + +sudo apt-get update +sudo apt-get install -y rocm-hip-sdk hipcc + +echo ">>> 4. Eksportowanie zmiennych środowiskowych..." +export ROCM_PATH=/opt/rocm +export HIP_PATH=/opt/rocm/hip +export PATH=$PATH:$ROCM_PATH/bin:$HIP_PATH/bin + +sudo apt-get update +sudo apt-get install -y pybind11-dev +sudo apt-get install -y bison flex +sudo apt-get install -y libclang-18-dev libclang-dev +echo "--- Środowisko gotowe do kompilacji ---" diff --git a/src/include/OSL/llvm_util.h b/src/include/OSL/llvm_util.h index c10f161ae3..ffacd60c06 100644 --- a/src/include/OSL/llvm_util.h +++ b/src/include/OSL/llvm_util.h @@ -1031,6 +1031,12 @@ class OSLEXECPUBLIC LLVM_Util { bool ptx_compile_group(llvm::Module* lib_module, const std::string& name, std::string& out); + + // NEW + llvm::TargetMachine* amdgpu_target_machine(const std::string& gpu_arch); + bool amdgpu_compile_group(llvm::Module* module, const std::string& name, std::string& out_code, const std::string& arch); + + /// Convert all functions in module's bitcode to a string. std::string bitcode_string(llvm::Module* module); @@ -1076,6 +1082,8 @@ class OSLEXECPUBLIC LLVM_Util { llvm::ExecutionEngine* m_llvm_exec; TargetISA m_target_isa = TargetISA::UNKNOWN; llvm::TargetMachine* m_nvptx_target_machine; + // NEW + llvm::TargetMachine* m_amdgpu_target_machine = nullptr; std::vector m_return_block; // stack for func call std::vector m_loop_after_block; // stack for break diff --git a/src/liboslexec/CMakeLists.txt b/src/liboslexec/CMakeLists.txt index a1e2d18fa3..680cce335c 100644 --- a/src/liboslexec/CMakeLists.txt +++ b/src/liboslexec/CMakeLists.txt @@ -149,9 +149,41 @@ file (GLOB exec_headers "*.h") file (GLOB compiler_headers "../liboslcomp/*.h") FLEX_BISON ( osolex.l osogram.y oso lib_src exec_headers ) +# NEW - KB +set(OSL_AMDGPU_TARGET_ARCH "" CACHE STRING "Target AMD GPU architecture (e.g. gfx900, gfx1030)") +if(OSL_AMDGPU_TARGET_ARCH) + message(STATUS "====> AMDGPU CROSS-COMPILATION: Forced target architecture -> ${OSL_AMDGPU_TARGET_ARCH}") + add_definitions(-DOSL_AMDGPU_TARGET_ARCH="${OSL_AMDGPU_TARGET_ARCH}") +endif() set ( CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -D__STDC_LIMIT_MACROS -D__STDC_CONSTANT_MACROS" ) - +# NEW - Ka +if (LLVM_CONFIG AND EXISTS ${LLVM_CONFIG}) + execute_process(COMMAND ${LLVM_CONFIG} --targets-built + OUTPUT_VARIABLE LLVM_TARGETS + OUTPUT_STRIP_TRAILING_WHITESPACE) + + if (LLVM_TARGETS MATCHES "AMDGPU") + message(STATUS "AMDGPU Backend: Supported target found in LLVM.") + add_definitions(-DOSL_ENABLE_AMDGPU=1) + + # W LLVM komponenty podaje się bez prefiksów. + # Nazwa 'amdgpu' jest nadrzędna i zazwyczaj wystarcza. + # Jeśli OSL nie widzi llvm_map_components_to_libnames, + # używamy bezpośredniego zapytania do llvm-config: + execute_process(COMMAND ${LLVM_CONFIG} --libfiles amdgpu codegen asmparser + OUTPUT_VARIABLE AMDGPU_LLVM_LIBS + OUTPUT_STRIP_TRAILING_WHITESPACE) + + # Konwersja tekstu z llvm-config na listę CMake + string(REPLACE " " ";" AMDGPU_LLVM_LIBS "${AMDGPU_LLVM_LIBS}") + string(REPLACE "\n" ";" AMDGPU_LLVM_LIBS "${AMDGPU_LLVM_LIBS}") + + message(STATUS "AMDGPU Backend: Linked libraries: ${AMDGPU_LLVM_LIBS}") + else() + message(STATUS "AMDGPU Backend: Not supported by your LLVM installation.") + endif() +endif() macro ( REQUIRE_INF_NAN SRC ) if (CMAKE_COMPILER_IS_INTELCLANG) @@ -581,6 +613,8 @@ target_link_libraries (${local_lib} pugixml::pugixml ${CMAKE_DL_LIBS} ${LLVM_LIBRARIES} ${LLVM_LDFLAGS} ${LLVM_SYSTEM_LIBRARIES} + ${AMDGPU_LLVM_LIBS} # NEW - Ka + ) target_compile_features (${local_lib} diff --git a/src/liboslexec/backendllvm.cpp b/src/liboslexec/backendllvm.cpp index 161ec852ab..86683a73ea 100644 --- a/src/liboslexec/backendllvm.cpp +++ b/src/liboslexec/backendllvm.cpp @@ -3,12 +3,41 @@ // https://github.com/AcademySoftwareFoundation/OpenShadingLanguage +// NEW - : Nagłówki LLVM potrzebne do zapisu bitkodu dla AMDGPU --- +#include +#include +#include +#include +#include +#include +#include +// Nagłówki Targetu i Emisji Kodu dla LLVM 18 +#include +#include +#include +#include +#include + +#include +#include +#include +#include + #include #include #include "oslexec_pvt.h" #include "backendllvm.h" + + +#include + +#include +#include +// Globalny cache testowy +std::map g_amdgpu_temp_cache; + using namespace OSL; using namespace OSL::pvt; @@ -924,7 +953,275 @@ BackendLLVM::find_userdata_index(const Symbol& sym) return userdata_index; } +std::vector BackendLLVM::get_llvm_bitcode(llvm::Module* custom_mod) { + shadingsys().info("========================================="); + shadingsys().info("[AMD] ROZPOCZYNAM GENEROWANIE KODU DLA GPU"); + shadingsys().info("========================================="); + + // 1. WYBÓR WŁAŚCIWEGO MODUŁU: + llvm::Module* mod = custom_mod ? custom_mod : ll.module(); + + // 2. INICJALIZACJA BACKENDU AMDGPU W LLVM + LLVMInitializeAMDGPUTargetInfo(); + LLVMInitializeAMDGPUTarget(); + LLVMInitializeAMDGPUTargetMC(); + LLVMInitializeAMDGPUAsmPrinter(); + + // 3. DETEKCJA ARCHITEKTURY SPRZĘTOWEJ RDNA + std::string arch_str = shadingsys().amdgpu_architecture().string(); + if (arch_str.empty()) { + arch_str = "gfx1100"; // Fallback dla RDNA3 + } + + // 4. WYSZUKANIE TARGETU ORAZ KONFIGURACJA TARGET MACHINE + std::string llvm_error; + const llvm::Target* target = llvm::TargetRegistry::lookupTarget("amdgcn-amd-amdhsa", llvm_error); + if (!target) { + shadingsys().error(OIIO::Strutil::format("LLVM Error: Nie znaleziono targetu AMDGPU: %s", llvm_error.c_str())); + return std::vector(); + } + + llvm::TargetOptions opt; + llvm::TargetMachine* target_machine = target->createTargetMachine( + "amdgcn-amd-amdhsa", arch_str, "", opt, llvm::Reloc::PIC_, + llvm::CodeModel::Small, llvm::CodeGenOptLevel::None); + + if (!target_machine) { + shadingsys().error("LLVM Error: Nie udalo sie utworzyc TargetMachine dla AMDGPU"); + return std::vector(); + } + + // Dostosowanie układu pamięci modułu pod architekturę AMD + mod->setDataLayout(target_machine->createDataLayout()); + mod->setTargetTriple("amdgcn-amd-amdhsa"); + + // Wymuszenie wygenerowania HSA Kernel Descriptors przez LLVM + mod->addModuleFlag(llvm::Module::Error, "amdgpu_code_object_version", 500); + +// ==================== DEBUG AMD ==================== +std::cout << "\n[AMD DEBUG] === ROZPOCZĘCIE ZRZUTU FUNKCJI W MODULE ===\n"; +std::cout << "[AMD DEBUG] Nazwa modułu: " << (mod->getModuleIdentifier()) << "\n"; +int func_counter = 0; +for (llvm::Function& F : *mod) { + func_counter++; + std::cout << " [" << func_counter << "] Nazwa: " << F.getName().str() + << " | Deklaracja: " << (F.isDeclaration() ? "TAK" : "NIE"); + if (!F.isDeclaration()) { + std::cout << " | Instrukcji: " << F.getInstructionCount(); + } + std::cout << "\n"; +} +std::cout << "[AMD DEBUG] === KONIEC ZRZUTU (Razem funkcji: " << func_counter << ") ===\n\n"; +// ============================================================== +// 5. LOKALIZACJA FUNKCJI OSL (Inicjalizacja + Warstwa) + shadingsys().info("[LLVM AMDGPU] Budowanie Megakernela GPU: szukanie init i layer..."); + + llvm::Function* osl_init_func = nullptr; + llvm::Function* osl_layer_func = nullptr; + + for (llvm::Function &F : *mod) { + if (!F.isDeclaration()) { + std::string fname = F.getName().str(); + if (fname.find("osl_init_group") != std::string::npos) { + osl_init_func = &F; + } else if (fname.find("osl_layer_group") != std::string::npos) { + osl_layer_func = &F; + } + } + } + + llvm::Function* ref_func = osl_init_func ? osl_init_func : osl_layer_func; + if (!ref_func) { + shadingsys().error("[LLVM AMDGPU] BLAD! Nie znaleziono zadnej funkcji OSL do owrapowania!"); + return std::vector(); + } + + shadingsys().info("[LLVM AMDGPU] Znaleziono funkcje OSL. Tworzenie wrappera..."); + + // 6. DYNAMICZNA BUDOWA WRAPPERA (LAUNCHER KERNEL) O NAZWIE "osl_kernel" + llvm::LLVMContext &ctx = mod->getContext(); + llvm::FunctionType *orig_type = ref_func->getFunctionType(); + std::vector wrapper_param_types; + + for (llvm::Type *param_type : orig_type->params()) { + if (param_type->isPointerTy()) { + wrapper_param_types.push_back(llvm::PointerType::get(ctx, 1)); + } else { + wrapper_param_types.push_back(param_type); + } + } + wrapper_param_types.push_back(llvm::Type::getInt32Ty(ctx)); // + int width + + llvm::FunctionType *wrapper_type = llvm::FunctionType::get(llvm::Type::getVoidTy(ctx), wrapper_param_types, false); + llvm::Function *wrapper_func = llvm::Function::Create(wrapper_type, llvm::GlobalValue::ExternalLinkage, "osl_kernel", mod); + wrapper_func->setCallingConv(llvm::CallingConv::AMDGPU_KERNEL); + wrapper_func->addFnAttr("amdgpu-flat-work-group-size", "1,1024"); + + llvm::BasicBlock *entry_block = llvm::BasicBlock::Create(ctx, "entry", wrapper_func); + llvm::IRBuilder<> builder(entry_block); + + // 7. WSTRZYKIWANIE KOORDYNATÓW WĄTKÓW + llvm::FunctionType *i32_fn_type = llvm::FunctionType::get(builder.getInt32Ty(), false); + llvm::FunctionCallee wg_id_x_c = mod->getOrInsertFunction("llvm.amdgcn.workgroup.id.x", i32_fn_type); + llvm::FunctionCallee wg_id_y_c = mod->getOrInsertFunction("llvm.amdgcn.workgroup.id.y", i32_fn_type); + llvm::FunctionCallee wi_id_x_c = mod->getOrInsertFunction("llvm.amdgcn.workitem.id.x", i32_fn_type); + llvm::FunctionCallee wi_id_y_c = mod->getOrInsertFunction("llvm.amdgcn.workitem.id.y", i32_fn_type); + + llvm::Value *wg_id_x = builder.CreateCall(wg_id_x_c); + llvm::Value *wg_id_y = builder.CreateCall(wg_id_y_c); + llvm::Value *wi_id_x = builder.CreateCall(wi_id_x_c); + llvm::Value *wi_id_y = builder.CreateCall(wi_id_y_c); + + llvm::Value *block_dim = builder.getInt32(16); + llvm::Value *global_x = builder.CreateAdd(builder.CreateMul(wg_id_x, block_dim), wi_id_x); + llvm::Value *global_y = builder.CreateAdd(builder.CreateMul(wg_id_y, block_dim), wi_id_y); + + llvm::Argument *width_arg = wrapper_func->getArg(ref_func->arg_size()); + llvm::Value *pixel_index = builder.CreateAdd(builder.CreateMul(global_y, width_arg), global_x); + + // 8. OBLICZANIE WSKAŹNIKÓW I WYWOŁANIE FUNKCJI + std::vector call_args; + auto orig_arg_it = ref_func->arg_begin(); + size_t arg_idx = 0; + + // Pobieramy rozmiar struktury GroupData (zabezpieczone min. 16 bajtów) + int gd_size = std::max((int)group().llvm_groupdata_size(), 16); + for (llvm::Argument &wrap_arg : wrapper_func->args()) { + if (arg_idx >= ref_func->arg_size()) break; -}; // namespace pvt + llvm::Value *final_val = &wrap_arg; + + if (wrap_arg.getType()->isPointerTy()) { + if (arg_idx == 1) { + llvm::Value *local_gd = builder.CreateAlloca(builder.getInt8Ty(), builder.getInt32(gd_size), "local_group_data"); + final_val = builder.CreateAddrSpaceCast(local_gd, orig_arg_it->getType(), "cast_to_flat"); + } else { + llvm::Value *offset = nullptr; + if (arg_idx == 0) { + offset = builder.CreateMul(pixel_index, builder.getInt32(256)); + } else if (arg_idx == 3) { + offset = nullptr; + } + + if (offset) { + final_val = builder.CreateGEP(builder.getInt8Ty(), &wrap_arg, offset); + } + final_val = builder.CreateAddrSpaceCast(final_val, orig_arg_it->getType(), "cast_to_flat"); + } + } + + else if (arg_idx == 4) { + final_val = pixel_index; + } + + call_args.push_back(final_val); + orig_arg_it++; + arg_idx++; + } + + // SEKWENCYJNE WYWOŁANIE FUNKCJI OSL + if (osl_init_func) { + builder.CreateCall(orig_type, osl_init_func, call_args); + } + if (osl_layer_func) { + builder.CreateCall(orig_type, osl_layer_func, call_args); + } + + builder.CreateRetVoid(); + + // 8.5 WYMUSZENIE WIDOCZNOŚCI I ATRYBUTÓW KERNELA (Dla LLVM 18+) + // 1. Zmieniamy całą tożsamość modułu na AMDGPU + mod->setTargetTriple(target_machine->getTargetTriple().str()); + mod->setDataLayout(target_machine->createDataLayout()); + + // 2. Oznaczamy nasz wrapper jako KERNEL sprzętowy + wrapper_func->setCallingConv(llvm::CallingConv::AMDGPU_KERNEL); + wrapper_func->setLinkage(llvm::GlobalValue::ExternalLinkage); + wrapper_func->setVisibility(llvm::GlobalValue::DefaultVisibility); + wrapper_func->addFnAttr("amdgpu-kernel"); + + shadingsys().info("[LLVM AMDGPU] ---> Tożsamość modułu zmieniona na AMDGPU. Wrapper ustawiony jako KERNEL."); + + // DIAGNOZA MODUŁU LLVM - SPRAWDZAMY, CZY FUNKCJE MAJĄ CIAŁA! + + llvm::errs() << "\n[!!! DIAGNOZA LLVM !!!]\n"; + llvm::errs() << "Liczba wszystkich funkcji w module: " << mod->getFunctionList().size() << "\n"; + + for (llvm::Function &F : *mod) { + // Wypisujemy tylko interesujące nas funkcje (nasz kernel i oryginalny kod shadera) + if (F.getName() == "osl_kernel" || F.getName().contains("group_") || F.getName().contains("shader")) { + llvm::errs() << " -> Funkcja: " << F.getName() + << " | Czy jest pusta (tylko deklaracja)? " << (F.empty() ? "TAK (Brak kodu!)" : "NIE (Ma kod!)") << "\n"; + } + } + llvm::errs() << "[!!! KONIEC DIAGNOZY !!!]\n\n"; + + // 8.6 ZRZUT WYGENEROWANEGO KODU IR DO PLIKU + + std::error_code EC; + llvm::raw_fd_ostream ir_file("/tmp/osl_ir_dump.ll", EC); + if (!EC) { + mod->print(ir_file, nullptr); + ir_file.close(); + shadingsys().info("[LLVM AMDGPU] Zapisano zrzut kodu LLVM IR do /tmp/osl_ir_dump.ll"); + } else { + shadingsys().error("[LLVM AMDGPU] Nie udalo sie zapisac pliku IR dump."); + } + + // Ukrywamy oryginalne funkcje, żeby LLVM skupił się na wyeksportowaniu tylko wrappera + if (osl_init_func) osl_init_func->setLinkage(llvm::GlobalValue::InternalLinkage); + if (osl_layer_func) osl_layer_func->setLinkage(llvm::GlobalValue::InternalLinkage); + // 8.7 USUWANIE ATRYBUTÓW CPU I WERYFIKACJA + for (llvm::Function &F : *mod) { + F.removeFnAttr("target-cpu"); + F.removeFnAttr("target-features"); + F.removeFnAttr("tune-cpu"); + + F.removeFnAttr(llvm::Attribute::OptimizeNone); + F.removeFnAttr(llvm::Attribute::NoInline); + } + + shadingsys().info("[LLVM AMDGPU] Atrybuty CPU zostały usunięte. Rozpoczynam weryfikację kodu IR..."); + + // 9. EMISJA KODU + llvm::SmallVector elf_buffer; + llvm::raw_svector_ostream dest(elf_buffer); + + llvm::legacy::PassManager code_gen_pm; + if (target_machine->addPassesToEmitFile(code_gen_pm, dest, nullptr, llvm::CodeGenFileType::ObjectFile)) { + shadingsys().error("LLVM Error: Backend kompilatora LLVM nie wspiera bezpośredniej emisji ELF dla AMDGPU"); + return std::vector(); + } + + std::cout << "[DEBUG-LLVM] 3. Uruchamiam code_gen_pm.run(*mod) -> To moze chwile potrwac..."<> g_amdgpu_elf_registry; + extern std::mutex g_amdgpu_registry_mutex; + + std::vector elf_blob(elf_buffer.begin(), elf_buffer.end()); + { + std::lock_guard lock(g_amdgpu_registry_mutex); + g_amdgpu_elf_registry[&group()] = elf_blob; + } + return elf_blob; +}} +// namespace pvt OSL_NAMESPACE_END diff --git a/src/liboslexec/backendllvm.h b/src/liboslexec/backendllvm.h index ed5f4fdeb9..63471c9c9d 100644 --- a/src/liboslexec/backendllvm.h +++ b/src/liboslexec/backendllvm.h @@ -44,7 +44,8 @@ class BackendLLVM final : public OSOProcessorBase { /// Set additional Module/Function options for the CUDA/OptiX target. void prepare_module_for_cuda_jit(); - +// NEW - wyciąganie binarnego bitkodu LLVM dla AMDGPU + std::vector get_llvm_bitcode(llvm::Module* custom_mod = nullptr); /// What LLVM debug level are we at? int llvm_debug() const; diff --git a/src/liboslexec/llvm_instance.cpp b/src/liboslexec/llvm_instance.cpp index 209799b622..3b4d285c63 100644 --- a/src/liboslexec/llvm_instance.cpp +++ b/src/liboslexec/llvm_instance.cpp @@ -2379,7 +2379,7 @@ BackendLLVM::run() // If we plan to call bitcode_string of a layer's function after // optimization it may not exist after optimization unless we // treat it as external. - if (f && (group().is_entry_layer(layer) || llvm_debug())) { + if (f && (group().is_entry_layer(layer) || llvm_debug() || shadingsys().use_amdgpu())) { external_functions.insert(f); } } @@ -2500,6 +2500,22 @@ BackendLLVM::run() } } else #endif +// NEW +if (shadingsys().use_amdgpu() || !shadingsys().amdgpu_architecture().string().empty()) { + // 1. Zlecamy NASZEJ klasie wyciągnięcie binarnego bitkodu z pamięci + std::vector bitcode = get_llvm_bitcode(init_func->getParent()); + + if (bitcode.empty()) { + OSL_ASSERT(0 && "Unable to generate AMDGPU Bitcode"); + } else { + std::cout << "[DEBUG-LLVM] Sukces: Artefakt AMDGPU dodany do grupy.\n"; // Dodaj to + // 2. Pakujemy gotowe bajty w nowy Artefakt i dodajemy do wektora + group().m_compiled_gpu_artifacts.push_back( + CompiledGPUArtifact(bitcode, shadingsys().amdgpu_architecture().string(), "llvm_bitcode", OSL_LLVM_VERSION) + ); + } + } else +// END NEW { // Force the JIT to happen now and retrieve the JITed function pointers // for the initialization and all public entry points. @@ -2539,9 +2555,9 @@ BackendLLVM::run() } ll.delete_func_body (init_func); #endif - // Free the exec and module to reclaim all the memory. This definitely // saves memory, and has almost no effect on runtime. + ll.execengine(NULL); // N.B. Destroying the EE should have destroyed the module as well. diff --git a/src/liboslexec/llvm_util.cpp b/src/liboslexec/llvm_util.cpp index 7a9db27bf2..9000772a8c 100644 --- a/src/liboslexec/llvm_util.cpp +++ b/src/liboslexec/llvm_util.cpp @@ -42,6 +42,7 @@ OSL_GCC_PRAGMA(GCC diagnostic ignored "-Wmaybe-uninitialized") #include #include #include +#include #include #include #include @@ -122,6 +123,11 @@ OSL_GCC_PRAGMA(GCC diagnostic ignored "-Wmaybe-uninitialized") #include #include +// NEW +#include +#include +#include + OSL_PRAGMA_WARNING_POP OSL_NAMESPACE_BEGIN @@ -1856,6 +1862,62 @@ LLVM_Util::nvptx_target_machine() return m_nvptx_target_machine; } +// NEW +llvm::TargetMachine* +LLVM_Util::amdgpu_target_machine(const std::string& gpu_arch) +{ + if (m_amdgpu_target_machine == nullptr) { + llvm::Triple ModuleTriple(module()->getTargetTriple()); + llvm::TargetOptions options; + + options.AllowFPOpFusion = llvm::FPOpFusion::Standard; +#if OSL_LLVM_VERSION < 220 + options.UnsafeFPMath = 1; +#endif + options.NoInfsFPMath = 1; + options.NoNaNsFPMath = 1; + options.HonorSignDependentRoundingFPMathOption = 0; + options.FloatABIType = llvm::FloatABI::Default; + options.AllowFPOpFusion = llvm::FPOpFusion::Fast; + options.NoZerosInBSS = 0; + options.GuaranteedTailCallOpt = 0; + options.UseInitArray = 0; + + std::string error; +#if OSL_LLVM_VERSION >= 220 + const llvm::Target* llvm_target + = llvm::TargetRegistry::lookupTarget(ModuleTriple, error); +#else + const llvm::Target* llvm_target + = llvm::TargetRegistry::lookupTarget(ModuleTriple.str(), error); +#endif + OSL_ASSERT(llvm_target + && "AMDGPU compile error: LLVM Target is not initialized"); + + m_amdgpu_target_machine = llvm_target->createTargetMachine( +#if OSL_LLVM_VERSION >= 210 + llvm::Triple(ModuleTriple.str()), +#else + ModuleTriple.str(), +#endif + gpu_arch, + "", + options, + llvm::Reloc::PIC_, + llvm::CodeModel::Small, +#if OSL_LLVM_VERSION >= 180 + llvm::CodeGenOptLevel::Default +#else + llvm::CodeGenOpt::Default +#endif + ); + + OSL_ASSERT(m_amdgpu_target_machine + && "Unable to create TargetMachine for AMDGPU"); + } + return m_amdgpu_target_machine; +} + void* @@ -6453,6 +6515,77 @@ LLVM_Util::ptx_compile_group(llvm::Module*, const std::string& name, #endif } +// NEW +bool +LLVM_Util::amdgpu_compile_group(llvm::Module* mod, const std::string& name, + std::string& out, const std::string& gpu_arch) +{ + + std::cout << "\n[LLVM AMDGPU] --- Szukanie funkcji w bitcode ---\n"; + llvm::Function* kernel_func = mod->getFunction(name); + if (!kernel_func) { + for (auto& F : *mod) { + if (!F.isDeclaration() && (F.getName().contains("group") || F.getName().contains("test"))) { + kernel_func = &F; + break; + } + } + } + + + if (kernel_func) { + // Wymuszamy publiczną, globalną widoczność w tablicy symboli ELF + kernel_func->setLinkage(llvm::GlobalValue::ExternalLinkage); + + // Ustawiamy konwencję wywołania jako KERNEL karty AMD + kernel_func->setCallingConv(llvm::CallingConv::AMDGPU_KERNEL); + + // Nadajemy jej sztywną nazwę, którą wczytamy w hipModuleGetFunction + kernel_func->setName("osl_kernel"); + + std::cout << "[LLVM Backend] Sukces: Wyeksportowano publiczny symbol 'osl_kernel' dla GPU!\n"; + } + std::cout << "[LLVM AMDGPU] ----------------------------------\n\n"; + + /* + for (llvm::Function &F : *module()) { + if (!F.isDeclaration()) { + std::string fname = F.getName().str(); + std::cout << "[LLVM AMDGPU] Zdefiniowano: " << fname << "\n"; + + if (fname.find("group") != std::string::npos || fname.find("exec") != std::string::npos) { + std::cout << "[LLVM AMDGPU] ---> ZNALEZIONO KERNEL! Ustawiam AMDGPU_KERNEL i ExternalLinkage\n"; + + F.setCallingConv(llvm::CallingConv::AMDGPU_KERNEL); + F.setLinkage(llvm::GlobalValue::ExternalLinkage); + } + } + } + std::cout << "[LLVM AMDGPU] ----------------------------------\n\n"; + // ---------------------------------------------------------------*/ + + llvm::TargetMachine* target_machine = amdgpu_target_machine(gpu_arch); + llvm::legacy::PassManager mpm; + + llvm::SmallString<4096> binary_buffer; + llvm::raw_svector_ostream binary_stream(binary_buffer); + +#if OSL_LLVM_VERSION >= 180 + target_machine->addPassesToEmitFile(mpm, binary_stream, + nullptr, + llvm::CodeGenFileType::ObjectFile); +#else + target_machine->addPassesToEmitFile(mpm, binary_stream, + nullptr, + llvm::CGFT_ObjectFile); +#endif + + mpm.run(*mod); + + out = binary_buffer.str().str(); + + return true; +} std::string diff --git a/src/liboslexec/oslexec_pvt.h b/src/liboslexec/oslexec_pvt.h index 9d7acc6175..92a3d0fb7d 100644 --- a/src/liboslexec/oslexec_pvt.h +++ b/src/liboslexec/oslexec_pvt.h @@ -66,7 +66,35 @@ namespace Strutil = OIIO::Strutil; OSL_NAMESPACE_BEGIN +// NEW +enum class GPUBackendKind { + None = 0, + NVPTX = 1, // NEW for CUDA/OptiX + AMDGPU = 2 // NEW for HIP/ROCm +}; + +struct GPUTargetDesc { + GPUBackendKind backend = GPUBackendKind::None; + std::string triple; + std::vector archs; + std::string artifact_format; + + GPUTargetDesc() = default; + GPUTargetDesc(GPUBackendKind b, const std::string& t, const std::vector& a, const std::string& f) + : backend(b), triple(t), archs(a), artifact_format(f) {} +}; + +struct CompiledGPUArtifact { + std::vector payload; + std::string architecture; + std::string format; + int llvm_version; + CompiledGPUArtifact() : llvm_version(0) {} + CompiledGPUArtifact(const std::vector& p, const std::string& arch, const std::string& fmt, int llvm_ver) + : payload(p), architecture(arch), format(fmt), llvm_version(llvm_ver) {} +}; +// END NEW struct PerThreadInfo { PerThreadInfo(); @@ -639,6 +667,13 @@ class ShadingSystemImpl { bool use_optix() const { return m_use_optix; } bool use_optix_cache() const { return m_use_optix_cache; } + + // NEW + bool use_amdgpu() const { return m_use_amdgpu; } + bool use_amdgpu_cache() const { return m_use_amdgpu_cache; } + ustring amdgpu_architecture() const { return m_amdgpu_architecture; } + // END NEW + bool debug_nan() const { return m_debugnan; } bool debug_uninit() const { return m_debug_uninit; } bool lockgeom_default() const { return m_lockgeom_default; } @@ -965,6 +1000,12 @@ class ShadingSystemImpl { int m_compile_report; ///< Print compilation report? bool m_use_optix; ///< This is an OptiX-based renderer bool m_use_optix_cache; ///< Renderer-enabled caching for OptiX ptx + // NEW + bool m_use_amdgpu = false; + bool m_use_amdgpu_cache = true; + ustring m_amdgpu_architecture; + void amdgpu_cache_unwrap(const std::string& cache_value, ShaderGroup& group); + // END NEW int m_max_optix_groupdata_alloc; ///< Maximum OptiX groupdata buffer allocation bool m_buffer_printf; ///< Buffer/batch printf output? bool m_no_noise; ///< Substitute trivial noise calls @@ -984,9 +1025,12 @@ class ShadingSystemImpl { int m_optix_no_inline_thresh; ///< Disable inlining for functions larger than the threshold int m_optix_force_inline_thresh; ///< Force inling for functions smaller than the threshold + ustring m_colorspace; ///< What RGB colors mean ShadingStateUniform m_shading_state_uniform; + + OSL::GPUTargetDesc m_gpu_target; //NEW std::shared_ptr m_colorconfig; ///< OIIO/OCIO color configuration @@ -1855,6 +1899,13 @@ class ShaderGroup { void generate_optix_cache_key(string_view code); std::string optix_cache_key() const { return m_optix_cache_key; } + // NEW + std::string amdgpu_cache_key() const { + return m_name.string(); + } + // END NEW + + std::string serialize() const; void lock() const { m_mutex.lock(); } @@ -2062,6 +2113,10 @@ class ShaderGroup { // PTX assembly for compiled ShaderGroup std::string m_llvm_ptx_compiled_version; + + // NASZ KONTENER NA BINARIA AMDGPU + std::vector m_amdgpu_elf; //NEWNEW + std::vector m_compiled_gpu_artifacts; //NEWNEW ParamValueList m_pending_params; // Pending Parameter() values std::vector m_pending_hints; // ParamHints of pending params diff --git a/src/liboslexec/shadingsys.cpp b/src/liboslexec/shadingsys.cpp index 671b396fcb..7890a91044 100644 --- a/src/liboslexec/shadingsys.cpp +++ b/src/liboslexec/shadingsys.cpp @@ -29,6 +29,18 @@ #include "opcolor.h" +// NEW + +#include +#include + +// Odniesienie do mapy z backendllvm.cpp +extern std::map g_amdgpu_temp_cache; + +// NEW +#include +// END NEW + using namespace OSL; using namespace OSL::pvt; @@ -58,6 +70,29 @@ extern unsigned char shadeops_cuda_ptx_compiled_ops_block[]; OSL_NAMESPACE_BEGIN +// NEW +namespace pvt { + +// NEW +std::map> g_amdgpu_elf_registry; +std::mutex g_amdgpu_registry_mutex; + +} + +// NEW +static OSL::GPUTargetDesc make_amdgpu_target(const std::vector& archs) { + return OSL::GPUTargetDesc(OSL::GPUBackendKind::AMDGPU, + "amdgcn-amd-amdhsa", + archs, + "llvm_bitcode"); +} + +static OSL::GPUTargetDesc make_nvptx_target() { + return OSL::GPUTargetDesc(OSL::GPUBackendKind::NVPTX, + "nvptx64-nvidia-cuda", + std::vector{"sm_50"}, + "ptx"); +} ShadingSystem::ShadingSystem(RendererServices* renderer, @@ -1155,6 +1190,38 @@ ShadingSystemImpl::ShadingSystemImpl(RendererServices* renderer, , m_stat_inst_merge_time(0) , m_stat_max_llvm_local_mem(0) { + + // NEW +if (renderer && (renderer->supports("AMDGPU") || renderer->supports("HIP"))) { + std::string hip_arch_str = "gfx1100"; + renderer->get_attribute(nullptr, false, OSL::ustring(), TypeDesc::STRING, OSL::ustring("amdgpu_architecture"), &hip_arch_str); + + // 3. Rozbijamy string (np. "gfx1030,gfx1100") na listę architektur + std::vector archs; + std::stringstream ss(hip_arch_str); + std::string item; + + while (std::getline(ss, item, ',')) { + // Usuwamy ewentualne białe znaki, żeby uniknąć błędów + item.erase(std::remove_if(item.begin(), item.end(), ::isspace), item.end()); + if (!item.empty()) { + archs.push_back(item); + } + } + + // Zabezpieczenie: jeśli string był pusty, wrzucamy domyślną + if (archs.empty()) { + archs.push_back("gfx1100"); + } + + // 4. Przekazujemy gotową listę do Twojej nowej funkcji! + m_gpu_target = make_amdgpu_target(archs); +} +else if (m_use_optix) { + m_gpu_target = make_nvptx_target(); +} + + m_shading_state_uniform.m_commonspace_synonym = Strings::world; m_shading_state_uniform.m_unknown_coordsys_error = true; m_shading_state_uniform.m_max_warnings_per_thread = 100; @@ -1577,7 +1644,22 @@ osl_simd_caps() // clang-format on } - +// NEW +void +ShadingSystemImpl::amdgpu_cache_unwrap(const std::string& cache_value, ShaderGroup& group) +{ + // 1. Zamieniamy string z cache na wektor bajtów + std::vector data(cache_value.begin(), cache_value.end()); + + // 2. Synchronizacja: zapisujemy do pola, z którego korzysta rejestr + group.m_amdgpu_elf = data; + + // 3. Pakujemy to do kontenera OSL + group.m_compiled_gpu_artifacts.push_back( + CompiledGPUArtifact(data, m_amdgpu_architecture.string(), "llvm_bitcode", OSL_LLVM_VERSION) + ); +} +// END NEW bool ShadingSystemImpl::attribute(string_view name, TypeDesc type, const void* val) @@ -1604,6 +1686,11 @@ ShadingSystemImpl::attribute(string_view name, TypeDesc type, const void* val) } lock_guard guard(m_mutex); // Thread safety + // NEW + ATTR_SET("amdgpu", int, m_use_amdgpu); + ATTR_SET("amdgpu_cache", int, m_use_amdgpu_cache); + ATTR_SET_STRING("amdgpu_architecture", m_amdgpu_architecture); + // END NEW ATTR_SET("statistics:level", int, m_statslevel); ATTR_SET("debug", int, m_debug); ATTR_SET("lazylayers", int, m_lazylayers); @@ -2092,6 +2179,58 @@ ShadingSystemImpl::getattribute(ShaderGroup* group, string_view name, if (!group) return false; + // NEW + if (group) { + if (name == "gpu_num_artifacts" && type == TypeDesc::INT) { + if (val) + *(int*)val = (int)group->m_compiled_gpu_artifacts.size(); + return true; + } + + if (Strutil::starts_with(name, "gpu_artifact:")) { + int index = 0; + char prop[32] = {0}; + + // NEW + std::string name_str(name); + if (sscanf(name_str.c_str(), "gpu_artifact:%d:%31s", &index, prop) == 2) { + + // if (sscanf(name.c_str(), "gpu_artifact:%d:%31s", &index, prop) == 2) { + if (index >= 0 && index < (int)group->m_compiled_gpu_artifacts.size()) { + const auto& artifact = group->m_compiled_gpu_artifacts[index]; + string_view prop_str(prop); + + if (prop_str == "data" && type == TypeDesc::PTR) { + if (val) + *(const void**)val = artifact.payload.data(); + return true; + } + // NEW + else if (prop_str == "size") { + if (type == TypeDesc::INT) { + if (val) *(int*)val = (int)artifact.payload.size(); + return true; + } else if (type == TypeDesc::INT64) { + if (val) *(int64_t*)val = (int64_t)artifact.payload.size(); + return true; + } + } // END NEW + + else if (prop_str == "architecture" && type == TypeDesc::STRING) { + if (val) + *(ustring*)val = ustring(artifact.architecture); + return true; + } + else if (prop_str == "format" && type == TypeDesc::STRING) { + if (val) + *(ustring*)val = ustring(artifact.format); + return true; + } + } + } + } + } + if (name == "groupname" && type == TypeString) { *(ustring*)val = group->name(); return true; @@ -3714,7 +3853,11 @@ ShadingSystemImpl::group_post_jit_cleanup(ShaderGroup& group) } } - + //NEWNEW +void register_amdgpu_artifact_in_registry(const void* group_ptr, const std::vector& elf_data) { + std::lock_guard lock(g_amdgpu_registry_mutex); + g_amdgpu_elf_registry[group_ptr] = elf_data; +} void ShadingSystemImpl::optimize_group(ShaderGroup& group, ShadingContext* ctx, @@ -3832,8 +3975,69 @@ ShadingSystemImpl::optimize_group(ShaderGroup& group, ShadingContext* ctx, group.m_llvm_groupdata_size); } } + // NEW + else if (use_amdgpu_cache()) { + // 1. Pobieramy unikalny klucz dla tej grupy shaderów + std::string cache_key = group.amdgpu_cache_key() + + "_AMDGPU_" + + m_amdgpu_architecture.string(); + + std::string cache_value; + // 2. Szukamy w systemowym cache OSL pod kluczem "amdgpu_bc" + std::string filename = "/tmp/" + cache_key + ".bin"; + std::ifstream in_file(filename, std::ios::binary); + + if (in_file.good()) { + // Wczytujemy cały plik do stringa + std::stringstream buffer; + buffer << in_file.rdbuf(); + cache_value = buffer.str(); + cached = true; + amdgpu_cache_unwrap(cache_value, group); + // NEW + register_amdgpu_artifact_in_registry(&group, group.m_amdgpu_elf); + + static std::once_flag cache_logged_flag; + std::call_once(cache_logged_flag, [this]() { + this->info("INFO: Zastosowano AMDGPU Cache z dysku..."); + }); + } + } if (!cached) { + // NEW + // Sprawdzamy, czy użytkownik zażądał kompilacji na AMDGPU +// if (use_amdgpu_cache() || !m_amdgpu_architecture.string().empty()) { + +// // 1. Tworzymy instancję backendu LLVM +// BackendLLVM lljitter(*this, group, ctx); + +// std::cout << "[DEBUG] Backend utworzony. Wywołuję run()...\n"; + +// // 2. Czy lljitter wymaga jawnego uruchomienia? Spróbuj odkomentować/dodać: +// lljitter.run(); + +// // 3. Teraz pobierz bitcode +// std::vector amd_elf = lljitter.get_llvm_bitcode(); + +// std::cout << "[DEBUG] Rozmiar pobranego bitkodu: " << amd_elf.size() << " bajtów.\n"; + +// if (amd_elf.empty()) { +// error(OIIO::Strutil::format("Blad: Kompilacja do AMDGPU ELF dla grupy %s nie powiodla sie!", group.name().c_str())); +// } else { +// // 3. Zapisujemy wygenerowane artefakty w grupie shaderów +// group.m_amdgpu_elf = amd_elf; + +// // NEW +// register_amdgpu_artifact_in_registry(&group, amd_elf); + +// info(OIIO::Strutil::format("Sukces: Wygenerowano i zaladowano AMDGPU ELF dla %s", group.name().c_str())); +// } + +// // Ustawiamy flagę, że grupa została przetworzona, aby uniknąć standardowego czyszczenia CPU +// group.m_jitted = true; + +// } else { BackendLLVM lljitter(*this, group, ctx); lljitter.run(); @@ -3852,6 +4056,7 @@ ShadingSystemImpl::optimize_group(ShaderGroup& group, ShadingContext* ctx, } group.m_jitted = true; + spin_lock stat_lock(m_stat_mutex); m_stat_opt_locking_time += locking_time; m_stat_optimization_time += timer(); @@ -3862,6 +4067,7 @@ ShadingSystemImpl::optimize_group(ShaderGroup& group, ShadingContext* ctx, m_stat_llvm_jit_time += lljitter.m_stat_llvm_jit_time; m_stat_max_llvm_local_mem = std::max(m_stat_max_llvm_local_mem, lljitter.m_llvm_local_mem); + } } @@ -4831,3 +5037,93 @@ osl_incr_get_userdata_calls(void* sg_) ShaderGlobals* sg = (ShaderGlobals*)sg_; sg->context->incr_get_userdata_calls(); } + +//NEW + + +extern "C" OSLEXECPUBLIC bool OSL_get_amdgpu_artifact(const void* group_ptr, uint8_t* out_buffer, size_t* out_size) { + std::lock_guard lock(g_amdgpu_registry_mutex); + + // Szukamy używając wskaźnika! Bez konwersji na std::string + auto it = g_amdgpu_elf_registry.find(group_ptr); + + if (it == g_amdgpu_elf_registry.end()) { + return false; // Nie znaleziono artefaktu dla tej grupy + } + + // Jeśli podano bufor, kopiujemy dane + if (out_buffer != nullptr) { + std::memcpy(out_buffer, it->second.data(), it->second.size()); + } + + // Zwracamy rozmiar, żeby runtime wiedział ile pamięci zaalokować + if (out_size != nullptr) { + *out_size = it->second.size(); + } + + return true; +} + + +class DummyRenderer : public OSL::RendererServices { +public: + DummyRenderer() : OSL::RendererServices() {} +}; + +extern "C" OSLEXECPUBLIC bool OSL_run_gpu_backend(const std::string& oso_file, const std::string& arch) { + auto texturesys = OIIO::TextureSystem::create(); + auto renderer = new DummyRenderer(); + OSL::ShadingSystem* ss = new OSL::ShadingSystem(renderer, texturesys); + + ss->attribute("amdgpu_architecture", arch); + ss->attribute("amdgpu", 1); + + + int opt = 0; + ss->attribute("optimize", opt); + ss->attribute("opt_elide_unconnected_outputs", 0); + ss->attribute("lazyunconnected", 0); + + // Zabezpieczenie nr 2: Wymuszamy, by optymalizator uważał wyjścia `Ci` i `Cout` za potrzebne + const char* aovs[] = {"Ci", "Cout"}; + ss->attribute("renderer_outputs", OSL::TypeDesc(OSL::TypeDesc::STRING, 2), &aovs[0]); + + OSL::ShaderGroupRef group = ss->ShaderGroupBegin(); + if (!ss->Shader("surface", oso_file, "")) { + delete ss; delete renderer; return false; + } + ss->ShaderGroupEnd(*group); + OSL::SymLocationDesc symloc(OSL::ustring("Cout"), OSL::TypeDesc::TypeColor, false, OSL::SymArena::Outputs, 0, 12); + std::vector symlocs = { symloc }; + ss->add_symlocs(group.get(), symlocs); + + std::cout << "[DEBUG-SHADINGSYS] Odpalam optymalizacje i JIT LLVM...\n"; + // Zlecenie wygenerowania kodu i obiektu .o przez backend + ss->optimize_group(group.get(), nullptr, true); + std::cout << "[DEBUG-SHADINGSYS] Optymalizacja zrobiona. LLVM zakonczyl prace.\n"; + + // Sprawdzamy, czy LLVM faktycznie wypuścił plik. Jeśli nie, logujemy ładny błąd zamiast crasza Clanga. + std::string temp_o_file = "/tmp/osl_temp_shader.o"; + if (!OIIO::Filesystem::exists(temp_o_file)) { + std::cerr << "FATAL: Plik .o nie zostal wygenerowany przez LLVM!\n"; + std::cerr << "Prawdopodobna przyczyna: OSL uznał, że shader nic nie robi (does_nothing == true).\n"; + delete ss; delete renderer; return false; + } + + std::cout << "[DEBUG-SHADINGSYS] Linkuje do .hsaco...\n"; + + std::string base_name = oso_file.substr(0, oso_file.find_last_of('.')); + std::string hsaco_file = base_name + ".hsaco"; + + std::string link_cmd = "clang -target amdgcn-amd-amdhsa -mcpu=" + arch + " -shared " + temp_o_file + " -o " + hsaco_file; + if (std::system(link_cmd.c_str()) != 0) { + std::cerr << "FATAL: Linker (clang) zwrócił błąd!\n"; + delete ss; delete renderer; return false; + } + + OIIO::Filesystem::remove(temp_o_file); + std::cout << "[oslc] SUKCES! Gotowy plik jądra zapisany jako: " << hsaco_file << "\n"; + + delete ss; delete renderer; + return true; +} \ No newline at end of file diff --git a/src/oslc/CMakeLists.txt b/src/oslc/CMakeLists.txt index 7750fd1413..cac68dc5fa 100644 --- a/src/oslc/CMakeLists.txt +++ b/src/oslc/CMakeLists.txt @@ -13,5 +13,5 @@ endif () add_executable ( oslc ${oslc_srcs} ) target_include_directories ( oslc BEFORE PRIVATE ${OpenImageIO_INCLUDES}) -target_link_libraries ( oslc PRIVATE oslcomp ${CMAKE_DL_LIBS}) +target_link_libraries ( oslc PRIVATE oslcomp oslexec ${CMAKE_DL_LIBS}) install_targets (oslc) diff --git a/src/oslc/oslcmain.cpp b/src/oslc/oslcmain.cpp index e74af47d8c..6fb733b984 100644 --- a/src/oslc/oslcmain.cpp +++ b/src/oslc/oslcmain.cpp @@ -15,6 +15,8 @@ #include #include +extern "C" bool OSL_run_gpu_backend(const std::string& oso_file, const std::string& arch); + using namespace OSL; @@ -43,7 +45,9 @@ usage() "\t-MD, -MMD Write a depfile containing headers used, to a file\n" "\t-M, -MM Like -MD, but write depfile to stdout\n" "\t-MF filename Specify the name of the depfile to output (for -MD, -MMD)\n" - "\t-MT target Specify a custom dependency target name for -M...\n"; + "\t-MT target Specify a custom dependency target name for -M...\n" + "\t--amdgpu Compile for AMDGPU backend\n" + "\t--device Specify target AMD device (e.g., gfx1030, gfx1100)\n"; } @@ -93,6 +97,7 @@ main(int argc, const char* argv[]) // Globally force classic "C" locale, and turn off all formatting // internationalization, for the entire oslc application. std::locale::global(std::locale::classic()); + #ifdef OIIO_HAS_STACKTRACE // Helpful for debugging to make sure that any crashes dump a stack @@ -112,6 +117,10 @@ main(int argc, const char* argv[]) bool compile_from_buffer = false; std::string shader_path; + // NEW + [[maybe_unused]] bool use_amdgpu = false; + [[maybe_unused]] std::string amd_device = "gfx1030"; + // Parse arguments from command line for (int a = 1; a < argc; ++a) { if (!strcmp(argv[a], "--help") | !strcmp(argv[a], "-h")) { @@ -154,6 +163,15 @@ main(int argc, const char* argv[]) && (argv[a][1] == 'D' || argv[a][1] == 'U' || argv[a][1] == 'I')) { args.emplace_back(argv[a]); + } else if (!strcmp(argv[a], "--amdgpu")) { + use_amdgpu = true; + // NEW + args.emplace_back(argv[a]); + } else if (!strcmp(argv[a], "--device") && a < argc - 1) { + args.emplace_back(argv[a]); + ++a; + amd_device = argv[a]; + args.emplace_back(argv[a]); } else if (!strcmp(argv[a], "-buffer")) { compile_from_buffer = true; } else { @@ -190,11 +208,23 @@ main(int argc, const char* argv[]) // Ordinary compile from file ok = compiler.compile(shader_path, args); } - + if (ok) { if (!quiet) std::cout << "Compiled " << shader_path << " -> " << compiler.output_filename() << "\n"; + if (use_amdgpu) { + std::cout << "[oslc] Generowanie jądra AMDGPU dla architektury: " << amd_device << "...\n"; + + // Przekazujemy ścieżkę do skompilowanego .oso do bezpiecznego backendu + bool success = OSL_run_gpu_backend(compiler.output_filename().data(), amd_device); + + if (!success) { + std::cerr << "FATAL: Backend GPU zwrócił błąd!\n"; + return EXIT_FAILURE; + } + std::cout << "[oslc] SUKCES! Gotowy plik .hsaco znajduje się na dysku.\n"; + } } else { std::cout << "FAILED " << shader_path << "\n"; return EXIT_FAILURE; diff --git a/src/shaders/matte.oso b/src/shaders/matte.oso new file mode 100644 index 0000000000..fc6161e902 --- /dev/null +++ b/src/shaders/matte.oso @@ -0,0 +1,18 @@ +OpenShadingLanguage 1.00 +# Compiled by oslc 1.16.0.0dev +# options: --amdgpu +surface matte %meta{string,help,"Lambertian diffuse material"} +param float Kd 1 %meta{string,help,"Diffuse scaling"} %meta{float,min,0} %meta{float,max,1} %read{1,1} %write{2147483647,-1} +param color Cs 1 1 1 %meta{string,help,"Base color"} %meta{float,min,0} %meta{float,max,1} %read{1,1} %write{2147483647,-1} +global normal N %read{0,0} %write{2147483647,-1} +global closure color Ci %read{2147483647,-1} %write{2,2} +temp closure color $tmp1 %read{2,2} %write{0,0} +const string $const1 "diffuse" %read{0,0} %write{2147483647,-1} +temp color $tmp2 %read{2,2} %write{1,1} +code ___main___ +# matte.osl:18 +# Ci = Kd * Cs * diffuse (N); + closure $tmp1 $const1 N %filename{"matte.osl"} %line{18} %argrw{"wrr"} + mul $tmp2 Kd Cs %argrw{"wrr"} + mul Ci $tmp1 $tmp2 %argrw{"wrr"} + end diff --git a/src/shaders/unnamed_group_1.ll b/src/shaders/unnamed_group_1.ll new file mode 100644 index 0000000000..1309d563e0 --- /dev/null +++ b/src/shaders/unnamed_group_1.ll @@ -0,0 +1,33 @@ +; Function Attrs: nocallback nofree nounwind willreturn memory(argmem: write) +declare void @llvm.memset.p0.i64(ptr nocapture writeonly, i8, i64, i1 immarg) #0 + +declare ptr @osl_allocate_closure_component(ptr, i32, i32) #1 + +define void @osl_init_group_unnamed_group_1(ptr %shaderglobals_ptr, ptr %groupdata_ptr, ptr %userdata_base_ptr, ptr %output_base_ptr, i32 %shadeindex, ptr %interactive_params_ptr) #1 !dbg !8 { +bb_osl_init_group_unnamed_group_1_0: + ret void, !dbg !12 +} + +define void @osl_layer_group_unnamed_group_1_name_matte_0(ptr %shaderglobals_ptr, ptr %groupdata_ptr, ptr %userdata_base_ptr, ptr %output_base_ptr, i32 %shadeindex, ptr %interactive_params_ptr) #1 !dbg !13 { +bb_osl_layer_group_unnamed_group_1_name_matte_0_1: + %0 = getelementptr %Groupdata, ptr %groupdata_ptr, i32 0, i32 0, i32 0, !dbg !14 + %1 = call ptr @osl_allocate_closure_component(ptr %shaderglobals_ptr, i32 3, i32 24), !dbg !15 + %2 = icmp ne ptr %1, null, !dbg !15 + br i1 %2, label %bb_non_null_closure_2, label %bb_3, !dbg !15 + +bb_non_null_closure_2: ; preds = %bb_osl_layer_group_unnamed_group_1_name_matte_0_1 + %3 = getelementptr %ClosureComponent, ptr %1, i32 0, i32 2, !dbg !15 + call void @llvm.memset.p0.i64(ptr align 4 %3, i8 0, i64 24, i1 false), !dbg !15 + %4 = getelementptr %ShaderGlobals, ptr %shaderglobals_ptr, i32 0, i32 3, !dbg !15 + call void @llvm.memcpy.p0.p0.i64(ptr align 4 %3, ptr align 4 %4, i64 12, i1 false), !dbg !15 + br label %bb_3, !dbg !15 + +bb_3: ; preds = %bb_non_null_closure_2, %bb_osl_layer_group_unnamed_group_1_name_matte_0_1 + %5 = getelementptr %ShaderGlobals, ptr %shaderglobals_ptr, i32 0, i32 23, !dbg !15 + store ptr %1, ptr %5, align 8, !dbg !15 + ret void, !dbg !15 +} + +; Function Attrs: nocallback nofree nounwind willreturn memory(argmem: readwrite) +declare void @llvm.memcpy.p0.p0.i64(ptr noalias nocapture writeonly, ptr noalias nocapture readonly, i64, i1 immarg) #2 + diff --git a/src/shaders/unnamed_group_1_O1.ll b/src/shaders/unnamed_group_1_O1.ll new file mode 100644 index 0000000000..2b20928d9c --- /dev/null +++ b/src/shaders/unnamed_group_1_O1.ll @@ -0,0 +1,33 @@ +; Function Attrs: mustprogress nocallback nofree nounwind willreturn memory(argmem: write) +declare void @llvm.memset.p0.i64(ptr nocapture writeonly, i8, i64, i1 immarg) #0 + +declare ptr @osl_allocate_closure_component(ptr, i32, i32) local_unnamed_addr #1 + +; Function Attrs: mustprogress nofree norecurse nosync nounwind willreturn memory(none) +define void @osl_init_group_unnamed_group_1(ptr nocapture readnone %shaderglobals_ptr, ptr nocapture readnone %groupdata_ptr, ptr nocapture readnone %userdata_base_ptr, ptr nocapture readnone %output_base_ptr, i32 %shadeindex, ptr nocapture readnone %interactive_params_ptr) local_unnamed_addr #2 !dbg !8 { +bb_osl_init_group_unnamed_group_1_0: + ret void, !dbg !12 +} + +define void @osl_layer_group_unnamed_group_1_name_matte_0(ptr %shaderglobals_ptr, ptr nocapture readnone %groupdata_ptr, ptr nocapture readnone %userdata_base_ptr, ptr nocapture readnone %output_base_ptr, i32 %shadeindex, ptr nocapture readnone %interactive_params_ptr) local_unnamed_addr #1 !dbg !13 { +bb_osl_layer_group_unnamed_group_1_name_matte_0_1: + %0 = tail call ptr @osl_allocate_closure_component(ptr %shaderglobals_ptr, i32 3, i32 24), !dbg !14 + %.not = icmp eq ptr %0, null, !dbg !14 + br i1 %.not, label %bb_3, label %bb_non_null_closure_2, !dbg !14 + +bb_non_null_closure_2: ; preds = %bb_osl_layer_group_unnamed_group_1_name_matte_0_1 + %1 = getelementptr %ClosureComponent, ptr %0, i64 0, i32 2, !dbg !14 + tail call void @llvm.memset.p0.i64(ptr noundef nonnull align 4 dereferenceable(24) %1, i8 0, i64 24, i1 false), !dbg !14 + %2 = getelementptr %ShaderGlobals, ptr %shaderglobals_ptr, i64 0, i32 3, !dbg !14 + tail call void @llvm.memcpy.p0.p0.i64(ptr noundef nonnull align 4 dereferenceable(12) %1, ptr noundef nonnull align 4 dereferenceable(12) %2, i64 12, i1 false), !dbg !14 + br label %bb_3, !dbg !14 + +bb_3: ; preds = %bb_non_null_closure_2, %bb_osl_layer_group_unnamed_group_1_name_matte_0_1 + %3 = getelementptr %ShaderGlobals, ptr %shaderglobals_ptr, i64 0, i32 23, !dbg !14 + store ptr %0, ptr %3, align 8, !dbg !14 + ret void, !dbg !14 +} + +; Function Attrs: mustprogress nocallback nofree nounwind willreturn memory(argmem: readwrite) +declare void @llvm.memcpy.p0.p0.i64(ptr noalias nocapture writeonly, ptr noalias nocapture readonly, i64, i1 immarg) #3 + diff --git a/src/testrender/CMakeLists.txt b/src/testrender/CMakeLists.txt index 6b01e7e2a4..430c2f56dd 100644 --- a/src/testrender/CMakeLists.txt +++ b/src/testrender/CMakeLists.txt @@ -8,7 +8,8 @@ set (testrender_srcs simpleraytracer.cpp scene.cpp bvh.cpp - testrender.cpp) + testrender.cpp + hip_raytracer.cpp) # New - Na find_package(Threads REQUIRED) @@ -75,16 +76,24 @@ if (CMAKE_COMPILER_IS_INTELCLANG) # model, although with slightly different numerical results. endif () +# --- INTEGRACJA AMD ROCm HIP --- +find_package(HIP REQUIRED) + +# Tworzenie celu wykonywalnego testrender (tylko RAZ w pliku) add_executable (testrender ${testrender_srcs}) target_include_directories (testrender BEFORE PRIVATE ${OpenImageIO_INCLUDES}) +# Dodajemy ścieżki do nagłówków HIP +target_include_directories (testrender PRIVATE ${HIP_INCLUDE_DIRS}) +# Linkowanie bibliotek rdzenia OSL oraz sterownika hosta HIP target_link_libraries (testrender PRIVATE oslexec oslquery oslcomp BSDL pugixml::pugixml - Threads::Threads) + Threads::Threads + hip::host) osl_optix_target (testrender) -install ( TARGETS testrender RUNTIME DESTINATION ${CMAKE_INSTALL_BINDIR} ) +install ( TARGETS testrender RUNTIME DESTINATION ${CMAKE_INSTALL_BINDIR} ) \ No newline at end of file diff --git a/src/testrender/cuda/rend_lib.cu b/src/testrender/cuda/rend_lib.cu index 681fae9579..8a1c3b1c84 100644 --- a/src/testrender/cuda/rend_lib.cu +++ b/src/testrender/cuda/rend_lib.cu @@ -540,5 +540,7 @@ osl_get_inverse_matrix(void* sg_, void* r, OSL::ustringhash_pod to_) return ok; } + + #undef MAT } diff --git a/src/testrender/gpu_raytracer.h b/src/testrender/gpu_raytracer.h new file mode 100644 index 0000000000..5985e8aac7 --- /dev/null +++ b/src/testrender/gpu_raytracer.h @@ -0,0 +1,22 @@ +// New - Na +#pragma once +#include +#include + +// Uniwersalna struktura do wczytywania bajtów +struct GPUShaderModuleDesc { + std::string architecture; + std::string format; + const void* data_ptr; + size_t data_size; +}; + +// Abstrakcyjna klasa bazowa +class GPURaytracer { +public: + virtual ~GPURaytracer() = default; + + virtual bool init() = 0; + virtual bool load_shader(const GPUShaderModuleDesc& desc) = 0; + virtual void render(int width, int height) = 0; +}; diff --git a/src/testrender/hip_raytracer.cpp b/src/testrender/hip_raytracer.cpp new file mode 100644 index 0000000000..08d6012693 --- /dev/null +++ b/src/testrender/hip_raytracer.cpp @@ -0,0 +1,308 @@ +// New - Ka +#include "hip_raytracer.h" +#include +#include +#include +#include +#include +#include +#include +#include +#include + +// Makro do wygodnego sprawdzania, czy funkcje HIP nie zwracają błędów +#define HIP_CHECK(command) \ +{ \ + hipError_t status = command; \ + if (status != hipSuccess) { \ + std::cerr << "HIP Error: " << hipGetErrorString(status) \ + << " w linii " << __LINE__ << " w pliku " << __FILE__ << std::endl; \ + return false; \ + } \ +} + +// NOWE makro dla funkcji zwracających void (render) +#define HIP_CHECK_VOID(command) \ +{ \ + hipError_t status = command; \ + if (status != hipSuccess) { \ + std::cerr << "HIP Error: " << hipGetErrorString(status) \ + << " w linii " << __LINE__ << " w pliku " << __FILE__ << std::endl; \ + return; \ + } \ +} + +// ------------------------------------------------------------------------- +// 1. Inicjalizacja środowiska i karty graficznej +// ------------------------------------------------------------------------- +bool HipRaytracer::init() { + return true; + std::cout << "\n=== [Inicjalizacja AMD HIP] ===\n"; + + int deviceCount = 0; + if (hipGetDeviceCount(&deviceCount) != hipSuccess || deviceCount == 0) { + std::cerr << "BŁĄD: Nie znaleziono żadnych urządzeń obsługujących HIP!" << std::endl; + return false; + } + + std::cout << "Znaleziono " << deviceCount << " urządzenie/a HIP.\n"; + + /*for (int i = 0; i < deviceCount; ++i) { + hipDeviceProp_t deviceProp; + HIP_CHECK(hipGetDeviceProperties(&deviceProp, i)); + + std::cout << "Urządzenie [" << i << "]: " << deviceProp.name << "\n"; + std::cout << " Architektura (GCN/RDNA): " << deviceProp.gcnArchName << "\n"; + std::cout << " Całkowita pamięć VRAM: " << deviceProp.totalGlobalMem / (1024 * 1024) << " MB\n"; + std::cout << " Max wątków na blok: " << deviceProp.maxThreadsPerBlock << "\n"; + }*/ + + // Wybieramy domyślną kartę (indeks 0) + HIP_CHECK(hipSetDevice(0)); + std::cout << "Pomyślnie podpięto do GPU 0.\n"; + std::cout << "===============================\n\n"; + + return true; +} + +// Destruktor (Zwalnianie pamięci z karty graficznej) + +HipRaytracer::~HipRaytracer() { + if (m_module) { + hipModuleUnload(m_module); + std::cout << "[HIP] Zwalnianie pamięci: Usunięto moduł z pamięci VRAM.\n"; + } +} +//NEW +// 2. Ładowanie skompilowanego modułu OSL (pliku binarnego HSACO dla AMD) +bool HipRaytracer::load_shader(const GPUShaderModuleDesc& desc) { + if (!desc.data_ptr || desc.data_size == 0) { + std::cerr << "[Błąd HIP] Pusty moduł przekazany do load_shader()!\n"; + return false; + } + + // W PODEJŚCIU AOT OTRZYMUJEMY JUŻ GOTOWY KOD HSACO + std::cout << "[HIP] Ładowanie gotowego modułu HSACO z pamięci (" << desc.data_size << " bajtów)...\n"; + + // Używamy hipModuleLoadData, żeby załadować kernel PROSTO Z PAMIĘCI RAM. + HIP_CHECK(hipModuleLoadData(&m_module, desc.data_ptr)); + + // Pobierz wskaźnik na kernel + std::cout << "[HIP] Szukam kernela: osl_kernel\n"; + hipError_t err = hipModuleGetFunction(&m_kernel, m_module, "osl_kernel"); + if (err != hipSuccess) { + std::cerr << "BŁĄD: Nie znaleziono kernela 'osl_kernel'. HIP: " << hipGetErrorString(err) << "\n"; + return false; + } + + std::cout << "[HIP] Załadowano kernel: osl_kernel\n"; + return true; +} + +// 3. Uruchomienie kernela renderującego na GPU + +void HipRaytracer::render(int width, int height) { + + std::cout << "[HIP] Rozpoczęcie renderowania. Rozdzielczość: " + + << width << "x" << height << "\n"; + + + + if (!m_kernel) { + + std::cerr << "[Błąd HIP] Kernel nie jest załadowany! Przerywam renderowanie.\n"; + + return; + + } + + + + // 1. Alokacja pamięci dla obrazka wyjściowego + + size_t buffer_size = width * height * 3 * sizeof(float); + + hipDeviceptr_t d_output; + + HIP_CHECK_VOID(hipMalloc((void**)&d_output, buffer_size)); + + + + // 2. Alokacja zerowych ShaderGlobals na GPU + + size_t sg_size = width * height * 256; + + void* d_shaderglobals = nullptr; + + HIP_CHECK_VOID(hipMalloc(&d_shaderglobals, sg_size)); + + HIP_CHECK_VOID(hipMemset(d_shaderglobals, 0, sg_size)); + + + + // 3. Skonfigurowanie bloków i siatki wątków + + dim3 blockSize(16, 16); + + dim3 gridSize((width + blockSize.x - 1) / blockSize.x, + + (height + blockSize.y - 1) / blockSize.y); + + + + // 4. PRZYGOTOWANIE ARGUMENTÓW ZGODNYCH Z NOWYM WRAPPEREM LLVM + + void* d_groupdata = nullptr; + + void* userdata_base_ptr = nullptr; + + int shadeindex = 0; + + void* interactive_params_ptr = nullptr; + + + + void* args[] = { + + &d_shaderglobals, // 1. ShaderGlobals* + + &d_groupdata, // 2. void* groupdata + + &userdata_base_ptr, // 3. void* userdata + + &d_output, // 4. void* output_base + + &shadeindex, // 5. int shadeindex + + &interactive_params_ptr, // 6. void* interactive_params + + &width // 7. int width + + }; + + + + // 5. Uruchomienie kernela + + HIP_CHECK_VOID(hipModuleLaunchKernel( + + m_kernel, + + gridSize.x, gridSize.y, gridSize.z, + + blockSize.x, blockSize.y, blockSize.z, + + 0, 0, // Pamięć współdzielona i strumień domyślny + + args, nullptr + + )); + + + + // 6. Czekamy aż karta skończy renderować + + HIP_CHECK_VOID(hipDeviceSynchronize()); + + std::cout << "[HIP] Kernel zakończył pracę.\n"; + + + + if (m_host_buffer != nullptr) { + + std::cout << "[HIP] Kopiowanie wyrenderowanego obrazu prosto do pamieci RAM hosta...\n"; + + HIP_CHECK_VOID(hipMemcpy(m_host_buffer, (void*)d_output, buffer_size, hipMemcpyDeviceToHost)); + +} else { + + std::cerr << "[HIP] OSTRZEZENIE: m_host_buffer jest pusty! Wyniki znikna w prozni.\n"; + +} + + + + // POBRANIE WYNIKÓW Z VRAM NA CPU + + std::cout << "[HIP] Kopiowanie wyrenderowanego obrazu do pamięci RAM...\n"; + + std::vector h_output(width * height * 3, 0.0f); // Wektor na dane na CPU + + + +HIP_CHECK_VOID(hipMemcpy(h_output.data(), (void*)d_output, buffer_size, hipMemcpyDeviceToHost)); + + + // Wypiszmy wartość pierwszego wyrenderowanego piksela w lewym górnym rogu + + std::cout << "\n=== [WYNIKI RENDEROWANIA] ===\n"; + + std::cout << "Piksel [0,0] RGB: (" + + << h_output[0] << ", " + + << h_output[1] << ", " + + << h_output[2] << ")\n"; + + + + // Wypiszmy wartość piksela ze środka ekranu + + int mid_x = width / 2; + + int mid_y = height / 2; + + int mid_idx = (mid_y * width + mid_x) * 3; + + + + std::cout << "Piksel [" << mid_x << "," << mid_y << "] RGB: (" + + << h_output[mid_idx] << ", " + + << h_output[mid_idx + 1] << ", " + + << h_output[mid_idx + 2] << ")\n"; + + std::cout << "=============================\n\n"; + + + + // --- WYMUSZONY ZAPIS OBRAZU NA DYSK PRZEZ OIIO --- + + std::string custom_output = "sukces_hip.png"; + + auto out_file = OIIO::ImageOutput::create(custom_output); + + if (out_file) { + + // Definiujemy strukturę obrazu: szerokość, wysokość, 3 kanały (RGB), typ float + + OIIO::ImageSpec spec(width, height, 3, OIIO::TypeDesc::FLOAT); + + out_file->open(custom_output, spec); + + out_file->write_image(OIIO::TypeDesc::FLOAT, h_output.data()); + + out_file->close(); + + std::cout << "[HIP SUCCESS] Obraz został pomyślnie zapisany w: " << custom_output << "\n"; + + } else { + + std::cerr << "[BŁĄD HIP] OIIO nie mogło utworzyć pliku: " << custom_output << "\n"; + + } + + + + // 7. Sprzątanie VRAM + + HIP_CHECK_VOID(hipFree((void*)d_output)); + + HIP_CHECK_VOID(hipFree(d_shaderglobals)); + +} + diff --git a/src/testrender/hip_raytracer.h b/src/testrender/hip_raytracer.h new file mode 100644 index 0000000000..c11bbb3e58 --- /dev/null +++ b/src/testrender/hip_raytracer.h @@ -0,0 +1,26 @@ +// New - Na +#pragma once +#include "gpu_raytracer.h" +#include // Nagłówek dla typów HIP + +class HipRaytracer : public GPURaytracer { +public: + HipRaytracer() = default; + + // Nadpisujemy destruktor, żeby zwalniał pamięć VRAM przy wyjściu z programu + ~HipRaytracer() override; + + bool init() override; + bool load_shader(const GPUShaderModuleDesc& desc) override; + void render(int width, int height) override; + + // --- NASZE BOCZNE DRZWI DLA PAMIĘCI --- + void set_host_buffer(float* buffer) { m_host_buffer = buffer; } + +private: + // Zmienna przechowująca skompilowany kod na karcie graficznej + hipModule_t m_module = nullptr; + hipFunction_t m_kernel = nullptr; + + float* m_host_buffer = nullptr; +}; diff --git a/src/testshade/CMakeLists.txt b/src/testshade/CMakeLists.txt index 12e1745a18..ea63d99afc 100644 --- a/src/testshade/CMakeLists.txt +++ b/src/testshade/CMakeLists.txt @@ -5,7 +5,8 @@ # The 'testshade' executable set ( testshade_srcs testshade.cpp - simplerend.cpp ) + simplerend.cpp + ../testrender/hip_raytracer.cpp) # New if (OSL_BUILD_BATCHED) list(APPEND testshade_srcs @@ -75,13 +76,20 @@ list(APPEND include_dirs ${OpenImageIO_INCLUDES}) EMBED_LLVM_BITCODE_IN_CPP ( "${rs_srcs}" "_host" "testshade_llvm_compiled_rs" testshade_srcs "-DOSL_HOST_RS_BITCODE=1" "${include_dirs}") +# --- INTEGRACJA AMD ROCm HIP --- +find_package(HIP REQUIRED) + +# Budowanie głównego programu wykonywalnego testshade add_executable ( testshade ${testshade_srcs} testshademain.cpp ) target_include_directories (testshade BEFORE PRIVATE ${OpenImageIO_INCLUDES}) +target_include_directories (testshade PRIVATE ${HIP_INCLUDE_DIRS}) target_link_libraries (testshade PRIVATE - oslexec oslquery oslcomp) + oslexec oslquery oslcomp + hip::host) + if (OSL_USE_OPTIX) add_dependencies(testshade testshade_ptx) endif () @@ -117,9 +125,13 @@ if (NOT CODECOV) -Wl,--version-script=${CMAKE_CURRENT_SOURCE_DIR}/testshade_symbols.map) endif () + # Linkowanie nagłówków i bibliotek HIP również do wersji bibliotecznej (DSO) + target_include_directories (libtestshade PRIVATE ${HIP_INCLUDE_DIRS}) target_link_libraries (libtestshade PRIVATE - oslexec oslquery oslcomp) + oslexec oslquery oslcomp + hip::host) + target_include_directories (libtestshade BEFORE PRIVATE ${OpenImageIO_INCLUDES}) set_target_properties (libtestshade PROPERTIES PREFIX "") @@ -134,4 +146,4 @@ if (NOT CODECOV) ${CMAKE_DL_LIBS} ) install (TARGETS testshade_dso RUNTIME DESTINATION ${CMAKE_INSTALL_BINDIR} ) osl_optix_target(libtestshade) -endif () +endif () \ No newline at end of file diff --git a/src/testshade/simplerend.cpp b/src/testshade/simplerend.cpp index 1563caa6d4..4265563b83 100644 --- a/src/testshade/simplerend.cpp +++ b/src/testshade/simplerend.cpp @@ -269,10 +269,12 @@ SimpleRenderer::SimpleRenderer() // Ensure destructor code gen happens in this .cpp SimpleRenderer::~SimpleRenderer() {} - -int +//NEW - Ka +int SimpleRenderer::supports(string_view feature) const { + if (feature == "AMDGPU" || feature == "HIP") //- dodanie obslugi AMD + return true; if (m_use_rs_bitcode && feature == "build_attribute_getter") return true; else if (m_use_rs_bitcode && feature == "build_interpolated_getter") diff --git a/src/testshade/testshade.cpp b/src/testshade/testshade.cpp index 9c1f999da8..22ccfe12ca 100644 --- a/src/testshade/testshade.cpp +++ b/src/testshade/testshade.cpp @@ -24,7 +24,7 @@ #include #include #include - +#include "../testrender/hip_raytracer.h" // New #include #include #include @@ -44,6 +44,8 @@ extern int testshade_llvm_compiled_rs_size; extern unsigned char testshade_llvm_compiled_rs_block[]; + + using namespace OSL; using OIIO::ParamValue; using OIIO::ParamValueList; @@ -129,6 +131,9 @@ static char* output_base_ptr = nullptr; static bool use_rs_bitcode = false; // use free function bitcode version of renderer services static int jbufferMB = 16; +// NEW +// Flagi dla trybu AOT GPU +static bool use_hip_runtime = false; // Testshade thread tracking and assignment. // Not recommended for production renderer but fine for testshade @@ -860,6 +865,10 @@ getargs(int argc, const char* argv[]) .help("Use free function bitcode Renderer services"); ap.arg("--jbufferMB %d:JBUFFER", &jbufferMB) .help("journal jbuffer size in MB"); + //NEW + // AMDGPU backend options: allow specifying the GPU target and saving the artifact + ap.arg("--hip-runtime", &use_hip_runtime) + .help("Uruchom renderer w trybie sprzętowym AMD HIP (wczytuje plik .hsaco o nazwie shadera)"); // clang-format on ap.parse_args(argc, argv); @@ -1983,6 +1992,7 @@ test_shade(int argc, const char* argv[]) // TextureSystem (note: passing nullptr just makes the ShadingSystem // make its own TS), and an error handler. shadingsys = new ShadingSystem(rend.get(), texturesys, &rend->errhandler()); + rend->init_shadingsys(shadingsys); // Register the layout of all closures known to this renderer @@ -2269,6 +2279,67 @@ test_shade(int argc, const char* argv[]) double runtime = timer.lap(); +// NEW + // AMDGPU AOT RUNTIME + if (use_hip_runtime) { + std::cout << "[Testshade] Tryb AOT: Przygotowanie renderera HIP...\n"; + + if (shadernames.empty()) { + std::cerr << "FATAL: Brak zadeklarowanego shadera do wczytania!\n"; + return EXIT_FAILURE; + } + + // 1. Inicjalizacja renderera + std::unique_ptr gpu_renderer = std::make_unique(); + if (!gpu_renderer->init()) { + std::cerr << "FATAL: Błąd inicjalizacji HIP!\n"; + return EXIT_FAILURE; + } + + // 2. Pobranie nazwy shadera (np. "moj_material") i doklejenie ".hsaco" + // Używamy OIIO, by pozbyć się rozszerzenia .oso (jeśli użytkownik je podał) + std::string base_name = OIIO::Filesystem::replace_extension(shadernames[0], ""); + std::string hsaco_filename = base_name + ".hsaco"; + + std::cout << "[Testshade] Szukam gotowego pliku jądra GPU: " << hsaco_filename << "\n"; + + std::vector hsaco_buffer; + size_t file_size = OIIO::Filesystem::file_size(hsaco_filename); + if (file_size > 0) { + hsaco_buffer.resize(file_size); + if (OIIO::Filesystem::read_bytes(hsaco_filename, (void*)hsaco_buffer.data(), file_size) != file_size) { + std::cerr << "FATAL: Nie udało się wczytać pliku .hsaco!\n"; + return EXIT_FAILURE; + } + } + + // 3. Pakujemy bajty i wysyłamy prosto do HIP-a + GPUShaderModuleDesc desc; + desc.architecture = "auto"; // W trybie AOT (hipModuleLoadData) to nie ma znaczenia + desc.format = "hsaco"; + desc.data_ptr = hsaco_buffer.data(); + desc.data_size = hsaco_buffer.size(); + + if (!gpu_renderer->load_shader(desc)) { + std::cerr << "FATAL: Nie udało się załadować pliku " << hsaco_filename << " do HIP-a!\n"; + return EXIT_FAILURE; + } + + // 4. Łączymy GPU z pamięcią RAM dla obrazu wyjściowego + if (OIIO::ImageBuf* main_img = rend->outputbuf(0)) { + auto* hip_renderer = static_cast(gpu_renderer.get()); + hip_renderer->set_host_buffer((float*)main_img->localpixels()); + } else { + std::cerr << "[HIP] OSTRZEŻENIE: Nie udało się pobrać głównego bufora obrazu!\n"; + } + + // 5. Renderujemy na sprzęcie + gpu_renderer->render(xres, yres); + } + + // This awkward condition preserves an output oddity from long ago, + // eliminating the need to update hundreds of ref outputs. + if (outputfiles.size() == 1 && outputfiles[0] == "null") // This awkward condition preserves an output oddity from long ago, // eliminating the need to update hundreds of ref outputs. if (outputfiles.size() == 1 && outputfiles[0] == "null") diff --git a/testsuite/render-cornell/Untitled.ipynb b/testsuite/render-cornell/Untitled.ipynb new file mode 100644 index 0000000000..363fcab7ed --- /dev/null +++ b/testsuite/render-cornell/Untitled.ipynb @@ -0,0 +1,6 @@ +{ + "cells": [], + "metadata": {}, + "nbformat": 4, + "nbformat_minor": 5 +}