Identifier "__HALF2_TO_UI" is undefined when using asm for cuda

czxu · October 13, 2023, 4:34pm

Hi Forum,
I am trying to use asm code to implement ldg128 and stg128 for cuda global memory access, and here is my main code for this part:

__device__ __forceinline__ void ldg128(const __half2* addr, __half2 &reg0, __half2 &reg1, __half2 &reg2, __half2 &reg3){
    asm volatile(
        "ld.global.nc.v4.b32 {%1, %2, %3, %4}, [%0];\n"
        : "=r"(__HALF2_TO_UI(reg0)),
          "=r"(__HALF2_TO_UI(reg1)),
          "=r"(__HALF2_TO_UI(reg2)),
          "=r"(__HALF2_TO_UI(reg3))
        : "l"(addr)
    );
}

__device__ __forceinline__ void stg128(__half2* addr, __half2 &reg0, __half2 &reg1, __half2 &reg2, __half2 &reg3) {
    asm volatile(
        "st.global.v4.b32 [%0], {%1, %2, %3, %4};\n"
        :
        : "l"(addr),
          "r"(__HALF2_TO_UI(reg0)),
          "r"(__HALF2_TO_UI(reg1)),
          "r"(__HALF2_TO_UI(reg2)),
          "r"(__HALF2_TO_UI(reg3))
    );
}

To compile, I have included

#include <cuda.h>
#include <cuda_runtime.h>
#include <cuda_fp16.h>
#include <cuda_fp16.hpp>

and add my cuda arch in the Cmake (so functions like __hfma. __hfma2 works fine). When I want to compile the code above, an error occured saying:

common.cu(41): error: identifier "__HALF2_TO_UI" is undefined

common.cu(41): error: an asm operand must have scalar type

common.cu(42): error: an asm operand must have scalar type

common.cu(43): error: an asm operand must have scalar type

common.cu(44): error: an asm operand must have scalar type

May I know where did I do wrong? I cat the /usr/local/cuda/include/cuda_fp16.hpp and found the definition of __HALF2_TO_UI, but seems that the code cannot find that…

Thank you for your help!
Chengzhe

Robert_Crovella · October 13, 2023, 4:46pm

That macro is not usable in PTX assembly. The header file you reference is for CUDA C++, not PTX.

Why not just do this operation in CUDA C++?

czxu · October 13, 2023, 4:49pm

Oh I see… how I could include PTX?
I am learning how to use those macros lol, the code is from a former working project, and I try to learn how to utilize the vector ldg/stgs

czxu · October 13, 2023, 4:58pm

The first version of my code is using __ldg()s, and trying to play around vector ldg

Robert_Crovella · October 13, 2023, 6:23pm

these are 32-bit registers. How about just using the variable itself?
instead of:

"=r"(__HALF2_TO_UI(reg0)),

do:

 "=r"(reg0),

according to my testing, that compiles cleanly for me

czxu · October 13, 2023, 6:28pm

Thank you!
I change my code to

__device__ __forceinline__ void ldg128(const __half2* addr, __half2 &reg0, __half2 &reg1, __half2 &reg2, __half2 &reg3){
    asm volatile(
        "ld.global.nc.v4.b32 {%1, %2, %3, %4}, [%0];\n"
        : "=r"(reg0),
          "=r"(reg1),
          "=r"(reg2),
          "=r"(reg3)
        : "l"(addr)
    );
}

__device__ __forceinline__ void stg128(__half2* addr, __half2 &reg0, __half2 &reg1, __half2 &reg2, __half2 &reg3) {
    asm volatile(
        "st.global.v4.b32 [%0], {%1, %2, %3, %4};\n"
        :
        : "l"(addr),
          "r"(reg0),
          "r"(reg1),
          "r"(reg2),
          "r"(reg3)
    );
}

but it says that error: an asm operand must have scalar type. Do you know what does that means?

Robert_Crovella · October 13, 2023, 6:32pm

~~Which CUDA version are you using? Please switch to CUDA 12.2 or newer.~~

czxu · October 13, 2023, 6:36pm

Oh, it’s cuda_11.4.r11.4… Got it, I will install CUDA 12.2! Thank you!!
// btw do we have some cuda 12.2 based dockers for orin developer kit lol?

Robert_Crovella · October 13, 2023, 6:39pm

cancel that, I made a mistake.

czxu · October 13, 2023, 6:44pm

you mean about the cuda version12.2? lol

Robert_Crovella · October 13, 2023, 6:57pm

This seems to work:

# cat t63.cu
#include <cuda_fp16.h>

__device__ __forceinline__ void ldg128(const __half2* addr, __half2 &reg0, __half2 &reg1, __half2 &reg2, __half2 &reg3){
    unsigned reg0u, reg1u, reg2u, reg3u;
    asm volatile(
        "ld.global.nc.v4.b32 {%0, %1, %2, %3}, [%4];\n"
        : "=r"(reg0u),
          "=r"(reg1u),
          "=r"(reg2u),
          "=r"(reg3u)
        : "l"(addr)
    );
    reg0 = *(reinterpret_cast<__half2 *>(&reg0u));
    reg1 = *(reinterpret_cast<__half2 *>(&reg1u));
    reg2 = *(reinterpret_cast<__half2 *>(&reg2u));
    reg3 = *(reinterpret_cast<__half2 *>(&reg3u));
}

__global__ void k(__half2* a, __half2 *b){
        __half2 reg0, reg1, reg2, reg3;
        ldg128(a, reg0, reg1, reg2, reg3);
        b[0] = reg0;
        b[1] = reg1;
        b[2] = reg2;
        b[3] = reg3;
}
int main(){

  __half2 *a = NULL;
  k<<<1,1>>>(a, a);
  cudaDeviceSynchronize();
}
# nvcc t63.cu -o t63
# cuobjdump -sass ./t63

Fatbin elf code:
================
arch = sm_52
code version = [1,7]
host = linux
compile_size = 64bit

        code for sm_52

Fatbin elf code:
================
arch = sm_52
code version = [1,7]
host = linux
compile_size = 64bit

        code for sm_52
                Function : _Z1kP7__half2S0_
        .headerflags    @"EF_CUDA_TEXMODE_UNIFIED EF_CUDA_64BIT_ADDRESS EF_CUDA_SM52 EF_CUDA_VIRTUAL_SM(EF_CUDA_SM52)"
                                                              /* 0x001fc800fe2007f6 */
        /*0008*/                   MOV R1, c[0x0][0x20] ;     /* 0x4c98078000870001 */
        /*0010*/                   MOV R2, c[0x0][0x140] ;    /* 0x4c98078005070002 */
        /*0018*/                   MOV R3, c[0x0][0x144] ;    /* 0x4c98078005170003 */
                                                              /* 0x001fc800fe2007b1 */
        /*0028*/                   LDG.E.CI.U.128 R8, [R2] ;  /* 0xeed7a00000070208 */
        /*0030*/                   MOV R4, c[0x0][0x148] ;    /* 0x4c98078005270004 */
        /*0038*/                   MOV R5, c[0x0][0x14c] ;    /* 0x4c98078005370005 */
                                                              /* 0x001fc400fe2107f1 */
        /*0048*/                   STG.E [R4], R8 ;           /* 0xeedc200000070408 */
        /*0050*/                   STG.E [R4+0x4], R9 ;       /* 0xeedc200000470409 */
        /*0058*/                   STG.E [R4+0x8], R10 ;      /* 0xeedc20000087040a */
                                                              /* 0x001fbc00fde007f1 */
        /*0068*/                   STG.E [R4+0xc], R11 ;      /* 0xeedc200000c7040b */
        /*0070*/                   NOP ;                      /* 0x50b0000000070f00 */
        /*0078*/                   NOP ;                      /* 0x50b0000000070f00 */
                                                              /* 0x001ffc00ffe007e2 */
        /*0088*/                   NOP ;                      /* 0x50b0000000070f00 */
        /*0090*/                   EXIT ;                     /* 0xe30000000007000f */
        /*0098*/                   BRA 0x98 ;                 /* 0xe2400fffff87000f */
                                                              /* 0x001f8000fc0007e0 */
        /*00a8*/                   NOP;                       /* 0x50b0000000070f00 */
        /*00b0*/                   NOP;                       /* 0x50b0000000070f00 */
        /*00b8*/                   NOP;                       /* 0x50b0000000070f00 */
                ..........



Fatbin ptx code:
================
arch = sm_52
code version = [8,2]
host = linux
compile_size = 64bit
compressed
#

You can see in the compiled code there aren’t any extra instructions.

czxu · October 13, 2023, 7:01pm

Thank you! I will try it!
So it seems that we need to pass registers as unsigned ints into the asm function, either use __HALF2_TO_UI or use reinterpret_cast to convert __half2, right?

Robert_Crovella · October 13, 2023, 7:02pm

I haven’t studied __HALF2_TO_UI

czxu · October 13, 2023, 7:02pm

btw is the cuobjdump converting c to sass? can it do PTX? super cool tools!

Robert_Crovella · October 13, 2023, 7:08pm

__HALF2_TO_UI is not in cuda_fp16.h, it is in cuda_fp16.hpp

I took a look in that file, and it seems to be doing what I did:

#define __HALF2_TO_UI(var) *(reinterpret_cast<unsigned int *>(&(var)))

The CUDA documentation is here. cuobjdump is part of the binary utilities, look for “CUDA Binary Utilities” on the left, and click on it.

czxu · October 13, 2023, 7:14pm

That’s really awesome. Thank you so much!

Robert_Crovella · October 13, 2023, 8:53pm

This is how I would do the same thing using CUDA C++. I don’t see a need to use inline PTX for this.

# cat t63.cu
#include <cuda_fp16.h>
__device__ __forceinline__ void ldg128(const __half2* addr, __half2 &reg0, __half2 &reg1, __half2 &reg2, __half2 &reg3){

    const int4 val = __ldg(reinterpret_cast<const int4 *>(addr));
    reg0 = *(reinterpret_cast<const __half2 *>(&val.x));
    reg1 = *(reinterpret_cast<const __half2 *>(&val.y));
    reg2 = *(reinterpret_cast<const __half2 *>(&val.z));
    reg3 = *(reinterpret_cast<const __half2 *>(&val.w));
}

__global__ void k(__half2* a, __half2 *b){
        __half2 reg0, reg1, reg2, reg3;
        ldg128(a, reg0, reg1, reg2, reg3);
        b[0] = reg0;
        b[1] = reg1;
        b[2] = reg2;
        b[3] = reg3;
}
int main(){

  __half2 *a = NULL;
  k<<<1,1>>>(a, a);
  cudaDeviceSynchronize();
}
# nvcc t63.cu -o t63
# cuobjdump -sass ./t63

Fatbin elf code:
================
arch = sm_52
code version = [1,7]
host = linux
compile_size = 64bit

        code for sm_52

Fatbin elf code:
================
arch = sm_52
code version = [1,7]
host = linux
compile_size = 64bit

        code for sm_52
                Function : _Z1kP7__half2S0_
        .headerflags    @"EF_CUDA_TEXMODE_UNIFIED EF_CUDA_64BIT_ADDRESS EF_CUDA_SM52 EF_CUDA_VIRTUAL_SM(EF_CUDA_SM52)"
                                                              /* 0x001fc800fe2007f6 */
        /*0008*/                   MOV R1, c[0x0][0x20] ;     /* 0x4c98078000870001 */
        /*0010*/                   MOV R2, c[0x0][0x140] ;    /* 0x4c98078005070002 */
        /*0018*/                   MOV R3, c[0x0][0x144] ;    /* 0x4c98078005170003 */
                                                              /* 0x001fc800fe2007b1 */
        /*0028*/                   LDG.E.CI.U.128 R8, [R2] ;  /* 0xeed7a00000070208 */
        /*0030*/                   MOV R4, c[0x0][0x148] ;    /* 0x4c98078005270004 */
        /*0038*/                   MOV R5, c[0x0][0x14c] ;    /* 0x4c98078005370005 */
                                                              /* 0x001fc400fe2107f1 */
        /*0048*/                   STG.E [R4], R8 ;           /* 0xeedc200000070408 */
        /*0050*/                   STG.E [R4+0x4], R9 ;       /* 0xeedc200000470409 */
        /*0058*/                   STG.E [R4+0x8], R10 ;      /* 0xeedc20000087040a */
                                                              /* 0x001fbc00fde007f1 */
        /*0068*/                   STG.E [R4+0xc], R11 ;      /* 0xeedc200000c7040b */
        /*0070*/                   NOP ;                      /* 0x50b0000000070f00 */
        /*0078*/                   NOP ;                      /* 0x50b0000000070f00 */
                                                              /* 0x001ffc00ffe007e2 */
        /*0088*/                   NOP ;                      /* 0x50b0000000070f00 */
        /*0090*/                   EXIT ;                     /* 0xe30000000007000f */
        /*0098*/                   BRA 0x98 ;                 /* 0xe2400fffff87000f */
                                                              /* 0x001f8000fc0007e0 */
        /*00a8*/                   NOP;                       /* 0x50b0000000070f00 */
        /*00b0*/                   NOP;                       /* 0x50b0000000070f00 */
        /*00b8*/                   NOP;                       /* 0x50b0000000070f00 */
                ..........



Fatbin ptx code:
================
arch = sm_52
code version = [8,2]
host = linux
compile_size = 64bit
compressed
#

This of course requires the __half pointer to be properly aligned for a 128 bit load, but the other variants require that also.

njuffa · October 13, 2023, 9:26pm

Side-remark: My understanding is that bit_cast (new in C++20) is generally preferable to reinterpret_cast, as discussed here, for example:

The CUDA Programming Guide seems to list excluded C++20 features rather than providing a positive list of included features, and I have not yet tried to use it myself.

Topic		Replies	Views
How to write inline asm function sts128 (store 128 bits to shared memory)? CUDA Programming and Performance cuda	13	1213	October 27, 2023
Problem manipulating half precision variables in CUDA kernel Jetson TX2	3	1595	November 28, 2017
Does static_cast can support half data type? CUDA NVCC Compiler	1	228	August 24, 2024
error when trying to use half (fp16) CUDA Programming and Performance	16	20637	October 13, 2015
error: identifier "__hadd2" is undefined Legacy PGI Compilers (archived)	2	4155	March 26, 2020
This half2 behavior needs attention CUDA Programming and Performance cuda	8	537	June 28, 2023
Half2 atomics generate unused code CUDA Programming and Performance	12	484	August 8, 2024
Error: identifier "__hdiv" is undefined when include cuda_fp16.h in .cu CUDA NVCC Compiler	4	1394	August 5, 2024
__half2uchar_rz won't compile CUDA NVCC Compiler cuda , jetson , jetson-nano	2	507	August 31, 2023
CUDA 4.0 asm operand constraints for 8/16-bit types CUDA Programming and Performance	3	4359	March 20, 2012

Identifier "__HALF2_TO_UI" is undefined when using asm for cuda

Related topics