1. Examples
Copy example files to your home directory.
1. Login to HPC 2. Copy
cp -r /project/xuehaiqi_652/cuda .
The hello.cu contains the CUDA implementation of HelloWorld.
- Login to HPC
- Setup MPI toolchain:
- Compile
nvcc -O3 hello.cu
- Run
The option -t specifies the limit of run time. Setting it as a small number will get your program scheduled earlier. The option –mem specifies the minimum memory requirement. Setting it as a small number will get your program scheduled earlier. However, you need it when you run sumArraysOnGPU and sumMatrixOnGPU. For more information on srun options, you can use man srun to find out.
- Profile (optional)
srun -n1 –gres=gpu:p100:1 –partition=debug nvprof ./a.out
- Allocate a machine
module purge
module load gcc /8.3.0 cuda /10.1.243
srun -n1 –gres=gpu:1 –mem=16G -t1 ./a.out
salloc -n1 –gres=gpu:1 –mem=16G -t10
// After the allocation, you will log on the machine and have 10 minutes to perform multiple operations
./a.out
// edit , compile , and run again without waiting for a new allocation
./a.out ./a.out
2. (40 points) In reduceInteger.cu, refer to the kernel reduceUnrolling8 and replace the following code segment:
PA5b— Spring 2021
// unrolling 8
int a1 = g_idata[idx]; int a2 = g_idata[idx + int a3 = g_idata[idx + int a4 = g_idata[idx + int b1 = g_idata[idx + int b2 = g_idata[idx + int b3 = g_idata[idx + int b4 = g_idata[idx +
blockDim.x];
2 * blockDim.x]; 3 * blockDim.x]; 4 * blockDim.x]; 5 * blockDim.x]; 6 * blockDim.x]; 7 * blockDim.x];
with the functionally equivalent code below:
int *ptr = g_idata + idx;
int tmp = 0;
// Increment tmp 8 times with values strided by blockDim.x for *int i = 0; i < 8; i++) {
tmp += *ptr;
ptr += blockDim.x; }
g_idata[idx] = tmp;
Compare the performance.
- (30 points) Refer to the kernel reduceInterleaved and the kernel reduceCompleteUnrollWraps8 and implement a version of each for floats. Compare their performance and explain
any differences (profiling with nvprof is optional). Are there any differences compared
to operating on integer data types? - (30 points) Refer to the file nestedHelloWorld.cu and implement a new kernel using the methods illustrated in Figure 3-30 (Change igrid to 2; Only one thread calls the kernel recursively). To compile it, add the following options.
nvcc -O3 -arch=sm_35 -rdc=true hello.cu
