代码之家 › 专栏 › 技术社区 › sof

cuobjdump不发出PTX算术指令

ptx cuda

0

sof · 技术社区 · 11 月前

为什么不 cuobjdump 发射PTX mul 下面的指令?有 nvcc 优化了 立方英寸 输出自己?结果是在编译时计算的吗?如果是这样,对于这个最简单的情况 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;

}

1 回复 | 直到 11 月前

1

Robert Crovella 11 月前

结果是在编译时计算的吗?

对。

编译器可以观察到结果将始终为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);
}

如果是这样,对于这种最简单的情况,nvcc可以合理地进一步优化输出,而不会在设备端生成任何指令。

任何适当的优化仍然必须在全局状态下具有相同的结果。因此,在这种情况下,用一个 cudaMemcpy 类型操作(或者,由于它是托管内存,可能只是一个 memcpy 或者其他一些内存设置操作),但我认为编译器从未尝试过这种优化。