From f2b92ba945a2ae55f11e7d412ea5b862e68b8dd9 Mon Sep 17 00:00:00 2001 From: Chao Liu Date: Wed, 12 Jun 2019 09:59:28 -0500 Subject: [PATCH] fixed build issue --- driver/CMakeLists.txt | 2 +- ...icit_gemm_v4_nchw_kcyx_nkhw_lds_double_buffer.hpp | 12 ++---------- 2 files changed, 3 insertions(+), 11 deletions(-) diff --git a/driver/CMakeLists.txt b/driver/CMakeLists.txt index c5c2a0fd72..4a7a81a36a 100644 --- a/driver/CMakeLists.txt +++ b/driver/CMakeLists.txt @@ -1,5 +1,5 @@ if(DEVICE_BACKEND STREQUAL "HIP") - set(DRIVER_SOURCE driver.hip.cpp) + set(DRIVER_SOURCE driver.cpp) elseif(DEVICE_BACKEND STREQUAL "CUDA") set(DRIVER_SOURCE driver.cu) endif() diff --git a/include/composable_kernel/kernel_algorithm/gridwise_convolution_implicit_gemm_v4_nchw_kcyx_nkhw_lds_double_buffer.hpp b/include/composable_kernel/kernel_algorithm/gridwise_convolution_implicit_gemm_v4_nchw_kcyx_nkhw_lds_double_buffer.hpp index 786d83db97..d6badfda8b 100644 --- a/include/composable_kernel/kernel_algorithm/gridwise_convolution_implicit_gemm_v4_nchw_kcyx_nkhw_lds_double_buffer.hpp +++ b/include/composable_kernel/kernel_algorithm/gridwise_convolution_implicit_gemm_v4_nchw_kcyx_nkhw_lds_double_buffer.hpp @@ -264,16 +264,8 @@ struct GridwiseConvolutionImplicitGemm_v4_nchw_kcyx_nkhw_lds_double_buffer // LDS double buffer: preload data into LDS { - Float p_in_register_clipboard[blockwise_in_copy.GetRegisterClipboardSize()]; - Float p_wei_register_clipboard[blockwise_wei_copy.GetRegisterClipboardSize()]; - - blockwise_in_copy.RunLoadRegisterClipboard(p_in_global, p_in_register_clipboard); - blockwise_wei_copy.RunLoadRegisterClipboard(p_wei_block_on_global, - p_wei_register_clipboard); - - blockwise_in_copy.RunStoreRegisterClipboard(p_in_register_clipboard, p_in_block_double); - blockwise_wei_copy.RunStoreRegisterClipboard(p_wei_register_clipboard, - p_wei_block_double); + blockwise_in_copy.Run(p_in_global, p_in_block_double); + blockwise_wei_copy.Run(p_wei_global, p_wei_block_double); } // LDS double buffer: main body