@@ -177,14 +177,31 @@ ifdef LLAMA_HIPBLAS
177177 ROCM_PATH ?= /opt/rocm
178178 CC := $(ROCM_PATH)/llvm/bin/clang
179179 CXX := $(ROCM_PATH)/llvm/bin/clang++
180- GPU_TARGETS = gfx803 gfx900 gfx906 gfx908 gfx90a gfx1030
180+ GPU_TARGETS = gfx803 gfx900 gfx906 gfx908 gfx90a gfx1030 gfx1100
181181 LLAMA_CUDA_DMMV_X ?= 64
182182 LLAMA_CUDA_DMMV_Y ?= 2
183+ LLAMA_CUDA_FORCE_DMMV = false
183184 CFLAGS += -DGGML_USE_HIPBLAS -DGGML_USE_CUBLAS $(shell $(ROCM_PATH)/bin/hipconfig -C)
184185 CXXFLAGS += -DGGML_USE_HIPBLAS -DGGML_USE_CUBLAS $(shell $(ROCM_PATH)/bin/hipconfig -C)
185186 LDFLAGS += -L/opt/rocm/lib -Wl,-rpath=$(ROCM_PATH)/lib -lhipblas -lamdhip64
186187 OBJS += ggml-cuda.o ggml_v2-cuda.o ggml_v2-cuda-legacy.o
187188
189+ ifdef LLAMA_CUDA_FORCE_DMMV
190+ CXXFLAGS += -DGGML_CUDA_FORCE_DMMV
191+ endif # LLAMA_CUDA_FORCE_DMMV
192+ ifdef LLAMA_CUDA_DMMV_X
193+ CXXFLAGS += -DGGML_CUDA_DMMV_X=$(LLAMA_CUDA_DMMV_X )
194+ else
195+ CXXFLAGS += -DGGML_CUDA_DMMV_X=32
196+ endif # LLAMA_CUDA_DMMV_X
197+ ifdef LLAMA_CUDA_MMV_Y
198+ CXXFLAGS += -DGGML_CUDA_MMV_Y=$(LLAMA_CUDA_MMV_Y )
199+ else ifdef LLAMA_CUDA_DMMV_Y
200+ CXXFLAGS += -DGGML_CUDA_MMV_Y=$(LLAMA_CUDA_DMMV_Y ) # for backwards compatibility
201+ else
202+ CXXFLAGS += -DGGML_CUDA_MMV_Y=1
203+ endif # LLAMA_CUDA_MMV_Y
204+
188205ifdef LLAMA_CUDA_KQUANTS_ITER
189206 CXXFLAGS += -DK_QUANTS_PER_ITERATION=$(LLAMA_CUDA_KQUANTS_ITER)
190207else
@@ -193,7 +210,10 @@ endif
193210
194211ggml-cuda.o : CXXFLAGS += $(addprefix --offload-arch=,$(GPU_TARGETS ) ) \
195212 -DGGML_CUDA_DMMV_X=$(LLAMA_CUDA_DMMV_X ) \
196- -DGGML_CUDA_DMMV_Y=$(LLAMA_CUDA_DMMV_Y )
213+ -DGGML_CUDA_DMMV_Y=$(LLAMA_CUDA_DMMV_Y ) \
214+ -DGGML_CUDA_FORCE_DMMV
215+
216+
197217# DGGML_CUDA_DMMV_F16 does not currently work with AMD.
198218ggml-cuda.o : ggml-cuda.cu ggml-cuda.h
199219 $(CXX ) $(CXXFLAGS ) -x hip -c -o $@ $<
0 commit comments