Update Makefile

Makefile

@@ -178,9 +178,9 @@ ifdef LLAMA_HIPBLAS
 	CC         := $(ROCM_PATH)/llvm/bin/clang
 	CXX        := $(ROCM_PATH)/llvm/bin/clang++
 	GPU_TARGETS = gfx803 gfx900 gfx906 gfx908 gfx90a gfx1030 gfx1100
-	LLAMA_CUDA_DMMV_X ?= 256
-	LLAMA_CUDA_DMMV_Y ?= 2
-
+	LLAMA_CUDA_DMMV_X ?= 64
+	LLAMA_CUDA_MMV_Y ?= 2
+	LLAMA_CUDA_FORCE_DMMV = true
 	CFLAGS     += -DGGML_USE_HIPBLAS -DGGML_USE_CUBLAS $(shell $(ROCM_PATH)/bin/hipconfig -C)
 	CXXFLAGS   += -DGGML_USE_HIPBLAS -DGGML_USE_CUBLAS $(shell $(ROCM_PATH)/bin/hipconfig -C)
 	LDFLAGS    += -L/opt/rocm/lib -Wl,-rpath=$(ROCM_PATH)/lib -lhipblas -lamdhip64
@@ -190,25 +190,25 @@ ifdef LLAMA_CUDA_DMMV_X
     CXXFLAGS += -DGGML_CUDA_DMMV_X=$(LLAMA_CUDA_DMMV_X)
 else
     CXXFLAGS += -DGGML_CUDA_DMMV_X=32
-endif # LLAMA_CUDA_DMMV_X
+endif 
+ifeq ($(LLAMA_CUDA_FORCE_DMMV), true)
+    CXXFLAGS += -DGGML_CUDA_FORCE_DMMV
+endif
 ifdef LLAMA_CUDA_MMV_Y
     CXXFLAGS += -DGGML_CUDA_MMV_Y=$(LLAMA_CUDA_MMV_Y)
 else ifdef LLAMA_CUDA_DMMV_Y
     CXXFLAGS += -DGGML_CUDA_MMV_Y=$(LLAMA_CUDA_DMMV_Y) # for backwards compatibility
 else
     CXXFLAGS += -DGGML_CUDA_MMV_Y=1
-endif # LLAMA_CUDA_MMV_Y
+endif
 
 ifdef LLAMA_CUDA_KQUANTS_ITER
 	CXXFLAGS += -DK_QUANTS_PER_ITERATION=$(LLAMA_CUDA_KQUANTS_ITER)
 else
 	CXXFLAGS += -DK_QUANTS_PER_ITERATION=2
 endif
 
-ggml-cuda.o: CXXFLAGS += $(addprefix --offload-arch=,$(GPU_TARGETS)) \
-				-DGGML_CUDA_DMMV_X=$(LLAMA_CUDA_DMMV_X) \
-				-DGGML_CUDA_MMV_Y=$(LLAMA_CUDA_DMMV_Y) \
-				-DGGML_CUDA_FORCE_DMMV
+ggml-cuda.o: CXXFLAGS += $(addprefix --offload-arch=,$(GPU_TARGETS))
 
 
 # DGGML_CUDA_DMMV_F16 does not currently work with AMD.