Merge branch 'main' into 324-linker-example

Author: leofang

cuda_bindings/examples/0_Introduction/clock_nvrtc_test.py

@@ -58,6 +58,10 @@
 NUM_THREADS = 256
 
 
+def elems_to_bytes(nelems, dt):
+    return nelems * np.dtype(dt).itemsize
+
+
 def main():
     print("CUDA Clock sample")
 
@@ -75,31 +79,31 @@ def main():
     kernelHelper = common.KernelHelper(clock_nvrtc, devID)
     kernel_addr = kernelHelper.getFunction(b"timedReduction")
 
-    dinput = checkCudaErrors(cuda.cuMemAlloc(np.dtype(np.float32).itemsize * NUM_THREADS * 2))
-    doutput = checkCudaErrors(cuda.cuMemAlloc(np.dtype(np.float32).itemsize * NUM_BLOCKS))
-    dtimer = checkCudaErrors(cuda.cuMemAlloc(np.dtype(np.int64).itemsize * NUM_BLOCKS * 2))
-    checkCudaErrors(cuda.cuMemcpyHtoD(dinput, hinput, np.dtype(np.float32).itemsize * NUM_THREADS * 2))
+    dinput = checkCudaErrors(cuda.cuMemAlloc(hinput.nbytes))
+    doutput = checkCudaErrors(cuda.cuMemAlloc(elems_to_bytes(NUM_BLOCKS, np.float32)))
+    dtimer = checkCudaErrors(cuda.cuMemAlloc(timer.nbytes))
+    checkCudaErrors(cuda.cuMemcpyHtoD(dinput, hinput, hinput.nbytes))
+
+    args = ((dinput, doutput, dtimer), (None, None, None))
+    shared_memory_nbytes = elems_to_bytes(2 * NUM_THREADS, np.float32)
 
-    arr = ((dinput, doutput, dtimer), (None, None, None))
+    grid_dims = (NUM_BLOCKS, 1, 1)
+    block_dims = (NUM_THREADS, 1, 1)
 
     checkCudaErrors(
         cuda.cuLaunchKernel(
             kernel_addr,
-            NUM_BLOCKS,
-            1,
-            1,  # grid dim
-            NUM_THREADS,
-            1,
-            1,  # block dim
-            np.dtype(np.float32).itemsize * 2 * NUM_THREADS,
+            *grid_dims,  # grid dim
+            *block_dims,  # block dim
+            shared_memory_nbytes,
             0,  # shared mem, stream
-            arr,
+            args,
             0,
         )
     )  # arguments
 
     checkCudaErrors(cuda.cuCtxSynchronize())
-    checkCudaErrors(cuda.cuMemcpyDtoH(timer, dtimer, np.dtype(np.int64).itemsize * NUM_BLOCKS * 2))
+    checkCudaErrors(cuda.cuMemcpyDtoH(timer, dtimer, timer.nbytes))
     checkCudaErrors(cuda.cuMemFree(dinput))
     checkCudaErrors(cuda.cuMemFree(doutput))
     checkCudaErrors(cuda.cuMemFree(dtimer))

cuda_bindings/examples/0_Introduction/simpleCubemapTexture_test.py

@@ -6,7 +6,6 @@
 # this software and related documentation outside the terms of the EULA
 # is strictly prohibited.
 import ctypes
-import math
 import sys
 import time
 
@@ -103,18 +102,11 @@ def main():
     num_faces = 6
     num_layers = 1
     cubemap_size = width * width * num_faces
-    size = cubemap_size * num_layers * np.dtype(np.float32).itemsize
-    h_data = np.zeros(cubemap_size * num_layers, dtype="float32")
-
-    for i in range(cubemap_size * num_layers):
-        h_data[i] = i
+    h_data = np.arange(cubemap_size * num_layers, dtype="float32")
+    size = h_data.nbytes
 
     # This is the expected transformation of the input data (the expected output)
-    h_data_ref = np.zeros(cubemap_size * num_layers, dtype="float32")
-
-    for layer in range(num_layers):
-        for i in range(cubemap_size):
-            h_data_ref[layer * cubemap_size + i] = -h_data[layer * cubemap_size + i] + layer
+    h_data_ref = np.repeat(np.arange(num_layers, dtype=h_data.dtype), cubemap_size) - h_data
 
     # Allocate device memory for result
     d_data = checkCudaErrors(cudart.cudaMalloc(size))
@@ -130,10 +122,11 @@ def main():
             cudart.cudaArrayCubemap,
         )
     )
+    width_nbytes = h_data[:width].nbytes
     myparms = cudart.cudaMemcpy3DParms()
     myparms.srcPos = cudart.make_cudaPos(0, 0, 0)
     myparms.dstPos = cudart.make_cudaPos(0, 0, 0)
-    myparms.srcPtr = cudart.make_cudaPitchedPtr(h_data, width * np.dtype(np.float32).itemsize, width, width)
+    myparms.srcPtr = cudart.make_cudaPitchedPtr(h_data, width_nbytes, width, width)
     myparms.dstArray = cu_3darray
     myparms.extent = cudart.make_cudaExtent(width, width, num_faces)
     myparms.kind = cudart.cudaMemcpyKind.cudaMemcpyHostToDevice
@@ -211,23 +204,21 @@ def main():
     print(f"{cubemap_size / ((stop - start + 1) / 1000.0) / 1e6:.2f} Mtexlookups/sec")
 
     # Allocate mem for the result on host side
-    h_odata = np.zeros(cubemap_size * num_layers, dtype="float32")
+    h_odata = np.empty_like(h_data)
     # Copy result from device to host
     checkCudaErrors(cudart.cudaMemcpy(h_odata, d_data, size, cudart.cudaMemcpyKind.cudaMemcpyDeviceToHost))
 
-    print("Comparing kernel output to expected data")
-    MIN_EPSILON_ERROR = 5.0e-3
-    for i in range(cubemap_size * num_layers):
-        d = h_odata[i] - h_data_ref[i]
-        if math.fabs(d) > MIN_EPSILON_ERROR:
-            print("Failed")
-            sys.exit(-1)
-    print("Passed")
-
     checkCudaErrors(cudart.cudaDestroyTextureObject(tex))
     checkCudaErrors(cudart.cudaFree(d_data))
     checkCudaErrors(cudart.cudaFreeArray(cu_3darray))
 
+    print("Comparing kernel output to expected data")
+    MIN_EPSILON_ERROR = 5.0e-3
+    if np.max(np.abs(h_odata - h_data_ref)) > MIN_EPSILON_ERROR:
+        print("Failed")
+        sys.exit(-1)
+    print("Passed")
+
 
 if __name__ == "__main__":
     main()

cuda_bindings/examples/0_Introduction/vectorAddDrv_test.py

@@ -38,7 +38,7 @@
 def main():
     print("Vector Addition (Driver API)")
     N = 50000
-    size = N * np.dtype(np.float32).itemsize
+    nbytes = N * np.dtype(np.float32).itemsize
 
     # Initialize
     checkCudaErrors(cuda.cuInit(0))
@@ -57,18 +57,18 @@ def main():
     _VecAdd_kernel = kernelHelper.getFunction(b"VecAdd_kernel")
 
     # Allocate input vectors h_A and h_B in host memory
-    h_A = np.random.rand(size).astype(dtype=np.float32)
-    h_B = np.random.rand(size).astype(dtype=np.float32)
-    h_C = np.random.rand(size).astype(dtype=np.float32)
+    h_A = np.random.rand(N).astype(dtype=np.float32)
+    h_B = np.random.rand(N).astype(dtype=np.float32)
+    h_C = np.random.rand(N).astype(dtype=np.float32)
 
     # Allocate vectors in device memory
-    d_A = checkCudaErrors(cuda.cuMemAlloc(size))
-    d_B = checkCudaErrors(cuda.cuMemAlloc(size))
-    d_C = checkCudaErrors(cuda.cuMemAlloc(size))
+    d_A = checkCudaErrors(cuda.cuMemAlloc(nbytes))
+    d_B = checkCudaErrors(cuda.cuMemAlloc(nbytes))
+    d_C = checkCudaErrors(cuda.cuMemAlloc(nbytes))
 
     # Copy vectors from host memory to device memory
-    checkCudaErrors(cuda.cuMemcpyHtoD(d_A, h_A, size))
-    checkCudaErrors(cuda.cuMemcpyHtoD(d_B, h_B, size))
+    checkCudaErrors(cuda.cuMemcpyHtoD(d_A, h_A, nbytes))
+    checkCudaErrors(cuda.cuMemcpyHtoD(d_B, h_B, nbytes))
 
     if True:
         # Grid/Block configuration
@@ -98,7 +98,7 @@ def main():
 
     # Copy result from device memory to host memory
     # h_C contains the result in host memory
-    checkCudaErrors(cuda.cuMemcpyDtoH(h_C, d_C, size))
+    checkCudaErrors(cuda.cuMemcpyDtoH(h_C, d_C, nbytes))
 
     for i in range(N):
         sum_all = h_A[i] + h_B[i]