Add support for CUDA 13.0

* Add initial support for CUDA 13

* More fixes for CUDA 13

* Fix comment

Signed-off-by: Matthew Cong <[email protected]>

Author: matthewdcong

nanovdb/nanovdb/cuda/UnifiedBuffer.h

@@ -62,8 +62,8 @@ class UnifiedBuffer
     {
         assert(mSize <= mCapacity);
         cudaCheck(cudaMallocManaged(&mPtr, mCapacity, cudaMemAttachGlobal));
-        cudaCheck(cudaMemAdvise(mPtr, size, cudaMemAdviseSetPreferredLocation, device));
-        cudaCheck(cudaMemPrefetchAsync(mPtr, size, device, stream));
+        cudaCheck(util::cuda::memAdvise(mPtr, size, cudaMemAdviseSetPreferredLocation, device));
+        cudaCheck(util::cuda::memPrefetchAsync(mPtr, size, device, stream));
     }
 
     /// @brief Constructor with a specified device
@@ -113,8 +113,8 @@ class UnifiedBuffer
     {
         const size_t capacity = (reference && reference->capacity()) ? reference->capacity() : size;
         UnifiedBuffer buffer(size, capacity);
-        cudaCheck(cudaMemAdvise(buffer.mPtr, size, cudaMemAdviseSetPreferredLocation, device));
-        cudaCheck(cudaMemPrefetchAsync(buffer.mPtr, size, device, stream));
+        cudaCheck(util::cuda::memAdvise(buffer.mPtr, size, cudaMemAdviseSetPreferredLocation, device));
+        cudaCheck(util::cuda::memPrefetchAsync(buffer.mPtr, size, device, stream));
         return buffer;
     }
 
@@ -184,7 +184,7 @@ class UnifiedBuffer
         } else {
             void *ptr = 0;
             cudaCheck(cudaMallocManaged(&ptr, size, cudaMemAttachGlobal));
-            if (dev > -2) for (auto a : list) cudaCheck(cudaMemAdvise(ptr, size, a, dev));
+            if (dev > -2) for (auto a : list) cudaCheck(util::cuda::memAdvise(ptr, size, a, dev));
             if (mSize > 0) {// copy over data from the old memory block
                 cudaCheck(cudaMemcpy(ptr, mPtr, std::min(mSize, size), cudaMemcpyDefault));
                 cudaCheck(cudaFree(mPtr));
@@ -201,7 +201,7 @@ class UnifiedBuffer
     /// @param adv advice to be applied to the resized range
     void advise(ptrdiff_t byteOffset, size_t size, int dev, cudaMemoryAdvise adv) const
     {
-        cudaCheck(cudaMemAdvise(util::PtrAdd(mPtr, byteOffset), size, adv, dev));
+        cudaCheck(util::cuda::memAdvise(util::PtrAdd(mPtr, byteOffset), size, adv, dev));
     }
 
     /// @brief Apply a list of advices to a memory block
@@ -212,7 +212,7 @@ class UnifiedBuffer
     void advise(ptrdiff_t byteOffset, size_t size, int dev, std::initializer_list<cudaMemoryAdvise> list) const
     {
         void *ptr = util::PtrAdd(mPtr, byteOffset);
-        for (auto a : list)  cudaCheck(cudaMemAdvise(ptr, size, a, dev));
+        for (auto a : list)  cudaCheck(util::cuda::memAdvise(ptr, size, a, dev));
     }
 
     /// @brief Prefetches data to the specified device, i.e. ensure the device has an up-to-date copy of the memory specified
@@ -222,7 +222,7 @@ class UnifiedBuffer
     /// @param stream  cuda stream
     void prefetch(ptrdiff_t byteOffset = 0, size_t size = 0, int dev = cudaCpuDeviceId, cudaStream_t stream = cudaStreamPerThread) const
     {
-        cudaCheck(cudaMemPrefetchAsync(util::PtrAdd(mPtr, byteOffset), size ? size : mSize, dev, stream));
+        cudaCheck(util::cuda::memPrefetchAsync(util::PtrAdd(mPtr, byteOffset), size ? size : mSize, dev, stream));
     }
 
     ///////////////////////////////////////////////////////////////////////
@@ -234,7 +234,7 @@ class UnifiedBuffer
     /// @note Legacy method included for compatibility with DeviceBuffer
     void deviceUpload(int device = 0, cudaStream_t stream = cudaStreamPerThread, bool sync = false) const
     {
-        cudaCheck(cudaMemPrefetchAsync(mPtr, mSize, device, stream));
+        cudaCheck(util::cuda::memPrefetchAsync(mPtr, mSize, device, stream));
         if (sync) cudaCheck(cudaStreamSynchronize(stream));
     }
     void deviceUpload(int device, void* stream, bool sync) const{this->deviceUpload(device, cudaStream_t(stream));}
@@ -256,7 +256,7 @@ class UnifiedBuffer
     /// @param sync if false the memory copy is asynchronous
     void deviceDownload(cudaStream_t stream = 0, bool sync = false) const
     {
-        cudaCheck(cudaMemPrefetchAsync(mPtr, mSize, cudaCpuDeviceId, stream));
+        cudaCheck(util::cuda::memPrefetchAsync(mPtr, mSize, cudaCpuDeviceId, stream));
         if (sync) cudaCheck(cudaStreamSynchronize(stream));
     }
 

nanovdb/nanovdb/tools/cuda/DistributedPointsToGrid.cuh

@@ -447,17 +447,17 @@ void DistributedPointsToGrid<BuildT>::countNodes(const PtrT coords, size_t coord
         uint64_t* deviceOutputKeys = mData->d_keys + deviceStripeOffset;
         uint32_t* deviceOutputIndices = mData->d_indx + deviceStripeOffset;
 
-        cudaMemAdvise(deviceCoords, deviceStripeCount * sizeof(nanovdb::Coord), cudaMemAdviseSetPreferredLocation, deviceId);
-        cudaMemAdvise(deviceCoords, deviceStripeCount * sizeof(nanovdb::Coord), cudaMemAdviseSetReadMostly, deviceId);
+        util::cuda::memAdvise(deviceCoords, deviceStripeCount * sizeof(nanovdb::Coord), cudaMemAdviseSetPreferredLocation, deviceId);
+        util::cuda::memAdvise(deviceCoords, deviceStripeCount * sizeof(nanovdb::Coord), cudaMemAdviseSetReadMostly, deviceId);
 
-        cudaMemAdvise(deviceInputKeys, deviceStripeCount * sizeof(uint64_t), cudaMemAdviseSetPreferredLocation, deviceId);
-        cudaMemAdvise(deviceInputIndices, deviceStripeCount * sizeof(uint32_t), cudaMemAdviseSetPreferredLocation, deviceId);
-        cudaMemAdvise(deviceOutputKeys, deviceStripeCount * sizeof(uint64_t), cudaMemAdviseSetPreferredLocation, deviceId);
-        cudaMemAdvise(deviceOutputIndices, deviceStripeCount * sizeof(uint32_t), cudaMemAdviseSetPreferredLocation, deviceId);
+        util::cuda::memAdvise(deviceInputKeys, deviceStripeCount * sizeof(uint64_t), cudaMemAdviseSetPreferredLocation, deviceId);
+        util::cuda::memAdvise(deviceInputIndices, deviceStripeCount * sizeof(uint32_t), cudaMemAdviseSetPreferredLocation, deviceId);
+        util::cuda::memAdvise(deviceOutputKeys, deviceStripeCount * sizeof(uint64_t), cudaMemAdviseSetPreferredLocation, deviceId);
+        util::cuda::memAdvise(deviceOutputIndices, deviceStripeCount * sizeof(uint32_t), cudaMemAdviseSetPreferredLocation, deviceId);
 
         uint32_t* devicePointsPerTile = mPointsPerTile + deviceStripeOffset;
-        cudaMemAdvise(devicePointsPerTile, deviceStripeCount * sizeof(uint32_t), cudaMemAdviseSetPreferredLocation, deviceId);
-        cudaMemAdvise(deviceNodeCount(deviceId), 3 * sizeof(uint32_t), cudaMemAdviseSetPreferredLocation, deviceId);
+        util::cuda::memAdvise(devicePointsPerTile, deviceStripeCount * sizeof(uint32_t), cudaMemAdviseSetPreferredLocation, deviceId);
+        util::cuda::memAdvise(deviceNodeCount(deviceId), 3 * sizeof(uint32_t), cudaMemAdviseSetPreferredLocation, deviceId);
     }
 
     // Radix sort the subset of keys assigned to each device in parallel
@@ -472,7 +472,7 @@ void DistributedPointsToGrid<BuildT>::countNodes(const PtrT coords, size_t coord
         uint64_t* deviceOutputKeys = mData->d_keys + deviceStripeOffset;
         uint32_t* deviceOutputIndices = mData->d_indx + deviceStripeOffset;
 
-        cudaMemPrefetchAsync(coords, coordCount * sizeof(nanovdb::Coord), deviceId, stream);
+        util::cuda::memPrefetchAsync(coords, coordCount * sizeof(nanovdb::Coord), deviceId, stream);
 
         nanovdb::util::cuda::offsetLambdaKernel<<<numBlocks(deviceStripeCount), mNumThreads, 0, stream>>>(deviceStripeCount, deviceStripeOffset, TileKeyFunctor<BuildT, PtrT>(), mData, coords, mKeys, mIndices);
 
@@ -661,7 +661,7 @@ void DistributedPointsToGrid<BuildT>::countNodes(const PtrT coords, size_t coord
         uint64_t* deviceOutputKeys = mData->d_keys + deviceStripeOffset;
         uint32_t* devicePointsPerTile = mPointsPerTile + deviceStripeOffset;
 
-        // cudaMemPrefetchAsync(deviceInputKeys, deviceStripeCount * sizeof(uint64_t), deviceId, stream);
+        // util::cuda::memPrefetchAsync(deviceInputKeys, deviceStripeCount * sizeof(uint64_t), deviceId, stream);
 
         CUB_LAUNCH(DeviceRunLengthEncode::Encode, mTempDevicePools[deviceId], stream, deviceInputKeys, deviceOutputKeys, devicePointsPerTile, deviceNodeCount(deviceId) + 2, deviceStripeCount);
         cudaCheck(cudaEventRecord(runLengthEncodeEvents[deviceId], stream));

nanovdb/nanovdb/util/cuda/Util.h

@@ -198,6 +198,43 @@ inline size_t blocksPerGrid(size_t numItems, size_t threadsPerBlock)
     return (numItems + threadsPerBlock - 1) / threadsPerBlock;
 }
 
+// CUDA 13.0 changes cudaMemPrefetchAsync and cudaMemPrefetch to use a cudaMemLocation as an argument as
+// opposed to an integer device id. This function provides compatibility by returning the corresponding
+// location in CUDA 13.0 and above while passing through the device in earlier versions.
+#if (CUDART_VERSION < 13000)
+/// @brief Compatbility wrapper for cudaMemAdvise/cudaMemAdvise
+inline cudaError_t memAdvise(const void* devPtr, size_t count, cudaMemoryAdvise advice, int device) {
+    return cudaMemAdvise(devPtr, count, advice, device);
+}
+
+/// @brief Compatbility wrapper for cudaMemPrefetchAsync/cudaMemPrefetchAsync
+inline cudaError_t memPrefetchAsync(const void* devPtr, size_t count, int dstDevice, cudaStream_t stream) {
+    return cudaMemPrefetchAsync(devPtr, count, dstDevice, stream);
+}
+#else
+/// @brief Helper function that converts a device id to a cudaMemLocation
+/// @param device Integer device id
+/// @return cudaMemLocation corresponding to the device id
+inline cudaMemLocation deviceToLocation(int device) {
+    if (device < cudaCpuDeviceId) {
+        return {cudaMemLocationTypeInvalid, device};
+    } else if (device == cudaCpuDeviceId) {
+        return {cudaMemLocationTypeHost, device};
+    } else {
+        return {cudaMemLocationTypeDevice, device};
+    }
+}
+
+/// @brief Compatbility wrapper for cudaMemAdvise/cudaMemAdvise
+inline cudaError_t memAdvise(const void* devPtr, size_t count, cudaMemoryAdvise advice, int device) {
+    return cudaMemAdvise(devPtr, count, advice, deviceToLocation(device));
+}
+
+/// @brief Compatbility wrapper for cudaMemPrefetchAsync/cudaMemPrefetchAsync
+inline cudaError_t memPrefetchAsync(const void* devPtr, size_t count, int dstDevice, cudaStream_t stream) {
+    return cudaMemPrefetchAsync(devPtr, count, deviceToLocation(dstDevice), 0u, stream);
+}
+#endif
 
 #if defined(__CUDACC__)// the following functions only run on the GPU!