diff --git a/README.md b/README.md index 62e8fbf..f46ed6c 100644 --- a/README.md +++ b/README.md @@ -28,7 +28,7 @@ cmake --build . --target mlir-doc ## LLVM Build Instructions -The repository depends on a build of LLVM including MLIR. The Open Earth Compiler build has been tested with LLVM commit 383b6501ffed using the following configuration: +The repository depends on a build of LLVM including MLIR. The Open Earth Compiler build has been tested with LLVM commit cc4244d55f98 using the following configuration: ``` cmake -G Ninja ../llvm -DLLVM_BUILD_EXAMPLES=OFF -DLLVM_TARGETS_TO_BUILD="host;NVPTX;AMDGPU" -DCMAKE_INSTALL_PREFIX= -DLLVM_ENABLE_PROJECTS='mlir;lld' -DLLVM_OPTIMIZED_TABLEGEN=ON -DLLVM_ENABLE_OCAMLDOC=OFF -DLLVM_ENABLE_BINDINGS=OFF -DLLVM_INSTALL_UTILS=ON -DCMAKE_LINKER= -DLLVM_PARALLEL_LINK_JOBS=2 ``` diff --git a/patches/runtime.patch b/patches/runtime.patch index 8568b6d..d941d54 100644 --- a/patches/runtime.patch +++ b/patches/runtime.patch @@ -376,37 +376,3 @@ index 162c2f4e838a..e665c825f037 100644 // CHECK: %[[src:.*]] = llvm.bitcast // CHECK: %[[dst:.*]] = llvm.bitcast // CHECK: llvm.call @mgpuMemcpy(%[[dst]], %[[src]], %[[size_bytes]], %[[t0]]) -diff --git a/mlir/tools/mlir-rocm-runner/rocm-runtime-wrappers.cpp b/mlir/tools/mlir-rocm-runner/rocm-runtime-wrappers.cpp -index 4f62f204f4a8..b8f55ddd8e83 100644 ---- a/mlir/tools/mlir-rocm-runner/rocm-runtime-wrappers.cpp -+++ b/mlir/tools/mlir-rocm-runner/rocm-runtime-wrappers.cpp -@@ -36,7 +36,7 @@ static auto InitializeCtx = [] { - HIP_REPORT_IF_ERROR(hipInit(/*flags=*/0)); - hipDevice_t device; - HIP_REPORT_IF_ERROR(hipDeviceGet(&device, /*ordinal=*/0)); -- hipContext_t context; -+ hipCtx_t context; - HIP_REPORT_IF_ERROR(hipCtxCreate(&context, /*flags=*/0, device)); - return 0; - }(); -@@ -110,17 +110,17 @@ extern "C" void mgpuEventRecord(hipEvent_t event, hipStream_t stream) { - - extern "C" void *mgpuMemAlloc(uint64_t sizeBytes, hipStream_t /*stream*/) { - void *ptr; -- HIP_REPORT_IF_ERROR(hipMemAlloc(&ptr, sizeBytes)); -+ HIP_REPORT_IF_ERROR(hipMalloc(&ptr, sizeBytes)); - return ptr; - } - - extern "C" void mgpuMemFree(void *ptr, hipStream_t /*stream*/) { -- HIP_REPORT_IF_ERROR(hipMemFree(ptr)); -+ HIP_REPORT_IF_ERROR(hipFree(ptr)); - } - - extern "C" void mgpuMemcpy(void *dst, void *src, uint64_t sizeBytes, - hipStream_t stream) { -- HIP_REPORT_IF_ERROR(hipMemcpyAsync(dst, src, sizeBytes, stream)); -+ HIP_REPORT_IF_ERROR(hipMemcpyAsync(dst, src, sizeBytes, hipMemcpyDefault, stream)); - } - - /// Helper functions for writing mlir example code