Large overhead due to device init code

[ndcontini]

I was profiling the performance of multiple zfp compressions/decompressions and noticed that the newer (unreleased) zfp code contains a "device init" function that is called whenever the execution mode is set for a zfp stream.

zfp/src/cuda/device.cuh

Lines 17 to 53 in cccbb9d

	// warm up device by launching dummy kernel
	bool device_init()
	{
	bool success = true;
	ErrorCheck error;
	// allocate device memory
	unsigned int* d_word = NULL;
	cudaMalloc(&d_word, sizeof(*d_word));
	success &= error.check("zfp device init - cudaMalloc");
	// launch a dummy kernel
	device_init_kernel<<<1, 1>>>(d_word);
	success &= error.check("zfp device init - kernel");
	// allocate host memory
	unsigned int* h_word = NULL;
	cudaMallocHost(&h_word, sizeof(*h_word));
	success &= error.check("zfp device init - cudaMallocHost");
	// copy from device to host
	if (success) {
	cudaMemcpy(h_word, d_word, sizeof(*h_word), cudaMemcpyDeviceToHost);
	success &= error.check("zfp device init - cudaMemcpy");
	success &= (*h_word == ZFP_MAGIC);
	}
	// free host memory
	cudaFreeHost(h_word);
	success &= error.check("zfp device init - cudaFreeHost");
	// free device memory
	cudaFree(d_word);
	success &= error.check("zfp device init - cudaFree");
	return success;
	}

called from here

zfp/src/cuda/interface.cu

Lines 20 to 42 in cccbb9d

	zfp_bool
	zfp_internal_cuda_init(zfp_exec_params_cuda* params)
	{
	// ensure GPU word size equals CPU word size
	if (sizeof(Word) != sizeof(bitstream_word))
	return false;
	// perform expensive query of device properties only once
	static bool initialized = false;
	static cudaDeviceProp prop;
	if (!initialized && cudaGetDeviceProperties(&prop, 0) != cudaSuccess)
	return zfp_false;
	initialized = true;
	// cache device properties
	params->processors = prop.multiProcessorCount;
	params->grid_size[0] = prop.maxGridSize[0];
	params->grid_size[1] = prop.maxGridSize[1];
	params->grid_size[2] = prop.maxGridSize[2];
	// launch device warm-up kernel
	return (zfp_bool)zfp::cuda::internal::device_init();
	}

Here is a list of events from a profile I took that shows two different compression calls:

For our use case, the data to be compressed is relatively small (8MB-32MB), so the compression call itself only runs for ~40us, while the other calls related to this device init code incur up to 100s of microseconds of latency. I was able to resolve this by only calling zfp::cuda::internal::device_init the first time zfp_internal_cuda_init is called. I think this is an acceptable approach, since I believe the device init function is just making sure that the application has correctly enabled the CUDA runtime and that the application is able to run the kernels. After device init is succeeds once, it is likely that it will succeed throughout the course of the application. Do I have a misconception about the purpose of this code or is my suggested code change valid?

[lindstro]

You are right that zfp::cuda::internal::device_init() should in principle be called only once per program execution. And I believe at one point we kept a "global" variable to indicate whether this function had (successfully) been called. We ultimately settled on a slightly modified design, for two reasons:

(1) There's no generic zfp_init() function that one might call only once and have it warm up the GPU. Even if there were, it's not clear that it should attempt to initialize CUDA for applications that do not use it. Also, we wanted to include this initialization in a function that would return an error code if it failed (e.g., if you have no GPU), and preferably in a function that serves to change/initialize state rather than deep in the bowels of zfp_compress(), which would provide little indication of why the call failed.

(2) We were moderately concerned about caching the result, including the queried CUDA device properties, in a global (or local static) variable without knowing if CUDA might allow such properties to change over time across an execution, or if there might be race conditions in applications that compress on many separate threads, or anything else that we had not thought of. This was perhaps a bit of laziness on our part, but it seemed reasonable to have each zfp_stream_set_execution() call ensure that the execution policy was in fact supported before switching to it.

Our expectation for most applications that use zfp CUDA compression is that they would allocate and initialize a zfp_stream once (including setting the execution policy), and that this object would then persist throughout the execution. (Consider a simulation that initializes a zfp_stream once and then generates and writes compressed data to disk once per minute or so.) I realize that your use case is quite different, and that there are cases when it may not be practical to create only a single persistent zfp_stream. I think your code changes, as I understand them, likely make sense, and I'd be open to any suggestions on how to handle this more cleanly and efficiently before we release the next zfp version.

[ndcontini]

The way I approached this (see below for actual code) was to use the existing static variable initialized in zfp_internal_cuda_init and wrap the entirety of the function in an if (!initialized) block to avoid rerunning the code upon calling the function again. The result of the first call is cached another static variable res. The assumption here is that if the function returns true the first time, it will almost always return true if it were run again. I make the same assumption for false as well. It's a pretty reasonable assumption considering the device init function is just testing to see if it can allocate device and host memory and run a kernel. I'm sure there are edge cases where this assumption is incorrect, but for my use case where I have to create many streams over the course of an application run, it is worth the performance benefit.

Technically, multiple threads could run this code simultaneously and init would be run twice. However each thread would write the exact same values, so I think ultimately no state would be corrupted as a result, i.e. this race condition is benign. That said, these functions could easily change in the future where this race condition would become malignant, so it is probably bad coding practice. I can't think of a way off the top of my head to make it safer outside of simply adding a mutex.

I'd be curious, is there any way to have a pool of zfp streams that I can reuse? For example, in the MPI use case, it is very likely we would determine that we will only every have n compressions/decompressions queued at once. Instead of having to create new streams and calling zfp_set_execution (which calls these init functions underneath) on each, is there a way to reuse zfp streams, changing the input and output buffers between operations? I suspect the answer is "no" as this doesn't seem how they are intended to be used. However if this were supported behavior, this would be a way to resolve the issue from the MVAPICH side without changing the zfp code.

diff --git a/src/cuda/interface.cu b/src/cuda/interface.cu
index d66fa27..89f19ea 100644
--- a/src/cuda/interface.cu
+++ b/src/cuda/interface.cu
@@ -20,16 +20,28 @@
 zfp_bool
 zfp_internal_cuda_init(zfp_exec_params_cuda* params)
 {
+  static bool initialized = false;
+  static zfp_bool res = zfp_false;
+  static cudaDeviceProp prop;
+
+  if (!initialized) {
+    initialized = zfp_true;
+
     // ensure GPU word size equals CPU word size
-  if (sizeof(Word) != sizeof(bitstream_word))
-    return false;
+    if (sizeof(Word) != sizeof(bitstream_word)) {
+        res = zfp_false;
+        return res;
+    }
 
     // perform expensive query of device properties only once
-  static bool initialized = false;
-  static cudaDeviceProp prop;
-  if (!initialized && cudaGetDeviceProperties(&prop, 0) != cudaSuccess)
-    return zfp_false;
-  initialized = true;
+    if (cudaGetDeviceProperties(&prop, 0) != cudaSuccess) {
+        res = zfp_false;
+        return res;
+    }
+
+    // launch device warm-up kernel
+    res = (zfp_bool)zfp::cuda::internal::device_init();
+  }
 
   // cache device properties
   params->processors = prop.multiProcessorCount;
@@ -37,8 +49,7 @@ zfp_internal_cuda_init(zfp_exec_params_cuda* params)
   params->grid_size[1] = prop.maxGridSize[1];
   params->grid_size[2] = prop.maxGridSize[2];
 
-  // launch device warm-up kernel
-  return (zfp_bool)zfp::cuda::internal::device_init();
+  return res;
 }
 
 size_t

[lindstro]

I should have looked more closely--I didn't realize we were still using this static initialized variable. So I guess what you're doing is as safe as what what we already are. :-)

I'd be curious, is there any way to have a pool of zfp streams that I can reuse?

We don't provide such a pool, but there's no reason you couldn't manage one yourself. Using an existing zfp_stream, you can update the compressed memory buffer it references as follows:

  stream = stream_open(buffer, bufsize);
  zfp_stream_set_bit_stream(zfp, stream);

You'll of course want to stream_close() any previously associated bitstream and deallocate unused memory buffers before calling zfp_stream_set_bit_stream() with a new memory buffer, but there's no reason you cannot reuse the zfp_stream object with a new memory buffer.

We have a related issue with thread-safe access to compressed-array objects, where we currently create thread-local zfp_stream objects on the fly:

zfp/include/zfp/codec/zfpcodec.hpp

Lines 146 to 158 in 28d12a0

	// encode full contiguous block
	size_t encode_block(bitstream_offset offset, const Scalar* block) const
	{
	if (thread_safety) {
	// make a thread-local copy of zfp stream and bit stream
	zfp_stream zfp = clone_stream();
	size_t size = encode_block(&zfp, offset, block);
	stream_close(zfp.stream);
	return size;
	}
	else
	return encode_block(stream, offset, block);
	}