Converting a kernel from floats and ints to halfs is 6x slower

Accelerated Computing CUDA CUDA Programming and Performance
14

Hi there, unfortunately I have found the “—maxregcount=32” compiler argument to be vital for a different kernel in which not using it resulted in strange behaviour: Weird CUDA problem: changing += to /= in a loop causes a variable not to be set - #11 by AastaLLL
I tried compiling without the “—maxregcount=32” anyway and while this particular kernel ran fine, there was no speed up.

I would like to mention something in case it matters. The diff above still stands; applying that diff and only that diff causes the increased delay effects, but in order to get the code with halfs compiling I had to make sure I was not compiling for older architectures. Originally I had these in my nvcc args:

GENCODE_SM30    := -gencode arch=compute_30,code=sm_30
GENCODE_SM32    := -gencode arch=compute_32,code=sm_32
GENCODE_SM35    := -gencode arch=compute_35,code=sm_35
GENCODE_SM50    := -gencode arch=compute_50,code=sm_50
GENCODE_SMXX    := -gencode arch=compute_50,code=compute_50
GENCODE_SM53    := -gencode arch=compute_53,code=sm_53 -maxrregcount=32
GENCODE_SM62    := -gencode arch=compute_62,code=sm_62
GENCODE_SM72    := -gencode arch=compute_72,code=sm_72
GENCODE_SM_PTX  := -gencode arch=compute_72,code=compute_72

This caused compilation to fail, presumably as older architectures do not support halfs. This is the error I saw:

nvsample_cudaprocess.cu(368): error: more than one conversion function from "__half" to a built-in type applies:
            function "__half::operator float() const"
            function "__half::operator short() const"
            function "__half::operator unsigned short() const"
            function "__half::operator int() const"
            function "__half::operator unsigned int() const"
            function "__half::operator long long() const"
            function "__half::operator unsigned long long() const"
            function "__half::operator __nv_bool() const"

As for the SASS code, I dumped it but the assembly is 2000 lines long for the half implementation and 1000 lines long for floats/halfs. I’ve attached them in case they’re of help. These are the CUDA objdumps of the application version of the function before and after applying the patch I posted above.

kernel_floatint.txt (85.5 KB)
kernel_half.txt (183.5 KB)

Cheers.

(edited formatting)