diff a/hat/backends/ffi/cuda/src/main/native/cpp/squares.cpp b/hat/backends/ffi/cuda/src/main/native/cpp/squares.cpp
--- /dev/null
+++ b/hat/backends/ffi/cuda/src/main/native/cpp/squares.cpp
@@ -0,0 +1,171 @@
+/*
+ * Copyright (c) 2024, Oracle and/or its affiliates. All rights reserved.
+ * DO NOT ALTER OR REMOVE COPYRIGHT NOTICES OR THIS FILE HEADER.
+ *
+ * This code is free software; you can redistribute it and/or modify it
+ * under the terms of the GNU General Public License version 2 only, as
+ * published by the Free Software Foundation.  Oracle designates this
+ * particular file as subject to the "Classpath" exception as provided
+ * by Oracle in the LICENSE file that accompanied this code.
+ *
+ * This code is distributed in the hope that it will be useful, but WITHOUT
+ * ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
+ * FITNESS FOR A PARTICULAR PURPOSE.  See the GNU General Public License
+ * version 2 for more details (a copy is included in the LICENSE file that
+ * accompanied this code).
+ *
+ * You should have received a copy of the GNU General Public License version
+ * 2 along with this work; if not, write to the Free Software Foundation,
+ * Inc., 51 Franklin St, Fifth Floor, Boston, MA 02110-1301 USA.
+ *
+ * Please contact Oracle, 500 Oracle Parkway, Redwood Shores, CA 94065 USA
+ * or visit www.oracle.com if you need additional information or have any
+ * questions.
+ */
+
+#include "cuda_backend.h"
+#include "../include/cuda_backend.h"
+
+class KernelContextWithBufferState{
+public:
+int x;
+int maxX;
+BufferState bufferState;
+};
+struct ArgArray_2 {
+    int argc;
+    u8_t pad12[12];
+    KernelArg argv[2];
+};
+
+
+struct S32Array1024WithBufferState {
+    int length;
+    int array[1024];
+    BufferState bufferState;
+};
+int main(int argc, char **argv) {
+    CudaBackend cudaBackend(0
+            | Backend::Config::Config::INFO_BIT
+            | Backend::Config::Config::TRACE_CALLS_BIT
+            | Backend::Config::Config::TRACE_COPIES_BIT
+    );
+    CudaSource* cudaSource = new CudaSource((char *) R"(
+       #define NDRANGE_CUDA
+       #define __global
+       typedef char s8_t;
+       typedef char byte;
+       typedef char boolean;
+       typedef unsigned char u8_t;
+       typedef short s16_t;
+       typedef unsigned short u16_t;
+       typedef unsigned int u32_t;
+       typedef int s32_t;
+       typedef float f32_t;
+       typedef long s64_t;
+       typedef unsigned long u64_t;
+       typedef struct KernelContext_s{
+         int x;
+         int maxX;
+       }KernelContext_t;
+       typedef struct S32Array_s{
+          int length;
+          int array[1];
+       }S32Array_t;
+
+       extern "C" __device__ inline int squareit(
+           int v
+       ){
+          return v*v;
+       }
+
+       extern "C" __global__ void squareKernel(
+           KernelContext_t *global_kc,  S32Array_t* s32Array
+       ){
+         KernelContext_t mine;
+         KernelContext_t* kc=&mine;
+         kc->x=blockIdx.x*blockDim.x+threadIdx.x;
+         kc->maxX=global_kc->maxX;
+         if(kc->x<kc->maxX){
+             int value = s32Array->array[(long)kc->x];
+             s32Array->array[(long)kc->x]=squareit(value);
+         }
+         return;
+        }
+    )");
+
+    PtxSource *ptxSource = new PtxSource((char *)(R"(
+       // Generated by NVIDIA NVVM Compiler
+       //
+       // Compiler Build ID: CL-33191640
+       // Cuda compilation tools, release 12.2, V12.2.140
+       // Based on NVVM 7.0.1
+       //
+
+       .version 8.2
+       .target sm_52
+       .address_size 64
+
+               // .globl        squareKernel
+
+       .visible .entry squareKernel(
+               .param .u64 squareKernel_param_0,
+               .param .u64 squareKernel_param_1
+       )
+       {
+               .reg .pred           %p<2>;
+               .reg .b32            %r<8>;
+               .reg .b64            %rd<7>;
+
+
+               ld.param.u64         %rd2, [squareKernel_param_0];
+               ld.param.u64         %rd1, [squareKernel_param_1];
+               cvta.to.global.u64         %rd3, %rd2;
+               mov.u32              %r2, %ntid.x;
+               mov.u32              %r3, %ctaid.x;
+               mov.u32              %r4, %tid.x;
+               mad.lo.s32           %r1, %r3, %r2, %r4;
+               ld.global.u32        %r5, [%rd3+4];
+               setp.ge.s32          %p1, %r1, %r5;
+               @%p1 bra             $L__BB0_2;
+
+               cvta.to.global.u64   %rd4, %rd1;
+               mul.wide.s32         %rd5, %r1, 4;
+               add.s64              %rd6, %rd4, %rd5;
+               ld.global.u32        %r6, [%rd6+4];
+               mul.lo.s32           %r7, %r6, %r6;
+               st.global.u32        [%rd6+4], %r7;
+
+       $L__BB0_2:
+               ret;
+       }
+    )"));
+    const int maxX = 32;
+
+    bool useCuda=false;
+    std::cerr<<"using " << (useCuda?"CUDA":"PTX")<<std::endl;
+    auto *module =useCuda?cudaBackend.compile(cudaSource):cudaBackend.compile(ptxSource);
+   //auto *module =cudaBackend.compile(ptxSource);
+
+     auto  *kernelContextWithBufferState = bufferOf<KernelContextWithBufferState>("kernelcontext");
+    kernelContextWithBufferState->x=0;
+    kernelContextWithBufferState->maxX=maxX;
+    auto *pS32Array1024WithBufferState = bufferOf<S32Array1024WithBufferState>("s32Arrayx1024");
+    pS32Array1024WithBufferState->length=maxX;
+    for (int i=0; i<pS32Array1024WithBufferState->length; i++){
+        pS32Array1024WithBufferState->array[i]=i;
+    }
+
+    ArgArray_2 args2Array{.argc = 2, .argv={
+            {.idx = 0, .variant = '&',.value = {.buffer ={.memorySegment = static_cast<void *>(kernelContextWithBufferState), .sizeInBytes = sizeof(KernelContextWithBufferState), .access = RO_BYTE}}},
+            {.idx = 1, .variant = '&',.value = {.buffer ={.memorySegment = static_cast<void *>(pS32Array1024WithBufferState), .sizeInBytes = sizeof(S32Array1024WithBufferState), .access = RW_BYTE}}}
+    }};
+    const auto kernel = module->getCudaKernel((char*)"squareKernel");
+    std::cout << kernel->name <<std::endl;
+    kernel->ndrange( &args2Array);
+    for (int i=0; i<pS32Array1024WithBufferState->length; i++){
+        const int sq = pS32Array1024WithBufferState->array[i];
+        std::cout << i << " sq="<<sq <<std::endl;
+    }
+}
+
