NVIDIA OptiX 7.1 API nvidia_logo_transpbg.gif Up
 All Classes Namespaces Files Functions Variables Typedefs Enumerations Enumerator Macros Groups
optixPaging.h
Go to the documentation of this file.
1 //
2 // Copyright (c) 2020 NVIDIA Corporation. All rights reserved.
3 //
4 // Redistribution and use in source and binary forms, with or without
5 // modification, are permitted provided that the following conditions
6 // are met:
7 // * Redistributions of source code must retain the above copyright
8 // notice, this list of conditions and the following disclaimer.
9 // * Redistributions in binary form must reproduce the above copyright
10 // notice, this list of conditions and the following disclaimer in the
11 // documentation and/or other materials provided with the distribution.
12 // * Neither the name of NVIDIA CORPORATION nor the names of its
13 // contributors may be used to endorse or promote products derived
14 // from this software without specific prior written permission.
15 //
16 // THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS ``AS IS'' AND ANY
17 // EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
18 // IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
19 // PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR
20 // CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
21 // EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
22 // PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR
23 // PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY
24 // OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
25 // (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
26 // OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
27 //
28 
29 #pragma once
30 
31 #include <cuda_runtime.h>
32 
33 #if !defined(__CUDACC_RTC__)
34 #include <utility>
35 #include <stdio.h>
36 #endif
37 
38 inline bool optixPagingCheckCudaError( cudaError_t err )
39 {
40  if( err != cudaSuccess )
41  {
42  printf( "CUDA error: %d\n", err );
43  return false;
44  }
45  return true;
46 }
47 
48 #if !defined( OPTIX_PAGING_CHECK_CUDA_ERROR )
49 #define OPTIX_PAGING_CHECK_CUDA_ERROR( err ) optixPagingCheckCudaError( err )
50 #endif
51 
52 const int MAX_WORKER_THREADS = 32;
53 
54 template <typename T>
55 __host__ __device__ T minimum( T lhs, T rhs )
56 {
57  return lhs < rhs ? lhs : rhs;
58 }
59 
60 template <typename T>
61 __host__ __device__ T maximum( T lhs, T rhs )
62 {
63  return lhs > rhs ? lhs : rhs;
64 }
65 
67 {
68  unsigned int id;
69  unsigned long long page;
70 };
71 
73 {
74  unsigned int pageTableSizeInBytes; // only one for all workers
75  unsigned int usageBitsSizeInBytes; // per worker
76 };
77 
79 {
80  unsigned int maxVaSizeInPages;
81  unsigned int initialVaSizeInPages;
82 };
83 
85 {
86  unsigned int maxVaSizeInPages;
87  unsigned int* usageBits; // also beginning of referenceBits. [ referenceBits | residencesBits ]
88  unsigned int* residenceBits; // located half way into usasgeBits.
89  unsigned long long* pageTable;
90 };
91 
92 #ifndef __CUDACC_RTC__
93 __host__ void optixPagingCreate( OptixPagingOptions* options, OptixPagingContext** context );
94 __host__ void optixPagingDestroy( OptixPagingContext* context );
95 __host__ void optixPagingCalculateSizes( unsigned int vaSizeInPages, OptixPagingSizes& sizes );
96 __host__ void optixPagingSetup( OptixPagingContext* context, const OptixPagingSizes& sizes, int numWorkers );
97 __host__ void optixPagingPullRequests( OptixPagingContext* context,
98  unsigned int* devRequestedPages,
99  unsigned int numRequestedPages,
100  unsigned int* devStalePages,
101  unsigned int numStalePages,
102  unsigned int* devEvictablePages,
103  unsigned int numEvictablePages,
104  unsigned int* devNumPagesReturned );
105 __host__ void optixPagingPushMappings( OptixPagingContext* context,
106  PageMapping* devFilledPages,
107  int filledPageCount,
108  unsigned int* devInvalidatedPages,
109  int invalidatedPageCount );
110 #endif
111 
112 #if defined( __CUDACC__ ) || defined( OPTIX_PAGING_BIT_OPS )
113 __device__ inline void atomicSetBit( unsigned int bitIndex, unsigned int* bitVector )
114 {
115  const unsigned int wordIndex = bitIndex >> 5;
116  const unsigned int bitOffset = bitIndex % 32;
117  const unsigned int mask = 1U << bitOffset;
118  atomicOr( bitVector + wordIndex, mask );
119 }
120 
121 __device__ inline void atomicUnsetBit( int bitIndex, unsigned int* bitVector )
122 {
123  const int wordIndex = bitIndex / 32;
124  const int bitOffset = bitIndex % 32;
125 
126  const int mask = ~( 1U << bitOffset );
127  atomicAnd( bitVector + wordIndex, mask );
128 }
129 
130 __device__ inline bool checkBitSet( unsigned int bitIndex, const unsigned int* bitVector )
131 {
132  const unsigned int wordIndex = bitIndex >> 5;
133  const unsigned int bitOffset = bitIndex % 32;
134  return ( bitVector[wordIndex] & ( 1U << bitOffset ) ) != 0;
135 }
136 
137 __device__ inline unsigned long long optixPagingMapOrRequest( unsigned int* usageBits, unsigned int* residenceBits, unsigned long long* pageTable, unsigned int page, bool* valid )
138 {
139  bool requested = checkBitSet( page, usageBits );
140  if( !requested )
141  atomicSetBit( page, usageBits );
142 
143  bool mapped = checkBitSet( page, residenceBits );
144  *valid = mapped;
145 
146  return mapped ? pageTable[page] : 0;
147 }
148 #endif
149