sync : ggml (new ops, tests, backend, etc.)

[ggerganov]

Careful look and testing required - very likely to have broken something during this merge

[ggerganov]

Need help with enabling CUDA separable compilation in the Makefile

[slaren]

I think we can ignore that for now, that was used for the device-side mul mat id implementation, but that path has been disabled for now. The kernel that requires dynamic parallelism should be in an #if 0 block.

[ggerganov]

I see - it seems still required for the dequantization kernels:

ggml-cuda.cu(5229): error: kernel launch from __device__ or __global__ functions requires separate compilation mode
      dequantize_block<qk, qr, dequantize_kernel><<<num_blocks, 256, 0, stream>>>(vx, y, k);
      ^

ggml-cuda.cu(5236): error: kernel launch from __device__ or __global__ functions requires separate compilation mode
      dequantize_block_q2_K<<<nb, 64, 0, stream>>>(vx, y);
      ^

ggml-cuda.cu(5246): error: kernel launch from __device__ or __global__ functions requires separate compilation mode
      dequantize_block_q3_K<<<nb, 64, 0, stream>>>(vx, y);
      ^

ggml-cuda.cu(5255): error: kernel launch from __device__ or __global__ functions requires separate compilation mode
      dequantize_block_q4_K<<<nb, 32, 0, stream>>>(vx, y);
      ^

ggml-cuda.cu(5262): error: kernel launch from __device__ or __global__ functions requires separate compilation mode
      dequantize_block_q5_K<<<nb, 64, 0, stream>>>(vx, y);
      ^

ggml-cuda.cu(5272): error: kernel launch from __device__ or __global__ functions requires separate compilation mode
      dequantize_block_q6_K<<<nb, 64, 0, stream>>>(vx, y);
      ^

6 errors detected in the compilation of "ggml-cuda.cu".
make: *** [Makefile:451: ggml-cuda.o] Error 1
make: *** Waiting for unfinished jobs....

[slaren]

That should fix it.

[slaren]

Are you going to sync the tests too? It's in the PR title, but not in the files.

[ggerganov]

Yes, I forgot - will sync them in a few

[ggerganov]

Seems like when linking to libllama the backend registers do not work. Happens with the cmake build. I think it needs some "whole-archive" flags, but I've forgotten how this works

[slaren]

Ok, I was afraid that this way of self-registering backends was going to be more trouble than it is worth. It may be better to scratch it and just have a list of backends in the registry depending on the compilation flags. The disadvantage is that every time that a new backend is added, it will need to be added to this list manually rather than self-registering with GGML_CONSTRUCTOR, but it will cause less problems.

[ggerganov]

Yeah, my experience with such kind of logic is also not very positive - various platform compatibility problems arise. List of backends based on flags should work, despite the extra manual work to maintain the list