[QUESTION] CUDA API?

[EJainDev]

Why is the CUDA API still supported? NVIDIA launched CUTLASS to be an open-source version of cuBLAS and by now, CUDA and OpenCL have diverged so much that it is nearly impossible to upgrade the API to newer versions. Would you consider a major release of CLBlast that removes CUDA compatibility (sometime in the future) enabling contributors to make more OpenCL specific changes such as clCompileProgram + clLinkProgram, features in OpenCL C like async_work_group_copy without any checks, etc.

[EJainDev]

Hi, can you please let me know if you are willing to remove CUDA support so that I can work towards doing that if I have the time.

[CNugteren]

I'm not sure how many people actually used this. I myself found it useful when running on NVIDIA hardware to find out how much of the performance difference between cuBLAS comes from the fact that it is OpenCL and how much from the actual kernel implementation.

But I guess it would make sense to remove it to simplify the library. Unless people are using the backend, not sure? Let's wait a few months to see if there is anyone here that replies to this issue as a CUDA-backend user of CLBlast. And then remove it if there is no response.

[EJainDev]

@CNugteren do you want to continue on this removal or should we wait a few more months?

[CNugteren]

If you believe this is the path forward, then let's go ahead. However, I would suggest to do this only at a point where you don't have any other work ongoing or open PRs, otherwise this might yield too many conflicts.

[vdualb]

Today I found out that it's possible to get ptx out of compiled opencl program and use it in cuda. I only tested it on one simple kernel yet tho. Would that help to simplify this project?

[EJainDev]

Hi @RocketRide9, could you share how it works? Might be useful to keep the CUDA API and simplify the codebase.

[vdualb]

Basically, one needs to query CL_PROGRAM_BINARIES by calling clProgramGetInfo. On Nvidia platform, this function returns string containing ptx. I don't know well how cuda c api works, but in managedCuda I passed ptx into this function: https://github.com/kunzmi/managedCuda/blob/f39a6822c7e9e4450984cf8b1112eb963e64190b/ManagedCUDA/CudaContext.cs#L1063 and then was able to execute the kernel. I can upload full program that demonstrates the process tomorrow.

[EJainDev]

Hi RocketRide9, thanks for your help. I have a few thoughts on this that could be experimented with.

1. Wouldn't that essentially defeat the purpose of the CUDA API? The reason for the CUDA API is that NVIDIA's CUDA compiler can generate better IR and Assembly for CUDA than it's OpenCL compiler can for OpenCL.
2. Could we do PTX generation through OpenCL, then use NVIDIA's PTX compiler for PTX compilation and potentially use that for speed with the OpenCL clCreateProgramFromBinaries function?
3. If the PTX generated by compiling from CUDA can be loaded as a program binary for OpenCL, that could provide a consistent interface, less management, and still have performance benefits from CUDA. This would be experimental for now since it isn't thoroughly tested, but maybe it works. @CNugteren what are your thoughts on this?

[vdualb]

CUDA doesn't have enqueue offset and it can be an issue. See: https://stackoverflow.com/questions/13803882/cuda-ptx-code-envreg32-special-registers. For 1D work, envreg3 is initialised by OpenCL driver and is used to store an offset. I assume this project doesn't rely on offsets and if CUDA initialises envreg3 as 0, then using PTX generated by OpenCL compiler might work.

Here is how get_global_id(0) looks in PTX:

mov.b32 	%r12, %envreg3;  <- offset
mov.u32 	%r14, %ntid.x;   <- local size
mov.u32 	%r13, %ctaid.x;  <- group id
mov.u32 	%r15, %tid.x;    <- local id
add.s32 	%r16, %r15, %r12;
mad.lo.s32 	%r17, %r14, %r13, %r16; <- multiply and add to get global id

I thought a possible solution could be to use something else instead of get_global_id(x), like get_group_id(x) * get_local_size(x) + get_local_id(x). But it compiles to something like (envreg0 + ctaid.x) * ntid.x + tid.x.

Then I checked what work item functions are compiled to. I got these results:

get_group_id({0,1,2}): envreg{0,1,2} + ctaid{x,y,z}
get_local_size({0,1,2}): ntid{x,y,z}
get_enqueued_local_size({0,1,2}): ntid{x,y,z}
get_global_offset({0,1,2}): envreg{3,4,5} - (envreg{0,1,2} * ntid{x,y,z})
get_local_id(): tid{x,y,z}
get_global_size({0,1,2}): ntid{x,y,z} * envreg{6,7,8}
get_num_groups({0,1,2}): envreg{6,7,8}
get_work_dim(): envreg9

I don't see how can to replace get_global_id. PTX of get_group_id and get_global_offset doesn't make sense to me.

Wouldn't that essentially defeat the purpose of the CUDA API?

Nsight dropped support for OpenCL. CUDA backend can be useful to get profiling info for your kernels with Nsight.

[CNugteren]

Interesting, but I agree with @EJainDev that this idea would defeat the purpose of the current CLBlast CUDA API. I would propose to either keep it as is (fully CUDA) or remove it (if it is in the way).