forked from ggml-org/llama.cpp
-
Notifications
You must be signed in to change notification settings - Fork 1
Commit
This commit does not belong to any branch on this repository, and may belong to a fork outside of the repository.
- Loading branch information
Showing
11 changed files
with
384 additions
and
44 deletions.
There are no files selected for viewing
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -1,19 +1,46 @@ | ||
| if(NOT SOC_VERSION) | ||
| set(SOC_VERSION "ascend910b3") | ||
| if (NOT SOC_TYPE) | ||
| set (SOC_TYPE "Ascend910B3") | ||
| endif() | ||
| set(ASCEND_CANN_PACKAGE_PATH ${CANN_INSTALL_DIR}) | ||
| set(RUN_MODE "npu" CACHE STRING "run mode: npu/sim/cpu") | ||
|
|
||
| if(EXISTS ${ASCEND_CANN_PACKAGE_PATH}/compiler/tikcpp/ascendc_kernel_cmake) | ||
| set(ASCENDC_CMAKE_DIR ${ASCEND_CANN_PACKAGE_PATH}/compiler/tikcpp/ascendc_kernel_cmake) | ||
| elseif(EXISTS ${ASCEND_CANN_PACKAGE_PATH}/ascendc_devkit/tikcpp/samples/cmake) | ||
| set(ASCENDC_CMAKE_DIR ${ASCEND_CANN_PACKAGE_PATH}/ascendc_devkit/tikcpp/samples/cmake) | ||
| file(GLOB SRC_FILES | ||
| dequantize_q4_0.cpp | ||
| quantize_q4_0.cpp | ||
| ) | ||
|
|
||
| string(TOLOWER "${CMAKE_BUILD_TYPE}" lowercase_CMAKE_BUILD_TYPE) | ||
| if(${lowercase_CMAKE_BUILD_TYPE} STREQUAL "debug") | ||
| if (NOT DEFINED ENV{CMAKE_PREFIX_PATH}) | ||
| set(CMAKE_PREFIX_PATH ${CANN_INSTALL_DIR}/tools/tikicpulib/lib/cmake) | ||
| endif() | ||
|
|
||
| find_package(tikicpulib REQUIRED) | ||
| add_library(cann_kernels ${SRC_FILES} ascendc_kernels.cpp) | ||
| target_link_libraries(cann_kernels PRIVATE | ||
| ascendcl | ||
| tikicpulib::ascend910B1 | ||
| ) | ||
|
|
||
| target_compile_features(cann_kernels PRIVATE cxx_std_17) | ||
|
|
||
| else() | ||
| message(FATAL_ERROR "ascendc_kernel_cmake does not exist, please check whether the compiler package is installed.") | ||
| endif() | ||
| string(TOLOWER ${SOC_TYPE} SOC_VERSION) | ||
| set(ASCEND_CANN_PACKAGE_PATH ${CANN_INSTALL_DIR}) | ||
| set(RUN_MODE "npu" CACHE STRING "run mode: npu/sim") | ||
|
|
||
| include(${ASCENDC_CMAKE_DIR}/ascendc.cmake) | ||
| if(EXISTS ${ASCEND_CANN_PACKAGE_PATH}/compiler/tikcpp/ascendc_kernel_cmake) | ||
| set(ASCENDC_CMAKE_DIR ${ASCEND_CANN_PACKAGE_PATH}/compiler/tikcpp/ascendc_kernel_cmake) | ||
| elseif(EXISTS ${ASCEND_CANN_PACKAGE_PATH}/ascendc_devkit/tikcpp/samples/cmake) | ||
| set(ASCENDC_CMAKE_DIR ${ASCEND_CANN_PACKAGE_PATH}/ascendc_devkit/tikcpp/samples/cmake) | ||
| else() | ||
| message(FATAL_ERROR "ascendc_kernel_cmake does not exist, please check whether the compiler package is installed.") | ||
| endif() | ||
| include(${ASCENDC_CMAKE_DIR}/ascendc.cmake) | ||
|
|
||
| ascendc_library(ascendc_kernels STATIC | ||
| threshold_opencv_kernel.cpp | ||
| ) | ||
| ascendc_library(ascendc_kernels STATIC | ||
| ${SRC_FILES} | ||
| ) | ||
| #ascendc_compile_definitions(ascendc_kernels PRIVATE -DASCENDC_DUMP) | ||
|
|
||
| add_library(cann_kernels STATIC ascendc_kernels.cpp) | ||
| target_link_libraries(cann_kernels PUBLIC ascendc_kernels) | ||
| endif() |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,56 @@ | ||
| #include "ascendc_kernels.h" | ||
|
|
||
| #ifdef __CCE_KT_TEST__ | ||
| #include "tikicpulib.h" | ||
| #else | ||
| #include "aclrtlaunch_ascendc_dequantize_q4_0.h" | ||
| #include "aclrtlaunch_ascendc_quantize_q4_0.h" | ||
| #endif | ||
|
|
||
|
|
||
| #ifdef __CCE_KT_TEST__ | ||
| #include <acl/acl.h> | ||
|
|
||
| uint8_t* to_gm(uint8_t* ptr, size_t size) { | ||
| uint8_t* gm = (uint8_t*)AscendC::GmAlloc(size); | ||
| aclrtMemcpy(gm, size, ptr, size, ACL_MEMCPY_DEVICE_TO_HOST); | ||
| return gm; | ||
| } | ||
|
|
||
| void free_gm(uint8_t* ptr) { | ||
| aclrtFree(ptr); | ||
| } | ||
|
|
||
| extern "C" __global__ __aicore__ void ascendc_dequantize_q4_0(GM_ADDR x, GM_ADDR y, GM_ADDR size); | ||
| extern "C" __global__ __aicore__ void ascendc_quantize_q4_0(GM_ADDR x, GM_ADDR y, GM_ADDR size); | ||
| #endif | ||
|
|
||
| void cann_dequantize_q4_0(uint32_t block_dim, void* stream, uint8_t* x, uint8_t* y, uint8_t* size) { | ||
| #ifdef __CCE_KT_TEST__ | ||
| uint8_t* size_host = to_gm(size, sizeof(size_t)); | ||
| uint8_t* x_host = to_gm(x, *((size_t*)size_host)); | ||
| uint8_t* y_host = to_gm(y, *((size_t*)size_host)); | ||
| AscendC::SetKernelMode(KernelMode::AIV_MODE); | ||
| ICPU_RUN_KF(ascendc_dequantize_q4_0, 1, x_host, y_host, size_host); | ||
| free_gm(size_host); | ||
| free_gm(x_host); | ||
| free_gm(y_host); | ||
| #else | ||
| aclrtlaunch_ascendc_dequantize_q4_0(block_dim, stream, x, y, size); | ||
| #endif | ||
| } | ||
|
|
||
| void cann_quantize_q4_0(uint32_t block_dim, void* stream, uint8_t* x, uint8_t* y, uint8_t* size) { | ||
| #ifdef __CCE_KT_TEST__ | ||
| uint8_t* size_host = to_gm(size, sizeof(size_t)); | ||
| uint8_t* x_host = to_gm(x, *((size_t*)size_host)); | ||
| uint8_t* y_host = to_gm(y, *((size_t*)size_host)); | ||
| AscendC::SetKernelMode(KernelMode::AIV_MODE); | ||
| ICPU_RUN_KF(ascendc_quantize_q4_0, 1, x_host, y_host, size_host); | ||
| free_gm(size_host); | ||
| free_gm(x_host); | ||
| free_gm(y_host); | ||
| #else | ||
| aclrtlaunch_ascendc_quantize_q4_0(block_dim, stream, x, y, size); | ||
| #endif | ||
| } |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -1,4 +1,10 @@ | ||
| #ifndef ASCENDC_KERNELS_H | ||
| #define ASCENDC_KERNELS_H | ||
|
|
||
|
|
||
| #include <stdint.h> | ||
|
|
||
| void cann_dequantize_q4_0(uint32_t block_dim, void* stream, uint8_t* x, uint8_t* y, uint8_t* size); | ||
| void cann_quantize_q4_0(uint32_t block_dim, void* stream, uint8_t* x, uint8_t* y, uint8_t* size); | ||
|
|
||
| #endif //ASCENDC_KERNELS_H |
Oops, something went wrong.