@@ -103,31 +103,25 @@ the program is compiled to target our local compute capability architecture with
103103
104104.. code-block :: python
105105
106- Initialize CUDA Driver API
107- ==========================
106+ # Initialize CUDA Driver API
108107 checkCudaErrors(driver.cuInit(0 ))
109108
110- Retrieve handle for device 0
111- ============================
109+ # Retrieve handle for device 0
112110 cuDevice = checkCudaErrors(driver.cuDeviceGet(0 ))
113111
114- Derive target architecture for device 0
115- ====================================== =
112+ # Derive target architecture for device 0
116113 major = checkCudaErrors(driver.cuDeviceGetAttribute(driver.CUdevice_attribute.CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR , cuDevice))
117114 minor = checkCudaErrors(driver.cuDeviceGetAttribute(driver.CUdevice_attribute.CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR , cuDevice))
118115 arch_arg = bytes (f ' --gpu-architecture=compute_ { major}{ minor} ' , ' ascii' )
119116
120- Create program
121- ==============
117+ # Create program
122118 prog = checkCudaErrors(nvrtc.nvrtcCreateProgram(str .encode(saxpy), b " saxpy.cu" , 0 , [], []))
123119
124- Compile program
125- ============== =
120+ # Compile program
126121 opts = [b " --fmad=false" , arch_arg]
127122 checkCudaErrors(nvrtc.nvrtcCompileProgram(prog, 2 , opts))
128123
129- Get PTX from compilation
130- ========================
124+ # Get PTX from compilation
131125 ptxSize = checkCudaErrors(nvrtc.nvrtcGetPTXSize(prog))
132126 ptx = b " " * ptxSize
133127 checkCudaErrors(nvrtc.nvrtcGetPTX(prog, ptx))
@@ -139,8 +133,7 @@ following code example, a handle for compute device 0 is passed to
139133
140134.. code-block :: python
141135
142- Create context
143- ==============
136+ # Create context
144137 context = checkCudaErrors(driver.cuCtxCreate(0 , cuDevice))
145138
146139 With a CUDA context created on device 0, load the PTX generated earlier into a
@@ -150,11 +143,9 @@ After loading into the module, extract a specific kernel with
150143
151144.. code-block :: python
152145
153- Load PTX as module data and retrieve function
154- ============================================ =
146+ # Load PTX as module data and retrieve function
155147 ptx = np.char.array(ptx)
156- Note: Incompatible -- gpu- architecture would be detected here
157- ============================================================
148+ # Note: Incompatible --gpu-architecture would be detected here
158149 module = checkCudaErrors(driver.cuModuleLoadData(ptx.ctypes.data))
159150 kernel = checkCudaErrors(driver.cuModuleGetFunction(module, b " saxpy" ))
160151
@@ -258,8 +249,7 @@ in the designated stream are finished.
258249
259250.. code-block :: python
260251
261- Assert values are same after running kernel
262- ========================================== =
252+ # Assert values are same after running kernel
263253 hZ = a * hX + hY
264254 if not np.allclose(hOut, hZ):
265255 raise ValueError (" Error outside tolerance for host-device vectors" )
@@ -416,18 +406,15 @@ Note how all three pointers are ``np.intp`` since the pointer values are always
416406Putting it all together:
417407.. code-block :: python
418408
419- Define a custom type
420- ====================
409+ # Define a custom type
421410 testStruct = np.dtype([(" value" , np.int32)], align = True )
422411
423- Allocate device memory
424- ======================
412+ # Allocate device memory
425413 pInt = checkCudaErrors(cudart.cudaMalloc(np.dtype(np.int32).itemsize))
426414 pFloat = checkCudaErrors(cudart.cudaMalloc(np.dtype(np.float32).itemsize))
427415 pStruct = checkCudaErrors(cudart.cudaMalloc(testStruct.itemsize))
428416
429- Collect all input kernel arguments into a single tuple for further processing
430- ============================================================================ =
417+ # Collect all input kernel arguments into a single tuple for further processing
431418 kernelValues = (
432419 np.array(1 , dtype = np.uint32),
433420 np.array([pInt], dtype = np.intp),
@@ -478,19 +465,16 @@ For this example the result becomes:
478465
479466.. code-block :: python
480467
481- Define a custom type
482- ====================
468+ # Define a custom type
483469 class testStruct (ctypes .Structure ):
484470 _fields_ = [(" value" , ctypes.c_int)]
485471
486- Allocate device memory
487- ======================
472+ # Allocate device memory
488473 pInt = checkCudaErrors(cudart.cudaMalloc(ctypes.sizeof(ctypes.c_int)))
489474 pFloat = checkCudaErrors(cudart.cudaMalloc(ctypes.sizeof(ctypes.c_float)))
490475 pStruct = checkCudaErrors(cudart.cudaMalloc(ctypes.sizeof(testStruct)))
491476
492- Collect all input kernel arguments into a single tuple for further processing
493- ============================================================================ =
477+ # Collect all input kernel arguments into a single tuple for further processing
494478 kernelValues = (
495479 1 ,
496480 pInt,
0 commit comments