Похоже, что вычисление CUDA с плавающей точкой двойной точности log включает (небольшое количество) вычислений с плавающей точкой одинарной точности. Я думаю, что это не должно быть поводом для беспокойства, само по себе.
Аппаратное обеспечение графического процессора не поддерживает собственное вычисление двойной точности log, поэтому при обнаружении компилятор заменяет инструкцию последовательностью операций (например, сложение, умножение и т. Д.), Предназначенных для выполнения вычисления. Мы можем подтвердить это простым примером:
$ cat t288.cu
#include <math.h>
#include <stdio.h>
__global__ void kernel(double *y)
{
double x = *y;
*y = log(x);
}
int main(){
double *d_y, h_y = 10.0;
cudaMalloc(&d_y, sizeof(double));
cudaMemcpy(d_y, &h_y, sizeof(double), cudaMemcpyHostToDevice);
kernel<<<1,1>>>(d_y);
cudaDeviceSynchronize();
cudaMemcpy(&h_y, d_y, sizeof(double), cudaMemcpyDeviceToHost);
printf("val = %f\n", h_y);
}
$ nvcc -arch=sm_35 -o t288 t288.cu
$ cuda-memcheck ./t288
========= CUDA-MEMCHECK
val = 2.302585
========= ERROR SUMMARY: 0 errors
$ CUDA_VISIBLE_DEVICES="1" nvprof --metrics flop_count_dp,flop_count_sp,flop_count_sp_special ./t288
==14909== NVPROF is profiling process 14909, command: ./t288
==14909== Some kernel(s) will be replayed on device 0 in order to collect all events/metrics.
Replaying kernel "kernel(double*)" (done)
val = 2.302585rnal events
==14909== Profiling application: ./t288
==14909== Profiling result:
==14909== Metric result:
Invocations Metric Name Metric Description Min Max Avg
Device "Tesla K40m (0)"
Kernel: kernel(double*)
1 flop_count_dp Floating Point Operations(Double Precision) 44 44 44
1 flop_count_sp Floating Point Operations(Single Precision) 0 0 0
1 flop_count_sp_special Floating Point Operations(Single Precision Special) 1 1 1
$ cuobjdump -sass ./t288
Fatbin elf code:
================
arch = sm_35
code version = [1,7]
producer = <unknown>
host = linux
compile_size = 64bit
code for sm_35
Fatbin elf code:
================
arch = sm_35
code version = [1,7]
producer = cuda
host = linux
compile_size = 64bit
code for sm_35
Function : _Z6kernelPd
.headerflags @"EF_CUDA_SM35 EF_CUDA_PTX_SM(EF_CUDA_SM35)"
/* 0x088010fc10a010ac */
/*0008*/ MOV R1, c[0x0][0x44]; /* 0x64c03c00089c0006 */
/*0010*/ MOV R2, c[0x0][0x140]; /* 0x64c03c00281c000a */
/*0018*/ MOV R3, c[0x0][0x144]; /* 0x64c03c00289c000e */
/*0020*/ LD.E.64 R12, [R2]; /* 0xc5800000001c0830 */
/*0028*/ MOV32I R0, 0xfffffc01; /* 0x747ffffe009fc002 */
/*0030*/ ISETP.GT.AND P0, PT, R13, c[0x2][0x0], PT; /* 0x5b481c40001c341e */
/*0038*/ MOV R4, R12; /* 0xe4c03c00061c0012 */
/* 0x08ac1080a0a0a4ac */
/*0048*/ MOV R5, R13; /* 0xe4c03c00069c0016 */
/*0050*/ @!P0 DMUL R4, R4, 1.80143985094819840000e+16; /* 0xc400021a80201011 */
/*0058*/ @!P0 MOV R13, R5; /* 0xe4c03c0002a00036 */
/*0060*/ IADD32I R6, R13, -0x1; /* 0x407fffffff9c3419 */
/*0068*/ ISETP.LT.U32.AND P1, PT, R6, c[0x2][0x4], PT; /* 0x5b101c40009c183e */
/*0070*/ @!P0 MOV R12, R4; /* 0xe4c03c0002200032 */
/*0078*/ @!P0 MOV32I R0, 0xfffffbcb; /* 0x747ffffde5a3c002 */
/* 0x08b88010a4a010ac */
/*0088*/ @P1 BRA 0xc0; /* 0x120000001804003c */
/*0090*/ MOV32I R8, 0x0; /* 0x74000000001fc022 */
/*0098*/ MOV32I R9, 0x7ff00000; /* 0x743ff800001fc026 */
/*00a0*/ DFMA R8, R4, +INF , R8; /* 0xb38023ff801c1021 */
/*00a8*/ FCMP.NEU R12, R8, RZ, R5; /* 0xdd6814007f9c2032 */
/*00b0*/ FCMP.NEU R13, R9, -QNAN , R5; /* 0xbd6817ff801c2435 */
/*00b8*/ BRA 0x240; /* 0x12000000c01c003c */
/* 0x08a010b010a0a010 */
/*00c8*/ LOP32I.AND R4, R13, 0x800fffff; /* 0x204007ffff9c3410 */
/*00d0*/ IMAD.U32.U32.HI R0, R13, 0x1000, R0; /* 0xa2000008001c3401 */
/*00d8*/ LOP32I.OR R5, R4, 0x3ff00000; /* 0x211ff800001c1014 */
/*00e0*/ ISETP.LT.AND P0, PT, R5, c[0x2][0x8], PT; /* 0x5b181c40011c141e */
/*00e8*/ MOV R4, R12; /* 0xe4c03c00061c0012 */
/*00f0*/ @!P0 IADD32I R7, R5, -0x100000; /* 0x407ff8000020141d */
/*00f8*/ @!P0 IADD32I R0, R0, 0x1; /* 0x4000000000a00001 */
/* 0x08a010a01080a010 */
/*0108*/ @!P0 MOV R5, R7; /* 0xe4c03c0003a00016 */
/*0110*/ LOP32I.XOR R12, R0, 0x80000000; /* 0x22400000001c0030 */
/*0118*/ DADD R6, R4, 1; /* 0xc38001ff801c1019 */
/*0120*/ MOV R8, RZ; /* 0xe4c03c007f9c0022 */
/*0128*/ DADD R4, R4, -1; /* 0xcb8001ff801c1011 */
/*0130*/ MUFU.RCP64H R9, R7; /* 0x84000000031c1c26 */
/*0138*/ MOV32I R14, 0x8b7a8b04; /* 0x7445bd45821fc03a */
/* 0x08a01080a4a4a4a4 */
/*0148*/ DFMA R6, -R6, R8, c[0x2][0x10]; /* 0x9b882040021c181a */
/*0150*/ DFMA R6, R6, R6, R6; /* 0xdb801800031c181a */
/*0158*/ DFMA R6, R8, R6, R8; /* 0xdb802000031c201a */
/*0160*/ DMUL R8, R6, R4; /* 0xe4000000021c1822 */
/*0168*/ DFMA R8, R6, R4, R8; /* 0xdb802000021c1822 */
/*0170*/ MOV32I R15, 0x3ed0ee25; /* 0x741f6877129fc03e */
/*0178*/ MOV32I R13, 0x43300000; /* 0x74219800001fc036 */
/* 0x08a080a080a4a4a4 */
/*0188*/ DMUL R10, R8, R8; /* 0xe4000000041c202a */
/*0190*/ DFMA R14, R10, c[0x2][0x18], R14; /* 0x5b803840031c283a */
/*0198*/ DFMA R14, R10, R14, c[0x2][0x20]; /* 0x9b803840041c283a */
/*01a0*/ DFMA R14, R10, R14, c[0x2][0x28]; /* 0x9b803840051c283a */
/*01a8*/ DADD R16, R4, -R8; /* 0xe3810000041c1042 */
/*01b0*/ DFMA R14, R10, R14, c[0x2][0x30]; /* 0x9b803840061c283a */
/*01b8*/ DADD R18, R16, R16; /* 0xe3800000081c404a */
/* 0x088880948880a080 */
/*01c8*/ DFMA R14, R10, R14, c[0x2][0x38]; /* 0x9b803840071c283a */
/*01d0*/ DADD R12, R12, c[0x2][0x50]; /* 0x638000400a1c3032 */
/*01d8*/ DFMA R16, R10, R14, c[0x2][0x40]; /* 0x9b803840081c2842 */
/*01e0*/ DFMA R14, R12, c[0x2][0x58], R8; /* 0x5b8020400b1c303a */
/*01e8*/ DFMA R4, -R4, R8, R18; /* 0xdb884800041c1012 */
/*01f0*/ DFMA R16, R10, R16, c[0x2][0x48]; /* 0x9b804040091c2842 */
/*01f8*/ DFMA R18, -R12, c[0x2][0x58], R14; /* 0x5b8838400b1c304a */
/* 0x08aca4a4a4a08094 */
/*0208*/ DMUL R4, R6, R4; /* 0xe4000000021c1812 */
/*0210*/ DMUL R10, R10, R16; /* 0xe4000000081c282a */
/*0218*/ DADD R18, -R8, R18; /* 0xe3880000091c204a */
/*0220*/ DFMA R4, R8, R10, R4; /* 0xdb801000051c2012 */
/*0228*/ DADD R18, R4, -R18; /* 0xe3810000091c104a */
/*0230*/ DFMA R12, R12, c[0x2][0x60], R18; /* 0x5b8048400c1c3032 */
/*0238*/ DADD R12, R14, R12; /* 0xe3800000061c3832 */
/* 0x080000000000b810 */
/*0248*/ ST.E.64 [R2], R12; /* 0xe5800000001c0830 */
/*0250*/ EXIT; /* 0x18000000001c003c */
/*0258*/ BRA 0x258; /* 0x12007ffffc1c003c */
/*0260*/ NOP; /* 0x85800000001c3c02 */
/*0268*/ NOP; /* 0x85800000001c3c02 */
/*0270*/ NOP; /* 0x85800000001c3c02 */
/*0278*/ NOP; /* 0x85800000001c3c02 */
......................
Fatbin ptx code:
================
arch = sm_35
code version = [6,2]
producer = cuda
host = linux
compile_size = 64bit
compressed
$
Прежде всего мы видим, что профилировщик сообщает об этой двойной точности log, используя в основном вычисления с двойной точностью, но также одну "специальную" операцию одинарной точности.
Глядя на дамп SASS, мы наблюдаем инструкцию с одинарной точностью:
/*0130*/ MUFU.RCP64H R9, R7; /* 0x84000000031c1c26 */
(мы можем подтвердить это с одинарной точностью, обратившись к документации )
Большая часть алгоритма использует плавающую точку двойной точности. Этот конкретный шаг может быть задействован в некоторой оценке, которая требует только одинарной точности.
Таким образом, «ожидается», что использование log может сообщить о некоторой (ненулевой) метрике flop_count_sp_special.
(Да, инструкции FCMP также имеют одинарную точность, но, по-видимому, они находятся в неиспользованном пути кода).