Repository: ollama/ollama
Pull request: llama: fix compiler flag differences

Adjust the flags for the new Go server to more closely match the generate flow

Still more refinement to do, but I think this gets us closer.  Comparisons on an Ryzen 9 7950X system (windows) (linux)

This PR:  llama3.2 tps
- cpu: 9.17(w) 4.07(l)
- cpu_avx: 14.19(w) 14.78(l)
- cpu_avx2: 14.3(w) 15.3(l)

C++ server (0.3.12 release)
- cpu: 5.9(w) 7.34(l)
- cpu_avx: 17.5(w) 21.47(l)
- cpu_avx2: 19 tps(w) 22.63(l)

Changed files:
filename: llama/llama.go
patch: @@ -1,8 +1,8 @@
 package llama
 
 /*
-#cgo CFLAGS: -std=c11 -DNDEBUG -DLOG_DISABLE_LOGS -DGGML_USE_LLAMAFILE
-#cgo CXXFLAGS: -std=c++11 -DNDEBUG -DLOG_DISABLE_LOGS -DGGML_USE_LLAMAFILE
+#cgo CFLAGS: -O2 -std=c11 -DGGML_BUILD=1 -DNDEBUG -DLOG_DISABLE_LOGS -DGGML_USE_LLAMAFILE
+#cgo CXXFLAGS: -O2 -std=c++11 -DGGML_BUILD=1 -DNDEBUG -DLOG_DISABLE_LOGS -DGGML_USE_LLAMAFILE
 #cgo darwin,arm64 CFLAGS: -DGGML_USE_METAL -DGGML_USE_ACCELERATE -DGGML_METAL_EMBED_LIBRARY -DACCELERATE_NEW_LAPACK -DACCELERATE_LAPACK_ILP64 -DGGML_USE_BLAS
 #cgo darwin,arm64 CXXFLAGS: -DGGML_USE_METAL -DGGML_USE_ACCELERATE -DGGML_METAL_EMBED_LIBRARY -DACCELERATE_NEW_LAPACK -DACCELERATE_LAPACK_ILP64 -DGGML_USE_BLAS
 #cgo darwin,arm64 LDFLAGS: -framework Foundation -framework Metal -framework MetalKit -framework Accelerate
@@ -22,12 +22,12 @@ package llama
 #cgo windows,amd64 LDFLAGS: -L${SRCDIR}/build/Windows/amd64
 #cgo avx CFLAGS: -mavx
 #cgo avx CXXFLAGS: -mavx
-#cgo avx2 CFLAGS: -mavx2 -mfma
-#cgo avx2 CXXFLAGS: -mavx2 -mfma
-#cgo cuda CFLAGS: -DGGML_USE_CUDA -DGGML_CUDA_DMMV_X=32 -DGGML_CUDA_PEER_MAX_BATCH_SIZE=128 -DGGML_CUDA_MMV_Y=1 -DGGML_BUILD=1
-#cgo cuda CXXFLAGS: -DGGML_USE_CUDA -DGGML_CUDA_DMMV_X=32 -DGGML_CUDA_PEER_MAX_BATCH_SIZE=128 -DGGML_CUDA_MMV_Y=1 -DGGML_BUILD=1
-#cgo rocm CFLAGS: -DGGML_USE_CUDA -DGGML_USE_HIPBLAS -DGGML_CUDA_DMMV_X=32 -DGGML_CUDA_PEER_MAX_BATCH_SIZE=128 -DGGML_CUDA_MMV_Y=1 -DGGML_BUILD=1
-#cgo rocm CXXFLAGS: -DGGML_USE_CUDA -DGGML_USE_HIPBLAS -DGGML_CUDA_DMMV_X=32 -DGGML_CUDA_PEER_MAX_BATCH_SIZE=128 -DGGML_CUDA_MMV_Y=1 -DGGML_BUILD=1
+#cgo avx2 CFLAGS: -mavx2 -mfma -mf16c
+#cgo avx2 CXXFLAGS: -mavx2 -mfma -mf16c
+#cgo cuda CFLAGS: -fPIE -DGGML_USE_CUDA -DGGML_CUDA_DMMV_X=32 -DGGML_CUDA_PEER_MAX_BATCH_SIZE=128 -DGGML_CUDA_MMV_Y=1
+#cgo cuda CXXFLAGS: -fPIE -DGGML_USE_CUDA -DGGML_CUDA_DMMV_X=32 -DGGML_CUDA_PEER_MAX_BATCH_SIZE=128 -DGGML_CUDA_MMV_Y=1
+#cgo rocm CFLAGS: -DGGML_USE_CUDA -DGGML_USE_HIPBLAS -DGGML_CUDA_DMMV_X=32 -DGGML_CUDA_PEER_MAX_BATCH_SIZE=128 -DGGML_CUDA_MMV_Y=1 -D__HIP_PLATFORM_AMD__=1 -D__HIP_ROCclr__=1
+#cgo rocm CXXFLAGS: -DGGML_USE_CUDA -DGGML_USE_HIPBLAS -DGGML_CUDA_DMMV_X=32 -DGGML_CUDA_PEER_MAX_BATCH_SIZE=128 -DGGML_CUDA_MMV_Y=1 -D__HIP_PLATFORM_AMD__=1 -D__HIP_ROCclr__=1
 #cgo rocm LDFLAGS: -L${SRCDIR} -lggml_rocm -lhipblas -lamdhip64 -lrocblas
 #cgo cuda_v11 LDFLAGS: -lggml_cuda_v11 -L/usr/local/cuda-11/lib64
 #cgo cuda_v12 LDFLAGS: -lggml_cuda_v12 -L/usr/local/cuda-12/lib64

filename: llama/make/Makefile.rocm
patch: @@ -47,7 +47,7 @@ ROCM_DIST_DEPS_LIBS = $(addprefix $(ROCM_DIST_DEPS_DIR)/,$(notdir $(ROCM_LIBS))
 ROCBLAS_DIST_DEP_MANIFEST = $(ROCM_DIST_DEPS_DIR)/rocblas/library/TensileManifest.txt
 
 ifeq ($(OS),linux)
-	GPU_COMPILER_FPIC := -fPIC -Wno-unused-function
+	GPU_COMPILER_FPIC := -fPIC -Wno-unused-function -std=gnu++11
 	GPU_RUNNER_ARCH_FLAGS := $(foreach arch, $(HIP_ARCHS_COMMON) $(HIP_ARCHS_LINUX), --offload-arch=$(arch))
 else ifeq ($(OS),windows)
 	GPU_COMPILER_FPIC := -Xclang --dependent-lib=msvcrt
@@ -69,11 +69,15 @@ GPU_COMPILER_CUFLAGS = \
 	-DGGML_USE_HIPBLAS \
 	-DGGML_USE_LLAMAFILE \
 	-DHIP_FAST_MATH \
+	-D__HIP_PLATFORM_AMD__=1 \
+	-D__HIP_ROCclr__=1 \
 	-DNDEBUG \
 	-DK_QUANTS_PER_ITERATION=2 \
 	-D_CRT_SECURE_NO_WARNINGS \
-	-DCMAKE_POSITION_INDEPENDENT_CODE=on \
 	-D_GNU_SOURCE \
+	-D_XOPEN_SOURCE=600 \
+	-mllvm=-amdgpu-early-inline-all=true \
+	-mllvm=-amdgpu-function-calls=false \
 	-Wno-expansion-to-defined \
 	-Wno-invalid-noreturn \
 	-Wno-ignored-attributes \

filename: llama/make/cuda.make
patch: @@ -21,27 +21,27 @@ GPU_COMPILER_CXXFLAGS_WIN = $(CXXFLAGS) -D_WIN32_WINNT=0x602
 GPU_COMPILER_CXXFLAGS_LINUX = $(CXXFLAGS) -Xcompiler -fPIC -D_GNU_SOURCE
 ifeq ($(OS),linux)
 	CUDA_PATH?=/usr/local/cuda
-	GPU_COMPILER_FPIC = -fPIC -Wno-unused-function
+	GPU_COMPILER_FPIC = -fPIC -Wno-unused-function -std=c++11
 endif
 GPU_RUNNER_ARCH_FLAGS := $(foreach arch,$(subst ;,$(space),$(CUDA_ARCHITECTURES)),--generate-code=arch=compute_$(arch)$(comma)code=[compute_$(arch)$(comma)sm_$(arch)]) \
-	-DGGML_CUDA_USE_GRAPHS=on
+	-DGGML_CUDA_USE_GRAPHS=1
 GPU_COMPILER_CUFLAGS = \
 	$(GPU_COMPILER_FPIC) \
 	-Xcompiler "$(addprefix $(CPU_FLAG_PREFIX),$(_OS_GPU_RUNNER_CPU_FLAGS))" \
 	-t2 \
 	-DGGML_CUDA_DMMV_X=32 \
+	-DGGML_CUDA_MMV_Y=1 \
 	-DGGML_CUDA_PEER_MAX_BATCH_SIZE=128 \
 	-DGGML_USE_CUDA=1 \
 	-DGGML_SHARED=1 \
 	-DGGML_BUILD=1 \
 	-DGGML_USE_LLAMAFILE \
+	-DK_QUANTS_PER_ITERATION=2 \
 	-DNDEBUG \
 	-D_GNU_SOURCE \
-	-DCMAKE_POSITION_INDEPENDENT_CODE=on \
+	-D_XOPEN_SOURCE=600 \
 	-Wno-deprecated-gpu-targets \
 	--forward-unknown-to-host-compiler \
 	-use_fast_math \
-	-link \
-	-shared \
 	-I. \
 	-O3