Hip kernel refactor

[reyna-abhyankar]

• Linear
• Layer Norm

Follow the same procedure as the respective ___kernels.cu file for each operator

[lockshaw]

@reyna-abhyankar Could we get a bit more detail on this issue? Maybe a couple links?

[Bob-Chen222]

I located the code in lib/kernels/src/cuda but am not sure what are the procedures need to be done here. Could you elaborate more on this @reyna-abhyankar? Thanks!

[reyna-abhyankar]

I located the code in lib/kernels/src/cuda but am not sure what are the procedures need to be done here. Could you elaborate more on this @reyna-abhyankar? Thanks!

Essentially, we want the _kernels.cpp file to mirror the _kernels.cu file for all the operators. This means the function signature should match what's defined in _kernels.h. There are some small differences between hip and cuda. We use hipStream_t instead of cudaStream_t, miopenXXXX or hipXXXX instead of cudaXXXX as prefixes for certain types or function calls. Most notably, launching the actual kernel is different. In cuda, we use triple bracket syntax for certain hyperparameters (like # of thread blocks and # of threads). In hip, these are just normal function parameters.

So in linear_kernels.cu, we call gelu_forward_kernel like so:

gelu_forward_kernel<<<GET_BLOCKS(elements), CUDA_NUM_THREADS>>>(
        elements, B, C, (float *)output_ptr);

whereas, in linear_kernels.cpp, we call it like so:

hipLaunchKernelGGL(gelu_forward_kernel,
                       GET_BLOCKS(elements),
                       CUDA_NUM_THREADS,
                       0,
                       0,
                       elements,
                       B,
                       C,
                       (float *)output_ptr);

So aside from these differences, we want to align the .cpp file to the .cu file. For example, the function signature for init_kernel is incorrect in the hip kernels when compared with the cuda version. Let me know if you have more questions! We can chat more over Slack as well.

[Bob-Chen222]

I located the code in lib/kernels/src/cuda but am not sure what are the procedures need to be done here. Could you elaborate more on this @reyna-abhyankar? Thanks!

Essentially, we want the _kernels.cpp file to mirror the _kernels.cu file for all the operators. This means the function signature should match what's defined in _kernels.h. There are some small differences between hip and cuda. We use hipStream_t instead of cudaStream_t, miopenXXXX or hipXXXX instead of cudaXXXX as prefixes for certain types or function calls. Most notably, launching the actual kernel is different. In cuda, we use triple bracket syntax for certain hyperparameters (like # of thread blocks and # of threads). In hip, these are just normal function parameters.

So in linear_kernels.cu, we call gelu_forward_kernel like so:

gelu_forward_kernel<<<GET_BLOCKS(elements), CUDA_NUM_THREADS>>>(
        elements, B, C, (float *)output_ptr);

whereas, in linear_kernels.cpp, we call it like so:

hipLaunchKernelGGL(gelu_forward_kernel,
                       GET_BLOCKS(elements),
                       CUDA_NUM_THREADS,
                       0,
                       0,
                       elements,
                       B,
                       C,
                       (float *)output_ptr);

So aside from these differences, we want to align the .cpp file to the .cu file. For example, the function signature for init_kernel is incorrect in the hip kernels when compared with the cuda version. Let me know if you have more questions! We can chat more over Slack as well.

Thanks Reyna! I will follow up via slack if having more questions.

[Bob-Chen222]

Hi @lockshaw Colin and @reyna-abhyankar Reyna, I have one question regarding hip and Cuda. I have noticed that all check functions are in the CUDA format in the .cpp file. For example, checkCUDNN(miopenSetStream(m->handle.dnn, stream)); is used in hip. So, the question is to confirm that checkCUDNN, checkCUDA, or similar functions are the correct functions to use in hip. Is this correct?

[reyna-abhyankar]

Hi @lockshaw Colin and @reyna-abhyankar Reyna, I have one question regarding hip and Cuda. I have noticed that all check functions are in the CUDA format in the .cpp file. For example, checkCUDNN(miopenSetStream(m->handle.dnn, stream)); is used in hip. So, the question is to confirm that checkCUDNN, checkCUDA, or similar functions are the correct functions to use in hip. Is this correct?

Yes, those are fine. They're defined in device.h and take in a status (so miopenSetStream() returns a status).