//
// PTX fixture with a .global variable — used to test cuModuleGetGlobal.
//
// Equivalent CUDA C:
//   __device__ unsigned int g_counter;      // uninit; init via host memcpy
//   extern "C" __global__ void read_counter(unsigned int* out) {
//       *out = g_counter;
//   }
//
// We keep the global uninitialized because PTX literal-initializer syntax
// varies across toolchains; the host side writes a known value via
// `cuMemcpyHtoD` + `cuModuleGetGlobal` before launching the kernel.
//

.version 7.4
.target sm_50
.address_size 64

.visible .global .align 4 .u32 g_counter;

.visible .entry read_counter(
    .param .u64 read_counter_param_0
)
{
    .reg .b64 %rd<4>;
    .reg .b32 %r<2>;

    ld.param.u64       %rd1, [read_counter_param_0];
    cvta.to.global.u64 %rd2, %rd1;

    // &g_counter (constant address, resolved at module load)
    mov.u64            %rd3, g_counter;
    ld.global.u32      %r1, [%rd3];

    st.global.u32      [%rd2], %r1;
    ret;
}
