CUDA: add softmax broadcast

[am17an]

Target #14435, note that is just the softmax. I also refactored the code a little bit to make it cleaner.

[JohannesGaessler]

Since you were asking me for things to do, consider also tackling this. It's not an immediate problem but if we ever want to do sampling on the GPU it will be. See

llama.cpp/ggml/src/ggml-cuda/mmq.cuh

Lines 3019 to 3026 in 343b6e9

	#if !(defined(GGML_USE_HIP) && defined(__HIP_PLATFORM_AMD__)) && !defined(GGML_USE_MUSA)
	static bool shared_memory_limit_raised[GGML_CUDA_MAX_DEVICES] = {false};
	if (!shared_memory_limit_raised[id]) {
	CUDA_CHECK(cudaFuncSetAttribute(mul_mat_q<type, mmq_x, MMQ_NWARPS, false>, cudaFuncAttributeMaxDynamicSharedMemorySize, nbytes_shared));
	CUDA_CHECK(cudaFuncSetAttribute(mul_mat_q<type, mmq_x, MMQ_NWARPS, true>, cudaFuncAttributeMaxDynamicSharedMemorySize, nbytes_shared));
	shared_memory_limit_raised[id] = true;
	}
	#endif // !(defined(GGML_USE_HIP) && defined(__HIP_PLATFORM_AMD__)) && !defined(GGML_USE_MUSA)

. (This can be a separate PR.)

[JohannesGaessler]

Also, I checked with NVIDIA engineers: it is not possible to raise the shared memory limit universally, you have to do it manually for each function. Yes, it's stupid.

[am17an]

Let me know if I understand this correctly - prior to calling any which uses shared mem, we should somehow call cudaFuncSetAttribute for that kernel. If so, can we create macro for this?

[JohannesGaessler]

This only matters for like 3 kernels that sometimes want to use more than 48 kB of shared memory. The desired behavior is to raise the shared memory limit once per function and device. Raising the limit multiple times does not result in an error but it makes the code slower (thus the static variable).

A macro could conceivably be used to solve this issue.

[am17an]

I'll take a look at this