Hmm, it looks like the underlying SASS does have 64-bit integer instructions now, but only with the 12.0 capability level in the recent Blackwell processors. Older versions emulate it via chained 32-bit instructions. Take this example kernel:

  __global__ void add(uint64_t *res, uint64_t x) {
      *res = x + 0x12345;
  }

Compiled with -arch=sm_120, I get the SASS:

  /*0000*/  LDC R1, c[0x0][0x37c]                      ?trans8;
  /*0010*/  LDC.64 R2, c[0x0][0x388]          &wr=0x0  ?trans1;
  /*0020*/  LDCU.64 UR4, c[0x0][0x358]        &wr=0x1  ?trans7;
  /*0030*/  LDC.64 R4, c[0x0][0x380]          &wr=0x1  ?trans1;
  /*0040*/  IADD.64 R2, R2, 0x12345           &req={0} ?WAIT6_END_GROUP;
  /*0050*/  STG.E.64 desc[UR4][R4.64], R2     &req={1} ?trans1;
  /*0060*/  EXIT                                       ?trans5;
  /*0070*/  BRA 0x70;

But with -arch=sm_100, the IADD.64 is broken up into a UIADD3 and a UIADD3.X (contrast with the IADD3 that a regular 32-bit addition would produce):

  /*0000*/  LDC R1, c[0x0][0x37c] ;
  /*0010*/  LDCU.64 UR4, c[0x0][0x388] ;
  /*0020*/  LDC.64 R2, c[0x0][0x380] ;
  /*0030*/  LDCU.64 UR6, c[0x0][0x358] ;
  /*0040*/  UIADD3 UR4, UP0, UPT, UR4, 0x12345, URZ ;
  /*0050*/  UIADD3.X UR5, UPT, UPT, URZ, UR5, URZ, UP0, !UPT ;
  /*0060*/  MOV R4, UR4 ;
  /*0070*/  MOV R5, UR5 ;
  /*0080*/  STG.E.64 desc[UR6][R2.64], R4 ;
  /*0090*/  EXIT ;
  /*00a0*/  BRA 0xa0;

So if you want real 64-bit support, have fun getting your hands on a 5070! But even on sm_120, things like 64-bit × immediate 32-bit take a UIMAD.WIDE.U32 + UIMAD + UIADD3 sequence, so the support isn't all that complete.

(I've been looking into the specifics of CUDA integer arithmetic for some time now, since I've had the mad idea of doing 'horizontal' 448-bit integer arithmetic by storing one word in each thread and using the warp-shuffle instructions to send carries up and down. Given that the underlying arithmetic is all 32-bit, it doesn't make any sense to store more than 31 bits per thread. Then again, I don't know whether this mad idea makes any sense in the first place, until I implement and profile it.)