From a09cf7520a7fe1dfe32b99b6fd2ad55d7951ffc9 Mon Sep 17 00:00:00 2001
From: 0cc4m <picard12@live.de>
Date: Tue, 13 Jun 2023 10:10:16 +0200
Subject: [PATCH] Fix SHUF_BLOCKSIZE value

---
 exllama_ext/cuda_func/column_remap.cu | 7 +++++++
 1 file changed, 7 insertions(+)

diff --git a/exllama_ext/cuda_func/column_remap.cu b/exllama_ext/cuda_func/column_remap.cu
index 65514c59..adc91fed 100644
--- a/exllama_ext/cuda_func/column_remap.cu
+++ b/exllama_ext/cuda_func/column_remap.cu
@@ -1,7 +1,14 @@
 #include "column_remap.cuh"
 #include "../util.cuh"
 
+// Using 1024 make me crash with "Memory access fault by GPU node-1 (Agent
+// handle: 0x012345678912) on address 0x012345678912. Reason: Page not present
+// or supervisor privilege."
+#if defined(USE_ROCM)
 const int SHUF_BLOCKSIZE_X = 256;
+#else
+const int SHUF_BLOCKSIZE_X = 1024;
+#endif
 const int SHUF_BLOCKSIZE_Y = 16;
 
 __global__ void column_remap_kernel