From a9cde5c63e9c61292fa7454e0d3045612417812e Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Johannes=20G=C3=A4=C3=9Fler?= Date: Mon, 10 Jun 2024 08:51:35 +0200 Subject: [PATCH] __builtin_assume -> GGML_CUDA_ASSUME --- ggml-cuda/mma.cuh | 24 ++++++++++++------------ 1 file changed, 12 insertions(+), 12 deletions(-) diff --git a/ggml-cuda/mma.cuh b/ggml-cuda/mma.cuh index 0afc731df..71e8e3429 100644 --- a/ggml-cuda/mma.cuh +++ b/ggml-cuda/mma.cuh @@ -9,15 +9,15 @@ struct mma_int_A_I16K8 { static __device__ __forceinline__ int get_i(const int l) { const int ret = (l%2) * (I/2) + threadIdx.x / (K/2); - __builtin_assume(ret >= 0); - __builtin_assume(ret < I); + GGML_CUDA_ASSUME(ret >= 0); + GGML_CUDA_ASSUME(ret < I); return ret; } static __device__ __forceinline__ int get_k(const int l) { const int ret = (l/2) * (K/2) + threadIdx.x % (K/2); - __builtin_assume(ret >= 0); - __builtin_assume(ret < K); + GGML_CUDA_ASSUME(ret >= 0); + GGML_CUDA_ASSUME(ret < K); return ret; } }; @@ -31,15 +31,15 @@ struct mma_int_B_J8K8 { static __device__ __forceinline__ int get_j(const int /* l */) { const int ret = threadIdx.x / (K/2); - __builtin_assume(ret >= 0); - __builtin_assume(ret < J); + GGML_CUDA_ASSUME(ret >= 0); + GGML_CUDA_ASSUME(ret < J); return ret; } static __device__ __forceinline__ int get_k(const int l) { const int ret = l * (K/2) + threadIdx.x % (K/2); - __builtin_assume(ret >= 0); - __builtin_assume(ret < K); + GGML_CUDA_ASSUME(ret >= 0); + GGML_CUDA_ASSUME(ret < K); return ret; } }; @@ -53,15 +53,15 @@ struct mma_int_C_I16J8 { static __device__ __forceinline__ int get_i(const int l) { const int ret = (l/2) * (I/2) + threadIdx.x / (J/2); - __builtin_assume(ret >= 0); - __builtin_assume(ret < I); + GGML_CUDA_ASSUME(ret >= 0); + GGML_CUDA_ASSUME(ret < I); return ret; } static __device__ __forceinline__ int get_j(const int l) { const int ret = 2 * (threadIdx.x % (J/2)) + l%2; - __builtin_assume(ret >= 0); - __builtin_assume(ret < J); + GGML_CUDA_ASSUME(ret >= 0); + GGML_CUDA_ASSUME(ret < J); return ret; }