@@ -199,28 +199,24 @@ ifdef LLAMA_HIPBLAS
199
199
CXX := $(ROCM_PATH)/llvm/bin/clang++
200
200
GPU_TARGETS = gfx803 gfx900 gfx906 gfx908 gfx90a gfx1030 gfx1100
201
201
LLAMA_CUDA_DMMV_X ?= 128
202
- LLAMA_CUDA_MMV_Y ?= 1
202
+ LLAMA_CUDA_MMV_Y ?= 2
203
203
LLAMA_CUDA_KQUANTS_ITER ?= 1
204
- LLAMA_CUDA_FORCE_DMMV ?= true
205
204
HIPFLAGS += -DGGML_USE_HIPBLAS -DGGML_USE_CUBLAS $(shell $(ROCM_PATH)/bin/hipconfig -C)
206
205
HIPLDFLAGS += -L$(ROCM_PATH)/lib -Wl,-rpath=$(ROCM_PATH)/lib -lhipblas -lamdhip64 -lrocblas
207
206
HIP_OBJS += ggml-cuda.o ggml_v2-cuda.o ggml_v2-cuda-legacy.o
208
207
ggml-cuda.o : HIPFLAGS += $(addprefix --offload-arch=,$(GPU_TARGETS ) ) \
209
208
-DGGML_CUDA_DMMV_X=$(LLAMA_CUDA_DMMV_X ) \
210
209
-DGGML_CUDA_MMV_Y =$(LLAMA_CUDA_MMV_Y ) \
211
- -DGGML_CUDA_FORCE_DMMV \
212
210
-DK_QUANTS_PER_ITERATION=$(LLAMA_CUDA_KQUANTS_ITER ) \
213
211
-DCC_TURING=1000000000
214
212
ggml_v2-cuda.o : HIPFLAGS += $(addprefix --offload-arch=,$(GPU_TARGETS ) ) \
215
213
-DGGML_CUDA_DMMV_X=$(LLAMA_CUDA_DMMV_X ) \
216
214
-DGGML_CUDA_MMV_Y =$(LLAMA_CUDA_MMV_Y ) \
217
- -DGGML_CUDA_FORCE_DMMV \
218
215
-DK_QUANTS_PER_ITERATION=$(LLAMA_CUDA_KQUANTS_ITER ) \
219
216
-DCC_TURING=1000000000
220
217
ggml_v2-cuda-legacy.o : HIPFLAGS += $(addprefix --offload-arch=,$(GPU_TARGETS ) ) \
221
218
-DGGML_CUDA_DMMV_X=$(LLAMA_CUDA_DMMV_X ) \
222
219
-DGGML_CUDA_MMV_Y =$(LLAMA_CUDA_MMV_Y ) \
223
- -DGGML_CUDA_FORCE_DMMV \
224
220
-DK_QUANTS_PER_ITERATION=$(LLAMA_CUDA_KQUANTS_ITER ) \
225
221
-DCC_TURING=1000000000 # DGGML_CUDA_DMMV_F16 does not currently work with AMD.
226
222
ggml-cuda.o: ggml-cuda.cu ggml-cuda.h
0 commit comments