happyz
/
llama.cpp
mirror of https://github.com/ggerganov/llama.cpp.git
1
0
Fork 0
Code Issues Packages Projects Releases Wiki Activity

further reduce index swizzling computation cycles

This commit is contained in:
bssrdf 2025-11-07 23:20:46 -05:00
parent 8809af79a8
commit 414bb8d9ed
2 changed files with 10 additions and 6 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

11
ggml/src/ggml-cuda/conv2d-implicit.cu
View File

@ -691,15 +691,18 @@ static __global__ void conv2d_implicit_kernel(const half * __restrict__ input,
#pragma unroll
for (int subk = 0; subk < WN / 4; ++subk){
const uint row = m_i_wn + subk*2;
uint idx = output_lds_addr + subk*2;
idx = idx ^ ((idx & 0b110000000000) >> 9);
idx = idx ^ ((idx & 0b1110000000) >> 4);
#pragma unroll
for (int j = 0; j < 4; ++j){
const uint gemm_i = n_idx + j*32;
const int n = fastdiv(gemm_i, param.OHOW_fastdiv);
const int col = fastmodulo(gemm_i, param.OHOW_fastdiv);
uint idx = output_lds_addr + subk*2 + j*32*BN/2;
idx = idx ^ ((idx & 0b110000000000) >> 9);
idx = idx ^ ((idx & 0b1110000000) >> 4);
uint32_t dst_ptr = *(reinterpret_cast<uint32_t*>(&smemoutput[idx]));
// uint idx = output_lds_addr + subk*2 + j*32*BN/2;
// idx = idx ^ ((idx & 0b110000000000) >> 9);
// idx = idx ^ ((idx & 0b1110000000) >> 4);
uint32_t dst_ptr = *(reinterpret_cast<uint32_t*>(&smemoutput[idx+j*32*BN/2]));
half (&res_)[2] = reinterpret_cast<half(&)[2]>(dst_ptr);
if (n < param.n && row < param.k && col < PQ) {
if constexpr (ksplit > 0) {

5
tests/test-conv2d.cpp
View File

@ -325,8 +325,8 @@ int main(void)
std::make_tuple(512,256,416,608,3,3),
std::make_tuple(256,128,832,1216,3,3),
std::make_tuple(256,256,832,1216,3,3),
// std::make_tuple(320,256,1024,1920)
std::make_tuple(32,64,58,58,3,3)
// std::make_tuple(320,256,1024,1920)
};
std::vector<std::tuple<int, int, int, int, int, int>> configs_sdxl_512 = {
//512x512
@ -648,7 +648,8 @@ int main(void)
int k = 0;
for (auto c : configs_sdxl_1024){
for (auto c : configs_sdxl_512){
// for (auto c : configs){
test_model model;
load_model(model, std::get<0>(c), std::get<1>(c), std::get<2>(c),
std::get<3>(c), std::get<4>(c), std::get<5>(c), true);