ManagedCuda.CudaException: 'ErrorInvalidPtx: This indicates that a PTX JIT compilation failed.'

[Echostorm44]

Hi there. New to this and having trouble troubleshooting this error:

CUDA Info

PTX contents

`.version 7.5 .target sm_35 .address_size 64

.visible .entry LucasLehmerPrimalityTest( .param .u64 exponent, .param .u64 s, .param .u64 result) { .reg .u64 a; .reg .u64 b; .reg .u64 c;

mov.u64 a, s;
mov.u64 b, a;
add.u64 c, a, a;
add.u64 c, c, -2;

mov.u64 result, 0;

pfor (idx = 0; idx < exponent - 2; idx += blockDim.x * gridDim.x) {
    if (idx + threadIdx.x < exponent - 2) {
        mul.wide.u32 a, a, a;
        add.u64 a, a, -2;
        rem.u64 a, a, b;

        if (a == 0) {
            add.u64 result, result, 1;
        }
    }
}

}

.visible .entry MillerRabinPrimalityTest( .param .u64 n, .param .u32 iterations, .param .u64 d, .param .u32 result) { .reg .u64 a; .reg .u64 x; .reg .u32 i;

mov.u32 result, 1;

for (i = 0; i < iterations; i++) {
    a = (ulong) rand() % (n - 3) + 2;

    x = pow_mod(a, d, n);

    if (x == 1 || x == n - 1) {
        continue;
    }

    for (unsigned j = 0; j < d - 1; j++) {
        x = pow_mod(x, 2, n);

        if (x == 1) {
            result = 0;
            return;
        }

        if (x == n - 1) {
            break;
        }
    }

    if (x != n - 1) {
        result = 0;
        return;
    }
}

}

.visible .func (.param .u64 a, .param .u64 b, .param .u64 mod) pow_mod( .param .u64 a, .param .u64 b, .param .u64 mod) { .reg .u64 res; .reg .u64 mul;

mov.u64 res, 1;

pfor (idx = 0; idx < 64; idx += blockDim.x * gridDim.x) {
    if (idx + threadIdx.x < 64) {
        if ((b & (1 << idx)) != 0) {
            mul.wide.u32 res, res, a;
            rem.u64 res, res, mod;
        }

        mul.wide.u32 a, a, a;
        rem.u64 a, a, mod;
    }
}

ret.u64 res;

}`

Any Ideas?

[kunzmi]

Hi,

I guess that the ' at the beginning and end of your PTX are just formatting issues of github? If not, they shouldn't be there. The other thing I'd guess is the fact that your PTX is targeted for sm_35 which support was dropped with cuda 12. Not sure if they still should be able to compile the PTX though. I'd try to compile the original kernel source file with the same cuda toolkit version as the cuda runtime version you use (cuda 12.1) which is apperently not the case and set a supported minimum target version of sm_5.

Cheers, Michael

[kunzmi]

And wait, why do you have lines like for (unsigned j = 0; j < d - 1; j++) { in your PTX file? This is not valid PTX so the error message seams right - is this really compiled PTX from cuda source code?

[Echostorm44]

This is my bad, thanks!