CUDA int128相乘是怎么实现的

一.指令对应关系

1.cuda C代码

clock_t t0=clock64();
__prof_trigger(0);
int128_t r4=r2*r3;  //二个int128_t相加
__prof_trigger(1);
clock_t t1=clock64();

2.PTX代码[二个int64]

mov.u64 %rd1, %clock64;
.loc    1 52 5
pmevent 0;
.loc    1 53 5
mul.hi.u64 %rd9, %rd8, %rd7;
mul.wide.u32 %rd10, %r5, %r5;
add.s64 %rd11, %rd9, %rd10;  //二个s64相加
mul.lo.s64 %rd12, %rd7, %rd7;
add.s64 %rd13, %rd11, %rd12; //
mul.lo.s64 %rd14, %rd8, %rd7;
.loc    1 54 5
pmevent 1;
.loc    1 55 5
mov.u64 %rd2, %clock64;

3.SASS指令[多个u32]

/*0070*/                   CS2R R12, SR_CLOCKLO ;
/*0080*/                   PMTRIG 0x1 ;
/*0090*/                   IMAD.WIDE.U32 R4, P0, R11, R9, RZ ;
/*00a0*/                   ULDC.64 UR4, c[0x0][0x118] ;
/*00b0*/                   IMAD.WIDE.U32 R2, R11, R8, RZ ;
/*00c0*/                   MOV R6, R5 ;
/*00d0*/                   IMAD.X R7, RZ, RZ, RZ, P0 ;
/*00e0*/                   IADD3 R0, P0, R3, R4, RZ ;
/*00f0*/                   IMAD.WIDE.U32.X R4, RZ, R9, R6, P0 ;
/*0100*/                   IADD3 R6, P0, R1, c[0x0][0x20], RZ ;
/*0110*/                   IMAD R7, R9, R8, RZ ;
/*0120*/                   IMAD.WIDE.U32 R4, R11, R11, R4 ;
/*0130*/                   IMAD R17, R8.reuse, R9, R7 ;
/*0140*/                   IMAD.WIDE.U32 R8, R8, R8, R4 ;
/*0150*/                   IMAD.X R7, RZ, RZ, c[0x0][0x24], P0 ;
/*0160*/                   IADD3 R17, R9, R17, RZ ;
/*0170*/                   PMTRIG 0x2 ;
/*0180*/                   CS2R R14, SR_CLOCKLO ;

二.复现过程

tee cuda_types.cu<<-'EOF'
#include <iostream>
#include <cuda_runtime.h>
#include <iostream>
#include <vector>
#include <stdio.h>
#include <assert.h>
#include <cstdio>
#include <cuda.h>
#include <cassert>
#include <cstdint>
#include <type_traits>
typedef __int128 int128_t;
typedef unsigned __int128 uint128_t;
__host__ __device__ constexpr uint128_t create_uint128(uint32_t le0,uint32_t le1,uint32_t le2,uint32_t le3) {return (uint128_t)le0 | (uint128_t)le1 << 32 | (uint128_t)le2 << 64 |(uint128_t)le3 << 96;
}
__host__ __device__ constexpr int128_t create_int128(uint32_t le0, uint32_t le1,uint32_t le2,uint32_t le3) {return (int128_t)create_uint128(le0, le1, le2, le3);
}    
__host__ __device__ constexpr uint128_t create_uint128(uint64_t le0,uint64_t le1) {return (uint128_t)le0 | (uint128_t)le1 << 64;
}    
__host__ __device__ constexpr int128_t create_int128(uint64_t le0,uint64_t le1) {return (int128_t)create_uint128(le0, le1);
}
#define CHECK_CUDA(call)                                           \do {                                                           \cudaError_t err = call;                                    \if (err != cudaSuccess) {                                  \std::cerr << "CUDA error at " << __FILE__ << ":" << __LINE__; \std::cerr << " code=" << err << " (" << cudaGetErrorString(err) << ")" << std::endl; \exit(EXIT_FAILURE);                                    \}                                                          \} while (0)__global__ void kernel(uint32_t *addr)
{unsigned int tid  = threadIdx.x + blockIdx.x * blockDim.x;uint64_t r0=tid*2;uint64_t r1=tid*r0;int128_t r2=create_int128(r0,r1);int128_t r3=create_int128(r1,r2);clock_t t0=clock64(); //用来标记ptx和sass指令__prof_trigger(0);    //用来标记ptx和sass指令int128_t r4=r2*r3;__prof_trigger(1);    //用来标记ptx和sass指令clock_t t1=clock64(); //用来标记ptx和sass指令addr[tid]=(uint32_t)r4;addr[tid+4]=(uint32_t)(r4>>32);addr[tid+8]=(uint32_t)(r4>>64);addr[tid+12]=(uint32_t)(r4>>96);printf("%lld\n",t1-t0);
}int main(int argc,char *argv[])
{int deviceid=0;cudaSetDevice(deviceid);  int block_count=28;int block_size=32*4;int thread_size=block_count*block_size;{uint32_t *addr;CHECK_CUDA(cudaMalloc(&addr, thread_size*4));kernel<<<block_count, block_size>>>(addr);CHECK_CUDA(cudaDeviceSynchronize());}    
}
EOF
/usr/local/cuda/bin/nvcc -std=c++17 -arch=sm_86 \-o cuda_types cuda_types.cu -I /usr/local/cuda/include -L /usr/local/cuda/lib64 -lcuda
/usr/local/cuda/bin/cuobjdump --dump-ptx ./cuda_types
/usr/local/cuda/bin/cuobjdump --dump-sass ./cuda_types