How to correctly pass a struct as a parameter to a kernel

[ChenxiZhou0619]

Hi, I'm trying to pass a struct parameter to a kernel function, the struct is like

struct MyStruct {
  hiprtInt2 xy;
  int       z;
};

the kernel is like

extern "C" __global__ void
ReturnConstantInt(hiprtGeometry geom, uint8_t* pixels, MyStruct myStruct) {
  const uint32_t x     = blockIdx.x * blockDim.x + threadIdx.x;
  const uint32_t y     = blockIdx.y * blockDim.y + threadIdx.y;
  const uint32_t index = x + y * 512;

  pixels[index * 4 + 0] = (uint8_t)(myStruct.xy.x);
  pixels[index * 4 + 1] = (uint8_t)(myStruct.xy.x);
  pixels[index * 4 + 2] = (uint8_t)(myStruct.xy.x);
  pixels[index * 4 + 3] = 255;
}

then i initialize the struct and launch the kernel in main function

  oroFunction func;
  hiprtApiFunction         kernel_out;

 CHECK_HIPRT(hiprtutils::build_trace_kernel(ctxt,
                                             "kernel_path",
                                             "ReturnConstantInt",
                                             kernel_out,
                                             additional_includes,
                                             kernel_options, // empty vector
                                             0,
                                             1,
                                             false));

  func = reinterpret_cast<oroFunction>(kernel_out);

// ...

 MyStruct myStruct;
  myStruct.xy = {50, 150};
  myStruct.z  = 250;

  CHECK_ORO(oroMalloc(reinterpret_cast<oroDeviceptr*>(&pixels), 512 * 512 * 4));

  void* args[] = {&geom, &pixels, &myStruct};
  launchKernel(func, 512, 512, args);
  writeImage("test.png", 512, 512, pixels);

the result is

when I change MyStruct to

struct MyStruct {
  int       z;
  hiprtInt2 xy;
};

the result image is

I suspect this is due to the alignment issue or something else I haven't take care, where can I find more hiprt programming guide?

[meistdan]

Hi, I don't see any obvious issue with the alignment. I suspect that the kernel was not recompiled. So, the host code uses the new layout, but the kernel still the previous one. Could you try to delete files in the cache directory?