Hipify latest CUDA updates

Author: lindstro

src/hip/decode1.h

@@ -132,7 +132,7 @@ decode1(
 #endif
 
   // launch GPU kernel
-  hipLaunchKernelGGL(HIP_KERNEL_NAME(decode1_kernel<Scalar>), grid_size, block_size, 0, 0, 
+  decode1_kernel<Scalar><<<grid_size, block_size>>>(
     d_data,
     size[0],
     stride[0],

src/hip/decode2.h

@@ -143,7 +143,7 @@ decode2(
 #endif
 
   // launch GPU kernel
-  hipLaunchKernelGGL(HIP_KERNEL_NAME(decode2_kernel<Scalar>), grid_size, block_size, 0, 0, 
+  decode2_kernel<Scalar><<<grid_size, block_size>>>(
     d_data,
     make_size2(size[0], size[1]),
     make_ptrdiff2(stride[0], stride[1]),

src/hip/decode3.h

@@ -153,7 +153,7 @@ decode3(
 #endif
 
   // launch GPU kernel
-  hipLaunchKernelGGL(HIP_KERNEL_NAME(decode3_kernel<Scalar>), grid_size, block_size, 0, 0, 
+  decode3_kernel<Scalar><<<grid_size, block_size>>>(
     d_data,
     make_size3(size[0], size[1], size[2]),
     make_ptrdiff3(stride[0], stride[1], stride[2]),

src/hip/device.h

@@ -27,7 +27,7 @@ bool device_init()
   success &= error.check("zfp device init - hipMalloc");
 
   // launch a dummy kernel
-  hipLaunchKernelGGL(device_init_kernel, 1, 1, 0, 0, d_word);
+  device_init_kernel<<<1, 1>>>(d_word);
   success &= error.check("zfp device init - kernel");
 
   // allocate host memory
@@ -194,25 +194,22 @@ Word* setup_device_index_decompress(zfp_stream* stream)
   return d_index;
 }
 
-bool setup_device_compact(size_t* chunk_size, unsigned long long** d_offsets, size_t* lcubtemp, void** d_cubtemp, uint processors)
+bool setup_device_compact(size_t* chunk_size, unsigned long long** d_offset, size_t* cubtmp_size, void** d_cubtmp, uint processors)
 {
+  // use 1K threads per SM for high occupancy (assumes one thread per zfp block)
   const size_t threads_per_sm = 1024;
-  // Assuming 1 thread = 1 ZFP block,
-  // launching 1024 threads per SM should give a decent occupancy
   *chunk_size = processors * threads_per_sm;
-  size_t size = (*chunk_size + 1) * sizeof(unsigned long long);
+
   // allocate and zero-initialize offsets
-  if (!device_calloc(d_offsets, size, "offsets"))
+  const size_t size = (*chunk_size + 1) * sizeof(unsigned long long);
+  if (!device_calloc(d_offset, size, "offsets"))
     return false;
 
-  // TODO : error handling for CUB
-  // Using CUB for the prefix sum. CUB needs a bit of temp memory too
-  size_t tempsize;
-  hipcub::DeviceScan::InclusiveSum(nullptr, tempsize, *d_offsets, *d_offsets, *chunk_size + 1);
-  *lcubtemp = tempsize;
-  if (!device_malloc(d_cubtemp, tempsize, "offsets")) {
-    device_free(d_offsets);
-    *d_offsets = NULL;
+  // allocate temporary memory for CUB prefix sum
+  if (hipcub::DeviceScan::InclusiveSum(nullptr, *cubtmp_size, *d_offset, *d_offset, *chunk_size + 1) != hipSuccess ||
+      !device_malloc(d_cubtmp, *cubtmp_size, "offsets")) {
+    device_free(d_offset);
+    *d_offset = NULL;
     return false;
   }
 
@@ -268,8 +265,7 @@ void* setup_device_field_decompress(const zfp_field* field, void*& d_begin)
 }
 
 // copy from device to host (if needed) and deallocate device memory
-// TODO: d_begin should be first argument, with begin = NULL as default
-void cleanup_device(void* begin, void* d_begin, size_t bytes = 0)
+void cleanup_device(void* d_begin, void* begin = 0, size_t bytes = 0)
 {
   if (d_begin != begin) {
     // copy data from device to host and free device memory

src/hip/encode1.h

@@ -111,7 +111,7 @@ encode1(
 #endif
 
   // launch GPU kernel
-  hipLaunchKernelGGL(HIP_KERNEL_NAME(encode1_kernel<Scalar>), grid_size, block_size, 0, 0, 
+  encode1_kernel<Scalar><<<grid_size, block_size>>>(
     d_data,
     size[0],
     stride[0],

src/hip/encode2.h

@@ -126,7 +126,7 @@ encode2(
 #endif
 
   // launch GPU kernel
-  hipLaunchKernelGGL(HIP_KERNEL_NAME(encode2_kernel<Scalar>), grid_size, block_size, 0, 0, 
+  encode2_kernel<Scalar><<<grid_size, block_size>>>(
     d_data,
     make_size2(size[0], size[1]),
     make_ptrdiff2(stride[0], stride[1]),

src/hip/encode3.h

@@ -73,7 +73,7 @@ encode3_kernel(
   if (block_idx >= blocks)
     return;
 
-  // logical position in 2d array
+  // logical position in 3d array
   size_t pos = block_idx;
   const ptrdiff_t x = (pos % bx) * 4; pos /= bx;
   const ptrdiff_t y = (pos % by) * 4; pos /= by;
@@ -139,7 +139,7 @@ encode3(
 #endif
 
   // launch GPU kernel
-  hipLaunchKernelGGL(HIP_KERNEL_NAME(encode3_kernel<Scalar>), grid_size, block_size, 0, 0, 
+  encode3_kernel<Scalar><<<grid_size, block_size>>>(
     d_data,
     make_size3(size[0], size[1], size[2]),
     make_ptrdiff3(stride[0], stride[1], stride[2]),

src/hip/interface.cpp

@@ -119,12 +119,12 @@ zfp_internal_hip_compress(zfp_stream* stream, const zfp_field* field)
   if (d_index) {
     const size_t size = zfp_field_blocks(field) * sizeof(ushort);
     // TODO: assumes index stores block sizes
-    zfp::hip::internal::cleanup_device(stream->index ? stream->index->data : NULL, d_index, size);
+    zfp::hip::internal::cleanup_device(d_index, stream->index ? stream->index->data : NULL, size);
   }
 
   // copy stream from device to host if needed and free temporary buffers
-  zfp::hip::internal::cleanup_device(stream->stream->begin, d_stream, stream_bytes);
-  zfp::hip::internal::cleanup_device(zfp_field_begin(field), d_begin);
+  zfp::hip::internal::cleanup_device(d_stream, stream->stream->begin, stream_bytes);
+  zfp::hip::internal::cleanup_device(d_begin, zfp_field_begin(field));
 
   // update bit stream to point just past produced data
   if (bits_written)
@@ -220,10 +220,10 @@ zfp_internal_hip_decompress(zfp_stream* stream, zfp_field* field)
 
   // copy field from device to host if needed and free temporary buffers
   size_t field_bytes = zfp_field_size_bytes(field);
-  zfp::hip::internal::cleanup_device(zfp_field_begin(field), d_begin, field_bytes);
-  zfp::hip::internal::cleanup_device(stream->stream->begin, d_stream);
+  zfp::hip::internal::cleanup_device(d_begin, zfp_field_begin(field), field_bytes);
+  zfp::hip::internal::cleanup_device(d_stream, stream->stream->begin);
   if (d_index)
-    zfp::hip::internal::cleanup_device(stream->index->data, d_index);
+    zfp::hip::internal::cleanup_device(d_index, stream->index->data);
 
   // update bit stream to point just past consumed data
   if (bits_read)