Weird behavior of jagged_vector_buffer with zero capacity

[beomki-yeo]

I have encountered weird behavior of jagged_vector_buffer when the input argument include the zero capacity.

Everything is fine when we use the device resource and host(managed) resource together:
vecmem::data::jagged_vector_buffer<int> input_data_device({0, 0, 0, 0, 0, 0 }, {0, 1, 200, 1, 100, 2}, device_resource, &managed_resource);

But when I hook the buffer only with managed memory resource, the zero capacity of {0, 1, 200, 1, 100, 2} caused weird output.
vecmem::data::jagged_vector_buffer<int> input_data_device({0, 0, 0, 0, 0, 0 }, {0, 1, 200, 1, 100, 2}, manged_resource);

Actually I am not even sure if it is legal to use only the managed resource for initiating the jagged_vector_buffer, So anybody please confirm that if it is the unwanted usage of jagged_vector_buffer

Following is the codes that I tested in vecmem test directory

test_cuda_jagged_vector_view.cpp

/// Test filling a resizable jagged vector                                            
TEST_F(cuda_jagged_vector_view_test, zero_capacity) {
    vecmem::cuda::device_memory_resource device_resource;
    // Helper object for performing memory copies.                                    
    vecmem::cuda::copy copy;

    // Create the output data on the device.                                          
    vecmem::data::jagged_vector_buffer<int> input_data_device(
            {0, 0, 0, 0, 0, 0 }, {0, 1, 200, 1, 100, 2}, m_mem);
    copy.setup(input_data_device);

    // Run the filtering.                                                             
    test(input_data_device);


    for (size_t i=0 ; i<input_data_device.m_size; i++){
	printf("%d \n", input_data_device.host_ptr()[i].size());
    }
}

test_cuda_jagged_vector_view_kernels.cu

/// Kernel performing a linear transformation using the vector helper types           
__global__ void testKernel(vecmem::data::jagged_vector_view<int> input) {

    vecmem::jagged_device_vector<int> inputvec(input);

    while (1){
	if (inputvec[blockIdx.x].size() < inputvec[blockIdx.x].capacity()){
            auto subvec = inputvec.at(blockIdx.x);
            subvec.push_back(1);
        }
        else {
            break;
	}
    }
}

void test(vecmem::data::jagged_vector_view<int>& input){
    // Launch the kernel                                                              
    testKernel<<< input.m_size, 1 >>>(input);

    // Check whether it succeeded to run.                                             
    VECMEM_CUDA_ERROR_CHECK(cudaGetLastError());
    VECMEM_CUDA_ERROR_CHECK(cudaDeviceSynchronize());
}

Expected print output

0
1
200
1
100
2

Actual print output

7
1
200
1
100
2

[krasznaa]

This is very interesting! Though not exactly for the reason that I first thought...

You are right, technically your code should work. Even if it is not recommended to write code like this. Since (jagged) vector buffers should not be interacted with directly. You should always create a host or device vector "on top of" vector buffers. This is what I did in #130. And as long as I do that, things do work correctly.

What happens in your code is that after updating the memory atomically on the device in "managed" memory, the host code tries to access that managed memory through a "system of pointers". And it seems that the just-in-time-copy that would be necessary to get the most recent version of the managed memory back from the device to the host, is not triggering correctly. Very strangely I observed the same weird 7 value that you have. Even though I was convinced that this should be a random value. 😕

When using vecmem::copy to fill a vecmem::jagged_vector with the contents of the buffer, the copy from the managed memory is triggered explicitly through the CUDA runtime API. And that seems to work correctly. But the implicit copy does seem to fail. 😕

For now I would suggest that you stick to only using the contents of buffers though host/device vectors. In such a setup the code should work correctly... As I wrote in #130, I'll do some further tests with this. As I'll want to see if my theory is indeed correct, and with HIP and SYCL your code would actually work. 🤔

[krasznaa]

😕 I don't know what to believe anymore...

Running the equivalent test using HIP on an AMD card, the vecmem::jagged_vector_buffer can be interrogated successfully, and it returns the expected values. So at this point I was pretty convinced that we indeed found a CUDA runtime error.

But now I implemented the test for SYCL as well, and just for the first element of the buffer the size differs there as well. 😵 While running on the "host backend"! (However vecmem::sycl::copy can fill a host-side vecmem::jagged_vector in this example just fine as well.)

So now I'm really confused. The code seems to work as expected with HIP, when running on an actual GPU. But when using just the host memory with SYCL there is a very similar problem to what CUDA exhibits. Am I really losing my mind here? 😢

[krasznaa]

Aaaaaargh!!!

I understand it finally!

https://github.com/acts-project/vecmem/blob/main/core/include/vecmem/utils/impl/copy.ipp#L137

I was trying to be super aggressive with minimising memory operations. 😛 Since the first "internal vector" in the jagged vector has a zero capacity, the vecmem::copy::setup(...) function does not bother zeroing the "atomic size" of that vector. Since in no practical terms should it matter what the size of the vector is. Its capacity is zero after all. So the first "internal vector" in the jagged vector points at a place in memory that is never explicitly initialised during the job to any value. (And in fact no value is ever written to it.)

My guess is that HIP implicitly zeroes out its memory allocations. Or maybe that just happened by chance in my tests.

What I don't manage to understand is how the hell I also saw a size of 7 for the first internal vector in my tests with CUDA. 😕 With SYCL it was a different value. But with CUDA I also consistently got this "magic" 7 value. Even after a restart of the machine that I used for the tests. 😕

So... long story short, you should not access the contents of the (jagged) vector buffers directly. As interpreting these values requires some amount of logic. One that is present in vecmem::copy and the device vector classes.

I'll finish setting up the new tests in #130 after lunch, but at this point I would consider this behaviour to be a feature, not a bug. 😛