diff --git a/CMakeLists.txt b/CMakeLists.txt index 65f8347a4..b264be8b6 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -3,7 +3,7 @@ # For GCC: `cmake -B build . && cmake --build build` # For MSVC: `cmake -B build . && cmake --build build --config Release` # You can also use the following options and variables -# - COMPUTE_BACKEND: Set to `cpu`, `cuda`, or `mps` to select the backend +# - COMPUTE_BACKEND: Set to `cpu`, `cuda`, `hip`, or `xpu` to select the backend # - CUDA_VERSION: The expected CUDA version, for sanity checking. The actual version # is whatever CMake finds on your path. # - COMPUTE_CAPABILITY: Which GPU Arch/Compute codes to provide to NVCC. @@ -62,14 +62,12 @@ set(CMAKE_CXX_EXTENSIONS OFF) # Define included source files set(CPP_FILES csrc/cpu_ops.cpp csrc/pythonInterface.cpp) set(GPU_FILES csrc/ops.cu csrc/kernels.cu) -set(MPS_FILES csrc/mps_ops.mm) -set(METAL_FILES csrc/mps_kernels.metal) set(XPU_FILES csrc/xpu_ops.cpp csrc/xpu_kernels.cpp) # C++ sources are always included list(APPEND SRC_FILES ${CPP_FILES}) -set(COMPUTE_BACKEND "cpu" CACHE STRING "The compute backend to use (cpu, cuda, hip, mps, xpu)") -set_property(CACHE COMPUTE_BACKEND PROPERTY STRINGS cpu cuda hip mps xpu) +set(COMPUTE_BACKEND "cpu" CACHE STRING "The compute backend to use (cpu, cuda, hip, xpu)") +set_property(CACHE COMPUTE_BACKEND PROPERTY STRINGS cpu cuda hip xpu) option(PTXAS_VERBOSE "Pass through -v flag to PTX Assembler" OFF) if(APPLE) @@ -86,33 +84,22 @@ if(${COMPUTE_BACKEND} STREQUAL "cuda") endif() set(BUILD_CUDA ON) set(BUILD_HIP OFF) - set(BUILD_MPS OFF) elseif(${COMPUTE_BACKEND} STREQUAL "hip") if(APPLE) message(FATAL_ERROR "HIP is not supported on macOS" ) endif() set(BUILD_CUDA OFF) set(BUILD_HIP ON) - set(BUILD_MPS OFF) -elseif(${COMPUTE_BACKEND} STREQUAL "mps") - if(NOT APPLE) - message(FATAL_ERROR "MPS is only supported on macOS" ) - endif() - set(BUILD_CUDA OFF) - set(BUILD_HIP OFF) - set(BUILD_MPS ON) elseif(${COMPUTE_BACKEND} STREQUAL "xpu") if(APPLE) message(FATAL_ERROR "XPU is not supported on macOS" ) endif() set(BUILD_CUDA OFF) set(BUILD_HIP OFF) - set(BUILD_MPS OFF) set(BUILD_XPU ON) else() set(BUILD_CUDA OFF) set(BUILD_HIP OFF) - set(BUILD_MPS OFF) set(BUILD_XPU OFF) set(BUILD_CPU ON) endif() @@ -312,25 +299,6 @@ elseif(BUILD_HIP) add_compile_definitions(__HIP_PLATFORM_AMD__) add_compile_definitions(__HIP_PLATFORM_HCC__) add_compile_definitions(BUILD_HIP) -elseif(BUILD_MPS) - if(NOT APPLE) - message(FATAL_ERROR "MPS is only supported on macOS" ) - endif() - - enable_language(OBJCXX) - - list(APPEND SRC_FILES ${MPS_FILES}) - - string(APPEND BNB_OUTPUT_NAME "_mps") - add_compile_definitions(BUILD_MPS) - file(MAKE_DIRECTORY "build") - add_custom_command(OUTPUT "bitsandbytes/bitsandbytes.metallib" - COMMAND xcrun metal -c -o "build/bitsandbytes.air" ${METAL_FILES} - COMMAND xcrun metallib "build/bitsandbytes.air" -o "bitsandbytes/bitsandbytes.metallib" - DEPENDS "${METAL_FILES}" - COMMENT "Compiling Metal kernels" - VERBATIM) - add_custom_target(metallib DEPENDS "bitsandbytes/bitsandbytes.metallib") elseif(BUILD_XPU) list(APPEND SRC_FILES ${XPU_FILES}) string(APPEND BNB_OUTPUT_NAME "_xpu") @@ -469,10 +437,6 @@ if(BUILD_HIP) target_link_libraries(bitsandbytes PUBLIC roc::hipblaslt) endif() endif() -if(BUILD_MPS) - add_dependencies(bitsandbytes metallib) - target_link_libraries(bitsandbytes objc "-framework Foundation" "-framework Metal" "-framework MetalPerformanceShaders" "-framework MetalPerformanceShadersGraph") -endif() if(BUILD_XPU) set(SYCL_LINK_FLAGS "-fsycl;--offload-compress;-fsycl-targets=spir64_gen,spir64;-Xs;-device pvc,xe-lpg,ats-m150 -options ' -cl-intel-enable-auto-large-GRF-mode -cl-poison-unsupported-fp64-kernels -cl-intel-greater-than-4GB-buffer-required'") set(SYCL_COMPILE_FLAGS "-fsycl;-fhonor-nans;-fhonor-infinities;-fno-associative-math;-fno-approx-func;-fno-sycl-instrument-device-code;--offload-compress;-fsycl-targets=spir64_gen,spir64;") diff --git a/csrc/mps_kernels.metal b/csrc/mps_kernels.metal deleted file mode 100644 index 63b3bf78c..000000000 --- a/csrc/mps_kernels.metal +++ /dev/null @@ -1,117 +0,0 @@ -#include -using namespace metal; - -#define HLF_MAX 65504 -#define TH 1024 -#define NUM 4 -#define NUM_BLOCK 4096 - -template -static unsigned char quantize_scalar( - float rand, - device float* code, - float x) -{ - int pivot = 127; - int upper_pivot = 255; - int lower_pivot = 0; - - float lower = -1.0f; - float upper = 1.0f; - - float val = code[pivot]; - // i>>=1 = {32, 16, 8, 4, 2, 1} - for(int i = 64; i > 0; i>>=1) - { - if(x > val) - { - lower_pivot = pivot; - lower = val; - pivot+=i; - } - else - { - upper_pivot = pivot; - upper = val; - pivot-=i; - } - val = code[pivot]; - } - - if(upper_pivot == 255) - upper = code[upper_pivot]; - if(lower_pivot == 0) - lower = code[lower_pivot]; - - if(!STOCHASTIC) - { - if(x > val) - { - float midpoint = (upper+val)*0.5f; - if(x > midpoint) - { - return upper_pivot; - } - else - return pivot; - } - else - { - float midpoint = (lower+val)*0.5f; - if(x < midpoint) - return lower_pivot; - else - return pivot; - } - } - else - { - if(x > val) - { - float dist_to_upper = fabs(upper-x); - float dist_full = upper-val; - if(rand >= dist_to_upper/dist_full) return upper_pivot; - else return pivot; - } - else - { - float dist_to_lower = fabs(lower-x); - float dist_full = val-lower; - if(rand >= dist_to_lower/dist_full) return lower_pivot; - else return pivot; - } - } -} - -kernel void quantize(device float* code [[buffer(0)]], - device float* A [[buffer(1)]], - device uchar* out [[buffer(2)]], - constant uint& n [[buffer(3)]], - uint id [[thread_position_in_grid]]) { - const uint n_full = (NUM_BLOCK * (n / NUM_BLOCK)) + (n % NUM_BLOCK == 0 ? 0 : NUM_BLOCK); - uint valid_items = (id / NUM_BLOCK + 1 == (n + NUM_BLOCK - 1) / NUM_BLOCK) ? n - (id / NUM_BLOCK * NUM_BLOCK) : NUM_BLOCK; - const uint base_idx = (id / NUM_BLOCK * NUM_BLOCK); - - float vals[NUM]; - uchar qvals[NUM]; - - for (uint i = base_idx; i < n_full; i += ((n + NUM_BLOCK - 1) / NUM_BLOCK) * NUM_BLOCK) { - valid_items = n - i > NUM_BLOCK ? NUM_BLOCK : n - i; - - threadgroup_barrier(mem_flags::mem_threadgroup); - - for (uint j = 0; j < valid_items; j++) { - vals[j] = A[i + j]; - } - - for (uint j = 0; j < valid_items; j++) { - qvals[j] = quantize_scalar(0.0f, code, vals[j]); - } - - threadgroup_barrier(mem_flags::mem_threadgroup); - - for (uint j = 0; j < valid_items; j++) { - out[i + j] = qvals[j]; - } - } -} diff --git a/csrc/mps_ops.mm b/csrc/mps_ops.mm deleted file mode 100644 index 85ed1b1e4..000000000 --- a/csrc/mps_ops.mm +++ /dev/null @@ -1,62 +0,0 @@ -#import - -#define HLF_MAX 65504 -#define TH 1024 -#define NUM 4 -#define NUM_BLOCK 4096 - -static inline MPSGraph* get_graph() { - static MPSGraph* cur = nil; - if (!cur) { - cur = [[MPSGraph alloc] init]; - } - return cur; -} - -static inline id get_device() { - NSError* error = nil; - static id device = nil; - if (!device) { - device = MTLCreateSystemDefaultDevice(); - } - if (!device) { - NSLog(@"Failed to get MPS device"); - abort(); - } - return device; -} - -static inline id get_library() { - NSError* error = nil; - static id library = nil; - if (!library) { - library = [get_device() newLibraryWithURL:[NSURL fileURLWithPath:@"bitsandbytes.metallib"] error:&error]; - } - if (!library) { - NSLog(@"Failed to load bitsandbytes.metallib"); - abort(); - } - return library; -} - -/*MPSGraphTensor* dequantize_mps(MPSGraphTensor* code, MPSGraphTensor* A, int n) -{ - id out = [get_graph() dequantizeTensor:(MPSGraphTensor*)A scaleTensor:(MPSGraphTensor*)code zeroPoint:0.0 -dataType:MPSDataTypeInt8 axis:0 name:@"out"]; return out; -}*/ - -// MPSGraph function for quantize -extern "C" MPSGraphTensor* quantize_mps(MPSGraph* graph, MPSGraphTensor* code, MPSGraphTensor* A, int n) { - id device = get_device(); - id library = get_library(); - static id kernel = nil; - if (!kernel) { - kernel = [library newFunctionWithName:@"quantize"]; - if (!kernel) { - NSLog(@"Failed to load bitsandbytes.metallib"); - abort(); - } - } - NSLog(@"Not implemented"); - return nil; -} diff --git a/csrc/pythonInterface.cpp b/csrc/pythonInterface.cpp index 2f9f0ff69..24431603b 100644 --- a/csrc/pythonInterface.cpp +++ b/csrc/pythonInterface.cpp @@ -10,9 +10,6 @@ #if BUILD_HIP #include #endif -#if BUILD_MPS -// #include -#endif #if BUILD_XPU #include #endif