|
| 1 | +// CUDA inline-PTX variant of the m041-or-shifted-square-fold ptxas reproducer. |
| 2 | +// |
| 3 | +// Build this same CUDA source twice and compare the printed output from the -O0 and -O2 binaries: |
| 4 | +// |
| 5 | +// nvcc -std=c++17 -O2 -arch=sm_103 -Xptxas -O0 \ |
| 6 | +// repro_nvcc_inline_ptx.cu -o repro_nvcc_inline_ptx_o0 |
| 7 | +// |
| 8 | +// nvcc -std=c++17 -O2 -arch=sm_103 -Xptxas -O2 \ |
| 9 | +// repro_nvcc_inline_ptx.cu -o repro_nvcc_inline_ptx_o2 |
| 10 | +// |
| 11 | +// Verified on 2026-05-15 with CUDA Toolkit 13.2.1 nvcc/ptxas |
| 12 | +// (`release 13.2, V13.2.78`, build `cuda_13.2.r13.2/compiler.37668154_0`). |
| 13 | + |
| 14 | +#include <cuda_runtime.h> |
| 15 | + |
| 16 | +#include <cstdint> |
| 17 | +#include <cstdio> |
| 18 | +#include <cstdlib> |
| 19 | + |
| 20 | +constexpr int kThreads = 32; |
| 21 | +constexpr int kInputWords = 32; |
| 22 | +constexpr int kOutputWords = 128; |
| 23 | +constexpr uint32_t kN = 32u; |
| 24 | +constexpr uint32_t kX = 0x00000000u; |
| 25 | +constexpr uint32_t kInput0 = 0x691a38a3u; |
| 26 | +constexpr uint32_t kSentinel = 0xa5a5a5a5u; |
| 27 | + |
| 28 | +static void check(cudaError_t err, const char* what) { |
| 29 | + if (err != cudaSuccess) { |
| 30 | + std::fprintf(stderr, "%s: %s\n", what, cudaGetErrorString(err)); |
| 31 | + std::exit(2); |
| 32 | + } |
| 33 | +} |
| 34 | + |
| 35 | +__global__ void repro_kernel(const uint32_t* in, uint32_t* out, uint32_t n, uint32_t x) { |
| 36 | + asm volatile( |
| 37 | + "{\n\t" |
| 38 | + "\n\t" |
| 39 | + ".reg .b32 r<22>;\n\t" |
| 40 | + ".reg .b64 rd<8>;\n\t" |
| 41 | + "\n\t" |
| 42 | + "mov.u64 rd0, %0;\n\t" |
| 43 | + "mov.u64 rd1, %1;\n\t" |
| 44 | + "mov.u32 r20, %%tid.x;\n\t" |
| 45 | + "cvta.to.global.u64 rd2, rd0;\n\t" |
| 46 | + "mul.wide.u32 rd3, r20, 4;\n\t" |
| 47 | + "add.s64 rd2, rd2, rd3;\n\t" |
| 48 | + "ld.global.u32 r2, [rd2];\n\t" |
| 49 | + "\n\t" |
| 50 | + "mov.u32 r9, 33554432;\n\t" |
| 51 | + "mul.lo.u32 r0, r2, r9;\n\t" |
| 52 | + "and.b32 r1, 4278255360, r20;\n\t" |
| 53 | + "mad.lo.u32 r3, r0, 16, r1;\n\t" |
| 54 | + "sub.u32 r1, 620912, r3;\n\t" |
| 55 | + "mul.lo.u32 r19, 524288, r1;\n\t" |
| 56 | + "mul.lo.u32 r0, r19, r19;\n\t" |
| 57 | + "or.b32 r0, r0, 55244;\n\t" |
| 58 | + "\n\t" |
| 59 | + "exit:\n\t" |
| 60 | + "cvta.to.global.u64 rd4, rd1;\n\t" |
| 61 | + "mul.wide.u32 rd5, r20, 16;\n\t" |
| 62 | + "add.s64 rd4, rd4, rd5;\n\t" |
| 63 | + "st.global.u32 [rd4 + 0], r0;\n\t" |
| 64 | + "}\n" |
| 65 | + : |
| 66 | + : "l"(in), "l"(out), "r"(n), "r"(x) |
| 67 | + : "memory"); |
| 68 | +} |
| 69 | + |
| 70 | +static uint64_t fnv1a(const uint32_t* words, int n) { |
| 71 | + uint64_t h = 1469598103934665603ull; |
| 72 | + for (int i = 0; i < n; ++i) { |
| 73 | + uint32_t v = words[i]; |
| 74 | + for (int b = 0; b < 4; ++b) { |
| 75 | + h ^= static_cast<unsigned char>(v >> (8 * b)); |
| 76 | + h *= 1099511628211ull; |
| 77 | + } |
| 78 | + } |
| 79 | + return h; |
| 80 | +} |
| 81 | + |
| 82 | +int main() { |
| 83 | + uint32_t h_in[kInputWords] = {}; |
| 84 | + uint32_t h_out[kOutputWords]; |
| 85 | + for (int i = 0; i < kOutputWords; ++i) { |
| 86 | + h_out[i] = kSentinel; |
| 87 | + } |
| 88 | + h_in[0] = kInput0; |
| 89 | + |
| 90 | + uint32_t* d_in = nullptr; |
| 91 | + uint32_t* d_out = nullptr; |
| 92 | + check(cudaMalloc(&d_in, sizeof(h_in)), "cudaMalloc input"); |
| 93 | + check(cudaMalloc(&d_out, sizeof(h_out)), "cudaMalloc output"); |
| 94 | + check(cudaMemcpy(d_in, h_in, sizeof(h_in), cudaMemcpyHostToDevice), "cudaMemcpy input"); |
| 95 | + check(cudaMemcpy(d_out, h_out, sizeof(h_out), cudaMemcpyHostToDevice), "cudaMemcpy output sentinel"); |
| 96 | + |
| 97 | + repro_kernel<<<1, kThreads>>>(d_in, d_out, kN, kX); |
| 98 | + check(cudaGetLastError(), "repro_kernel launch"); |
| 99 | + check(cudaDeviceSynchronize(), "cudaDeviceSynchronize"); |
| 100 | + check(cudaMemcpy(h_out, d_out, sizeof(h_out), cudaMemcpyDeviceToHost), "cudaMemcpy output"); |
| 101 | + check(cudaFree(d_out), "cudaFree output"); |
| 102 | + check(cudaFree(d_in), "cudaFree input"); |
| 103 | + |
| 104 | + std::printf("threads=%d n=%u x=0x%08x input0=0x%08x\n", kThreads, kN, kX, kInput0); |
| 105 | + bool any = false; |
| 106 | + for (int i = 0; i < kOutputWords; ++i) { |
| 107 | + if (h_out[i] != kSentinel) { |
| 108 | + any = true; |
| 109 | + std::printf("out[%d]=0x%08x\n", i, h_out[i]); |
| 110 | + } |
| 111 | + } |
| 112 | + if (!any) { |
| 113 | + std::printf("no output words changed\n"); |
| 114 | + } |
| 115 | + std::printf("hash=0x%016llx\n", static_cast<unsigned long long>(fnv1a(h_out, kOutputWords))); |
| 116 | + return 0; |
| 117 | +} |
0 commit comments