ilusha/pbrt-v4
1
0
Fork 0
mirror of https://github.com/mmp/pbrt-v4 synced 2026-09-26 16:20:07 +03:00
Code Issues Projects Releases Packages Wiki Activity Actions 18

Merge pull request #515 from SPolton/fix-cuda-13

Browse source 
Fix CUDA 13.x Compatibility in PBRT-v4
This commit is contained in:
Matt Pharr 2025-10-22 12:22:44 -07:00 • committed by GitHub
commit 689587986d
No known key found for this signature in database
GPG key ID: B5690EEEBB952194
4 changed files with 30 additions and 2 deletions
Download patch file Download diff file Expand all files Collapse all files

2
.gitignore vendored
View file

@ -4,5 +4,5 @@
src/build
.DS_Store
.ipynb_checkpoints/
build/
*build*/
.cache/

8
src/pbrt/gpu/memory.cpp
View file

@ -61,8 +61,16 @@ void CUDATrackedMemoryResource::PrefetchToGPU() const {
LOG_VERBOSE("Prefetching %d allocations to GPU memory", allocations.size());
size_t bytes = 0;
for (auto iter : allocations) {
#if CUDART_VERSION >= 13000
cudaMemLocation location = {};
location.type = cudaMemLocationTypeDevice;
location.id = deviceIndex;
CUDA_CHECK(
cudaMemPrefetchAsync(iter.first, iter.second, location, 0 /* stream */));
#else
CUDA_CHECK(
cudaMemPrefetchAsync(iter.first, iter.second, deviceIndex, 0 /* stream */));
#endif
bytes += iter.second;
}
CUDA_CHECK(cudaDeviceSynchronize());

10
src/pbrt/gpu/util.cpp
View file

@ -48,11 +48,19 @@ void GPUInit() {
CUDA_CHECK(cudaGetDeviceProperties(&deviceProperties, i));
CHECK(deviceProperties.canMapHostMemory);
#if CUDART_VERSION >= 13000
int clockRateKHz = 0;
cudaDeviceGetAttribute(&clockRateKHz, cudaDevAttrClockRate, i);
float clockRate = clockRateKHz;
#else
float clockRate = deviceProperties.clockRate;
#endif
std::string deviceString = StringPrintf(
"CUDA device %d (%s) with %f MiB, %d SMs running at %f MHz "
"with shader model %d.%d",
i, deviceProperties.name, deviceProperties.totalGlobalMem / (1024. * 1024.),
deviceProperties.multiProcessorCount, deviceProperties.clockRate / 1000.,
deviceProperties.multiProcessorCount, clockRate / 1000.,
deviceProperties.major, deviceProperties.minor);
LOG_VERBOSE("%s", deviceString);
devices += deviceString + "\n";

12
src/pbrt/wavefront/integrator.cpp
View file

@ -618,10 +618,22 @@ void WavefrontPathIntegrator::PrefetchGPUAllocations() {
// performance. (This makes it possible to use the values of things
// like WavefrontPathIntegrator::haveSubsurface to conditionally launch
// kernels according to what's in the scene...)
#if CUDART_VERSION >= 13000
cudaMemLocation location = {};
location.type = cudaMemLocationTypeDevice;
location.id = 0; // For ReadMostly: device ID is ignored
CUDA_CHECK(cudaMemAdvise(this, sizeof(*this), cudaMemAdviseSetReadMostly,
location));
location.id = deviceIndex;
CUDA_CHECK(cudaMemAdvise(this, sizeof(*this), cudaMemAdviseSetPreferredLocation,
location));
#else
CUDA_CHECK(cudaMemAdvise(this, sizeof(*this), cudaMemAdviseSetReadMostly,
/* ignored argument */ 0));
CUDA_CHECK(cudaMemAdvise(this, sizeof(*this), cudaMemAdviseSetPreferredLocation,
deviceIndex));
#endif
// Copy all of the scene data structures over to GPU memory. This
// ensures that there isn't a big performance hitch for the first batch