     cuStreamDestroy(stream);
 }
 
-// __global__ void kernel_twostructs(struct MyStruct *mystruct, struct MyStruct2 *mystruct2) {
-//     mystruct->p1[0] = 9.0f;
-//     mystruct->p2[0] = 10.0f;
-//     mystruct2->p1[0] = 11.0f;
-// }
-
-// void testtwostructs() {
-//     int N = 1024;
-
-//     CUstream stream;
-//     cuStreamCreate(&stream, 0);
-
-//     float *gpuFloats1;
-//     cudaMalloc((void**)(&gpuFloats1), N * sizeof(float));
-
-//     float *gpuFloats2;
-//     cudaMalloc((void**)(&gpuFloats2), N * sizeof(float));
-
-//     float *gpuFloats3;
-//     cudaMalloc((void**)(&gpuFloats3), N * sizeof(float));
-
-//     float *hostFloats1 = new float[N];
-//     float *hostFloats2 = new float[N];
-//     float *hostFloats3 = new float[N];
-
-//     struct MyStruct mystruct = {(float *)gpuFloats1, (float *)gpuFloats2};
-//     struct MyStruct2 mystruct2 = {(float *)gpuFloats3};
-
-//     kernel_twostructs<<<dim3(1,1,1), dim3(32,1,1), 0, stream>>>(&mystruct, &mystruct2);
-
-//     cudaMemcpy(hostFloats1, gpuFloats1, 4 * sizeof(float), cudaMemcpyDeviceToHost);
-//     cudaMemcpy(hostFloats2, gpuFloats2, 4 * sizeof(float), cudaMemcpyDeviceToHost);
-//     cudaMemcpy(hostFloats3, gpuFloats3, 4 * sizeof(float), cudaMemcpyDeviceToHost);
-
-//     cuStreamSynchronize(stream);
-
-//     cout << hostFloats1[0] << endl;
-//     cout << hostFloats2[0] << endl;
-//     cout << hostFloats3[0] << endl;
-
-//     assert(hostFloats1[0] == 9);
-//     assert(hostFloats2[0] == 10);
-//     assert(hostFloats3[0] == 11);
-
-//     cudaFree(gpuFloats1);
-//     cudaFree(gpuFloats2);
-//     cudaFree(gpuFloats3);
-
-//     delete[]hostFloats1;
-//     delete[]hostFloats2;
-//     delete[]hostFloats3;
-
-//     cuStreamDestroy(stream);
-// }
-
-struct Struct_1float {
-    float f1;
-};
+__global__ void kernel_twostructs(struct Struct_fp_fp_f_f mystruct, struct Struct_fp mystruct2) {
+    mystruct.p1[0] = 9.0f;
+    mystruct.p2[0] = 10.0f;
+    mystruct2.p1[0] = 11.0f;
+}
+
+void testtwostructs() {
+    int N = 1024;
+
+    CUstream stream;
+    cuStreamCreate(&stream, 0);
+
+    float *gpuFloats1;
+    cudaMalloc((void**)(&gpuFloats1), N * sizeof(float));
+
+    float *gpuFloats2;
+    cudaMalloc((void**)(&gpuFloats2), N * sizeof(float));
+
+    float *gpuFloats3;
+    cudaMalloc((void**)(&gpuFloats3), N * sizeof(float));
+
+    float *hostFloats1 = new float[N];
+    float *hostFloats2 = new float[N];
+    float *hostFloats3 = new float[N];
+
+    struct Struct_fp_fp_f_f mystruct = {(float *)gpuFloats1, (float *)gpuFloats2};
+    struct Struct_fp mystruct2 = {(float *)gpuFloats3};
+
+    kernel_twostructs<<<dim3(1,1,1), dim3(32,1,1), 0, stream>>>(mystruct, mystruct2);
+
+    cudaMemcpy(hostFloats1, gpuFloats1, 4 * sizeof(float), cudaMemcpyDeviceToHost);
+    cudaMemcpy(hostFloats2, gpuFloats2, 4 * sizeof(float), cudaMemcpyDeviceToHost);
+    cudaMemcpy(hostFloats3, gpuFloats3, 4 * sizeof(float), cudaMemcpyDeviceToHost);
+
+    cuStreamSynchronize(stream);
+
+    cout << hostFloats1[0] << endl;
+    cout << hostFloats2[0] << endl;
+    cout << hostFloats3[0] << endl;
+
+    assert(hostFloats1[0] == 9);
+    assert(hostFloats2[0] == 10);
+    assert(hostFloats3[0] == 11);
+
+    cudaFree(gpuFloats1);
+    cudaFree(gpuFloats2);
+    cudaFree(gpuFloats3);
+
+    delete[]hostFloats1;
+    delete[]hostFloats2;
+    delete[]hostFloats3;
+
+    cuStreamDestroy(stream);
+}
 
 // __global__ void kernel_structbyval_noptrs(struct Struct_1float mystruct1, float *out) {
 //     if(threadIdx.x == 0) {

    float *hostFloats3 = new float[N];
    float *hostOut = new float[N];

    struct MyStruct mystruct = {(float *)gpuFloats1, (float *)gpuFloats2, 3.0f, 8.0f};
    struct_byvalue<<<dim3(1,1,1), dim3(32,1,1), 0, stream>>>(mystruct, (float *)gpuOut);

    cudaMemcpy(hostFloats1, gpuFloats1, 4 * sizeof(float), cudaMemcpyDeviceToHost);

    cuStreamDestroy(stream);
}

__global__ void struct_aspointer(struct Struct_2floats *mystruct, float *out) {
    out[0] = mystruct->f1;
    out[1] = mystruct->f2;
