CUDA int128相乘是怎么实现的
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