@@ -177,14 +177,31 @@ ifdef LLAMA_HIPBLAS
177
177
ROCM_PATH ?= /opt/rocm
178
178
CC := $(ROCM_PATH)/llvm/bin/clang
179
179
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
181
181
LLAMA_CUDA_DMMV_X ?= 64
182
182
LLAMA_CUDA_DMMV_Y ?= 2
183
+ LLAMA_CUDA_FORCE_DMMV = false
183
184
CFLAGS += -DGGML_USE_HIPBLAS -DGGML_USE_CUBLAS $(shell $(ROCM_PATH)/bin/hipconfig -C)
184
185
CXXFLAGS += -DGGML_USE_HIPBLAS -DGGML_USE_CUBLAS $(shell $(ROCM_PATH)/bin/hipconfig -C)
185
186
LDFLAGS += -L/opt/rocm/lib -Wl,-rpath=$(ROCM_PATH)/lib -lhipblas -lamdhip64
186
187
OBJS += ggml-cuda.o ggml_v2-cuda.o ggml_v2-cuda-legacy.o
187
188
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
+
188
205
ifdef LLAMA_CUDA_KQUANTS_ITER
189
206
CXXFLAGS += -DK_QUANTS_PER_ITERATION=$(LLAMA_CUDA_KQUANTS_ITER)
190
207
else
@@ -193,7 +210,10 @@ endif
193
210
194
211
ggml-cuda.o : CXXFLAGS += $(addprefix --offload-arch=,$(GPU_TARGETS ) ) \
195
212
-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
+
197
217
# DGGML_CUDA_DMMV_F16 does not currently work with AMD.
198
218
ggml-cuda.o : ggml-cuda.cu ggml-cuda.h
199
219
$(CXX ) $(CXXFLAGS ) -x hip -c -o $@ $<
0 commit comments