cuobjdump 不发出 PTX 算术指令

Question

为什么 cuobjdump 不发出下面的 PTX mul 指令？ nvcc 是否优化了 cubin 输出本身？结果是编译时计算的吗？如果是这样，对于这个最简单的情况，nvcc可以合理地进一步优化输出，根本不需要在设备端生成任何指令。

mul.cu

#include <cuda_runtime.h>
#include <stdio.h>

__global__ void mul(float *res) {
    float x = 11.1, y = 22.2;
    *res = x * y;
}

int main() {
    float *res;
    cudaMallocManaged(&res, sizeof(float));
    mul<<<1, 1>>>(res);
    cudaDeviceSynchronize();
    printf("11.1 * 22.2 = %f\n", *res);
}

问题

$ nvcc mul.cu -o mul

$ ./mul

11.1 * 22.2 = 246.420013

$ cuobjdump -fun mul -ptx ./mul

...

.visible .entry _Z3mulPf(
.param .u64 _Z3mulPf_param_0
)
{
.reg .b32 %r<2>;
.reg .b64 %rd<3>;


ld.param.u64 %rd1, [_Z3mulPf_param_0];
cvta.to.global.u64 %rd2, %rd1;
mov.u32 %r1, 1131834246;
st.global.u32 [%rd2], %r1;
ret;

}

Answer 1

结果是编译时计算的吗？

是的。

编译器可以观察到结果始终为 11.1x22.2，因此它只是将该值（当

float

位模式视为十进制整数时：1131834246）放入结果位置。

如果您想查看 mul 指令，请将乘法输入值设置为内核参数：

#include <cuda_runtime.h>
#include <stdio.h>

__global__ void mul(float *res, float x, float y) {
    *res = x * y;
}

int main() {
    float *res;
    cudaMallocManaged(&res, sizeof(float));
    mul<<<1, 1>>>(res, 11.1, 22.2);
    cudaDeviceSynchronize();
    printf("11.1 * 22.2 = %f\n", *res);
}