nvcc<\/code> currently runs out of a singularity image that must be run on the node before your jobs can execute. If the singularity container does not already exist, it can be recreated by following the steps below. The singularity definition file can be found in the appendix<\/a><\/p>\n\n\n\n # This command will build a singularity image from the definition file\n # {singularity_def_file}.def\n sudo singularity build {image_name}.sif {singularity_def_file}.def\n\n # Once the image is built, it can be deployed in a standalone environment on any \n # system with singularity as below:\n singularity run {image_name}.sif\n<\/code><\/pre>\n\n\n\nOnce your singularity image is ready, we can now run a job. CHTC manages jobs using the HTCondor framework. To submit a job to the requested node, we run the following batch script:<\/p>\n\n\n\n
universe = vanilla\noutput = nsight-compute\/log\/ncu-$(CLUSTER).$(PROCESS).out\nerror = nsight-compute\/log\/ncu-$(CLUSTER).$(PROCESS).err\nlog = nsight-compute\/log\/ncu-$(CLUSTER).$(PROCESS).log\n\ntransfer_executable = True\nexecutable = {profiling-script}.sh\nshould_transfer_files = yes\ntransfer_input_files = {image_name}.sif\ntransfer_output_files = {gpu-profiling-report}\nwhen_to_transfer_output = ON_EXIT\n\n+WantGPULab = true\n+InteractiveJob = true\n+GPUJobLength = \"short\"\n\nrequest_gpus = 1\nrequest_cpus = 1\nrequest_memory = 2GB\nrequest_disk = 7GB\n\n+SingularityImage = \"{image_name}.sif\"\nRequirements = HasSingularity\nRequirements = GPUs_DeviceName == \"NVIDIA A100-SXM4-40GB\"\nrequirements = Machine == \"gpulab2004.chtc.wisc.edu\" && OpSysMajorVer == 8\n\n\nqueue 1\n<\/code><\/pre>\n\n\n\nInformation on managing a condor job can be found on the HTCondor Documentation Page<\/a>.<\/p>\n\n\n\nRunning on CHTC<\/h2>\n\n\n\n
At present, it is not possible to run profiling jobs on cmsgpu01<\/code> without sudo access.<\/p>\n\n\n\n
\n\n\n\n(A brief) Introduction to GPU Programming<\/h1>\n\n\n\n
This section assumes familiarity with the CUDA programming model and fundamentals of GPU hardware. It is also assumed that the reader is comfortable with programming in C\/C++. For a more comprehensive introduction to CUDA, check out CUDA Training Series by Oakland National Laboratory<\/a>.<\/p>\n\n\n\nThe CUDA Programming Model<\/h2>\n\n\n\n
In order to perform computation, the host launches a kernel that is executed simultaneously by many threads on the device. It is important to note that a GPU thread differs from a regular CPU thread. The launch configuration defines the number of thread blocks and the number of threads in each block that execute the kernel. Each thread is associated with an intrinsic index, which can be used to calculate and access memory locations. Each thread has its own context and set of private variables. All threads have access to the GPU global memory, but there is no general way to synchronize execution of arbitrary threads.<\/p>\n\n\n\n
At the hardware level, the CUDA programming model groups 32 threads in groups called warps. All members of a warp execute the same instruction in lockstep. If threads within a same warp take two different execution paths on encountering a branching condition, both branches are executed sequentially with non-participating threads in that branch being inactive. This condition is called branch divergence, and is a major source of latency.<\/p>\n\n\n\n
Memory access to the GPU memory occur in blocks of specific sizes (32B, 64B, 128B, etc.). To efficiently access memory with minimum latency, all threads in a warp need to access elements of data that lie in the same cache line<\/strong>. This is called coalesced memory access<\/em>.<\/p>\n\n\n\nThreads are further grouped into thread blocks, which reside on a piece of hardware called the Streaming Multiprocessor (SM). Each block is assigned to a single SM. A SM consists of many computing elements, schedulers, and an extremely fast on-chip memory called shared memory. It is possible to synchronize threads at a block level. For optimal performance, blocks should have more than one warp. If sufficient resources are available, it is possible for an SM to have multiple active thread blocks simultaneously; however, these blocks cannot share data with each other.<\/p>\n\n\n\n
In order to hide latencies it is recommended to \u201cover-subscribe\u201d the GPU. There should be many more blocks than SMs present on the device. Also in order to ensure a good occupancy of the CUDA cores there should be more warps active on a given SM than SIMT units. This way while some warps of threads are idle waiting for some memory operations to complete, others use the CUDA cores, thus ensuring a high occupancy of the GPU.<\/p>\n\n\n\n
\n\n\n\nNVIDIA NSight Compute<\/h1>\n\n\n\n
Nsight Compute<\/a> is a powerful profiling tool that provides detailed performance data and offers guidance for optimizing CUDA kernels. Nsight Compute offers both a GUI as well as a command line interaface (CLI), and provides full support for CUDA 10.0+ and Volta+ GPUs.<\/p>\n\n\n\nGetting Started<\/h2>\n\n\n\n
As a pedagogical exercise in learning how to use Nsight Compute, we\u2019re going to profile a CUDA kernel that does a matrix-matrix element-wise add operation using a 2D CUDA grid configuration.<\/p>\n\n\n\n
#include <iostream>\n\nconst size_t size_w = 1024;\nconst size_t size_h = 1024;\n\ntypedef unsigned mytype;\ntypedef mytype arr_t[size_w];\nconst mytype A_val = 1;\nconst mytype B_val = 2;\n\n__global__ void matrix_add_2D(const arr_t* __restrict__ A, const arr_t* __restrict__ B, arr_t* __restrict__ C, const size_t sw, const size_t sh)\n{\n size_t idx = threadIdx.x + blockDim.x * (size_t) blockIdx.x;\n size_t idy = threadIdx.y + blockDim.y * (size_t) blockIdx.y;\n if((idx < sh) && (idy < sw))\n C[idx][idy] = A[idx][idy] + B[idx][idy];\n}\n\nint main()\n{\n arr_t *A, *B, *C;\n size_t ds = size_w * size_h * sizeof(mytype);\n cudaError_t err = cudaMallocManaged(&A, ds);\n if (err != cudaSuccess)\n {\n std::cout << \"CUDA error: \" << cudaGetErrorString(err) << std::endl;\n return 0;\n }\n cudaMallocManaged(&B, ds);\n cudaMallocManaged(&C, ds);\n for (int x = 0; x < size_h; x++)\n {\n for (int y = 0; y < size_h; y++)\n {\n A[x][y] = A_val;\n B[x][y] = B_val;\n C[x][y] = 0;\n }\n }\n\n int attr = 0;\n cudaDeviceGetAttribute(&attr, cudaDevAttrConcurrentManagedAccess, 0);\n if (attr)\n {\n cudaMemPrefetchAsync(A, ds, 0);\n cudaMemPrefetchAsync(B, ds, 0);\n cudaMemPrefetchAsync(C, ds, 0);\n }\n dim3 threads(32, 32);\n dim3 blocks((size_w + threads.x - 1)\/threads.x, (size_h + threads.y - 1)\/threads.y);\n matrix_add_2D<<<blocks, threads>>>(A, B, C, size_w, size_h);\n cudaDeviceSynchronize();\n err = cudaGetLastError();\n if (err != cudaSuccess)\n {\n std::cout << \"CUDA error: \" << cudaGetErrorString(err) << std::endl;\n return 0;\n }\n\n for (int x = 0; x < size_h; x++)\n {\n for(int y = 0; y < size_w; y++)\n {\n if (C[x][y] != A_val + B_val)\n {\n std::cout << \"mismatch at: (\" << x << \", \" << y << \") was: \" << C[x][y] << \" should be: \" << A_val+B_val << std::endl;\n return 0; \n }\n }\n }\n std::cout << \"Success!\" << std::endl;\n return 0;\n}\n<\/code><\/pre>\n\n\n\nThe code should be fairly straightforward to understand, but there are two main features:<\/p>\n\n\n\n
\n- Managed memory:<\/strong> Data is allocated using managed allocations. For GPUs that support on-demand paging, data is prefetched to avoid performance overheads on the kernel.<\/li>\n\n\n\n
- 2D:<\/strong> We are launching a 2D grid of blocks, along with a 2D threadblock shape. This allows us to use simple 2D arrays, without requiring arrays of pointers.<\/li>\n<\/ol>\n\n\n\n
As a CUDA programmer, the two most important optimization priorities are one, to expose enough parallel work to the GPU to avoid any idle cycles, and two, making efficient use of the memory subystem. In this code, we do not have enough data to meet the first objective. We only use global memory, so we\u2019re interested in efficient use of global memory via coalesced memory accesses.<\/p>\n\n\n\n
Using the Nsight Compute CLI<\/h2>\n\n\n\n
For most purposes, we can get similar data using the CLI or the GUI, but the CLI is easier to use if you know specifically what data you are looking for, and\/or if you want to use command line automation.<\/p>\n\n\n\n
If the path to the profiler tool is not setup by default, you want to add \/usr\/local\/cuda\/bin\/<\/code> to your PATH<\/code> variable. You can verify this by running ncu --version<\/code>. Before running the profiler, you must first compile your CUDA code with nvcc compiler. To do so, run bash nvcc -o example.o example.cu --lineinfo<\/code><\/p>\n\n\n\nPlease note that Nsight Compute is not supported on devices with CC 6.0 and lower.<\/p>\n\n\n\n
The --lineinfo<\/code> flag allows Nsight to cross-reference written code with CUDA opcode and suggest targeted optimization strategies by isolating bottlenecks and inefficiencies in user code. For more compilation options, refer to the nvcc documentation<\/a>.<\/p>\n\n\n\nNow, we can run the profiler on the compiled object file.<\/p>\n\n\n\n
$ ncu .\/example\n==PROF== Connected to process 3896054 (\/afs\/hep.wisc.edu\/user\/<user>\/public\/gpu-profiling\/example.o)\n==PROF== Profiling \"matrix_add_2D\" - 0: 0%....50%....100% - 10 passes\nSuccess!\n==PROF== Disconnected from process 3896054\n[3896054] example.o@127.0.0.1\n matrix_add_2D(const unsigned int (*)[1024], const unsigned int (*)[1024], unsigned int (*)[1024], unsigned long, unsigned long) (32, 32, 1)x(32, 32, 1), Context 1, Stream 7, Device 0, CC 7.5\n Section: GPU Speed Of Light Throughput\n ----------------------- ----------- ------------\n Metric Name Metric Unit Metric Value\n ----------------------- ----------- ------------\n DRAM Frequency Ghz 4.92\n SM Frequency Mhz 584.90\n Elapsed Cycles cycle 186,106\n Memory Throughput % 43.35\n DRAM Throughput % 32.94\n Duration us 318.18\n L1\/TEX Cache Throughput % 86.70\n L2 Cache Throughput % 27.16\n SM Active Cycles cycle 177,709.58\n Compute (SM) Throughput % 4.42\n ----------------------- ----------- ------------\n\n OPT This kernel exhibits low compute throughput and memory bandwidth utilization relative to the peak performance \n of this device. Achieved compute throughput and\/or memory bandwidth below 60.0% of peak typically indicate \n latency issues. Look at Scheduler Statistics and Warp State Statistics for potential reasons. \n\n Section: Launch Statistics\n -------------------------------- --------------- ---------------\n Metric Name Metric Unit Metric Value\n -------------------------------- --------------- ---------------\n Block Size 1,024\n Function Cache Configuration CachePreferNone\n Grid Size 1,024\n Registers Per Thread register\/thread 16\n Shared Memory Configuration Size Kbyte 32.77\n Driver Shared Memory Per Block byte\/block 0\n Dynamic Shared Memory Per Block byte\/block 0\n Static Shared Memory Per Block byte\/block 0\n # SMs SM 40\n Threads thread 1,048,576\n Uses Green Context 0\n Waves Per SM 25.60\n -------------------------------- --------------- ---------------\n\n Section: Occupancy\n ------------------------------- ----------- ------------\n Metric Name Metric Unit Metric Value\n ------------------------------- ----------- ------------\n Block Limit SM block 16\n Block Limit Registers block 4\n Block Limit Shared Mem block 16\n Block Limit Warps block 1\n Theoretical Active Warps per SM warp 32\n Theoretical Occupancy % 100\n Achieved Occupancy % 73.27\n Achieved Active Warps Per SM warp 23.45\n ------------------------------- ----------- ------------\n\n OPT Est. Local Speedup: 26.73% \n The difference between calculated theoretical (100.0%) and measured achieved occupancy (73.3%) can be the \n result of warp scheduling overheads or workload imbalances during the kernel execution. Load imbalances can \n occur between warps within a block as well as across blocks of the same kernel. See the CUDA Best Practices \n Guide (https:\/\/docs.nvidia.com\/cuda\/cuda-c-best-practices-guide\/index.html#occupancy) for more details on \n optimizing occupancy. \n\n Section: GPU and Memory Workload Distribution\n -------------------------- ----------- ------------\n Metric Name Metric Unit Metric Value\n -------------------------- ----------- ------------\n Average DRAM Active Cycles cycle 515,437.50\n Total DRAM Elapsed Cycles cycle 12,516,352\n Average L1 Active Cycles cycle 177,709.58\n Total L1 Elapsed Cycles cycle 7,410,016\n Average L2 Active Cycles cycle 266,296.69\n Total L2 Elapsed Cycles cycle 8,703,264\n Average SM Active Cycles cycle 177,709.58\n Total SM Elapsed Cycles cycle 7,410,016\n Average SMSP Active Cycles cycle 170,211.94\n Total SMSP Elapsed Cycles cycle 29,640,064\n -------------------------- ----------- ------------\n<\/code><\/pre>\n\n\n\nThat\u2019s a lot of output! If your code has multiple kernel invocations, details page data will be gathered and displayed for each. We won\u2019t try and go through it all in detail, but notice there are major sections for SOL (speed of light \u2014 comparisons against the best possible behaviour supported on the device), compute analysis, memory analysis, scheduler, warp state, instruction, launch statistics, and occupancy analysis. Also notice section-specific optimization strategies are that suggested by Nsight. while these suggestions might not always be helpful, there are an excellent starting point.<\/p>\n\n\n\n
You can optionally select which of these sections are collected and displayed with command-line flags and parameters. The list of all possible command-line flags can be found in the NSight CLI documentation<\/a>.<\/p>\n\n\n\nIn our example, we\u2019re interested in studying global load\/store transaction. For global load transactions, we will use l1tex\u200b_\u200b_t\u200b_sectors\u200b_pipe\u200b_lsu\u200b_mem\u200b_global\u200b_op\u200b_ld.sum<\/code> and for global load requests we will use l1tex\u200b_\u200b_t\u200b_requests\u200b_pipe\u200b_lsu\u200b_mem\u200b_global\u200b_op\u200b_ld.sum<\/code>. At this point you might be wondering about the length of these metric names and naming convention. There is a method to the naming, and you can review it in the documentation. The naming convention is intended to make it easier to understand what a metric represents from its name. Briefly, the metric name preceding the period identifies where in the architecture the data is being collected, and the token after the period identifies mathematically how the number is gathered. For most base metric names on Volta and newer, suffixes (like .sum, .avg, \u2026) exist that together with the base name make up the actual metric name that can be collected. Once you understand this concept for one metric, you can easily apply it to almost any other available metric on this architecture.<\/p>\n\n\n\nWe gather these metrics by passing their names as a command line argument<\/p>\n\n\n\n
$ ncu --metrics l1tex__t_sectors_pipe_lsu_mem_global_op_ld.sum,l1tex__t_requests_pipe_lsu_mem_global_op_ld.sum .\/exa\nmple.o\n==PROF== Connected to process 3906982 (\/afs\/hep.wisc.edu\/user\/athete\/public\/gpu-profiling\/example.o)\n==PROF== Profiling \"matrix_add_2D\" - 0: 0%....50%....100% - 4 passes\nSuccess!\n==PROF== Disconnected from process 3906982\n[3906982] example.o@127.0.0.1\n matrix_add_2D(const unsigned int (*)[1024], const unsigned int (*)[1024], unsigned int (*)[1024], unsigned long, unsigned long) (32, 32, 1)x(32, 32, 1), Context 1, Stream 7, Device 0, CC 7.5\n Section: Command line profiler metrics\n ----------------------------------------------- ----------- ------------\n Metric Name Metric Unit Metric Value\n ----------------------------------------------- ----------- ------------\n l1tex__t_requests_pipe_lsu_mem_global_op_ld.sum 65,536\n l1tex__t_sectors_pipe_lsu_mem_global_op_ld.sum sector 2,097,152\n ----------------------------------------------- ----------- ------------\n<\/code><\/pre>\n\n\n\nThis first metric above represents the denominator (requests) of the desired measurement (transactions per request) and the second metric represents the numerator (transactions). If we divide these, we get 32 transactions per request. Therefore, each thread in the warp is generating a separate transaction. This is a good indication that our access pattern (reading, in this case) is not coalesced.<\/p>\n\n\n\n
Using the Nsight Compute GUI<\/h2>\n\n\n\n
If instead, we are interested in gathering and analyzing these metrics using the GUI, we need an active X11 session to run the GUI app version in.<\/p>\n\n\n\n
N.B. Ideally, once an active X11 session is established, you type ncu-ui<\/code> to open the GUI and everything should work, but unfortunately I have not been able to get it running by following these steps. Instead, you have to export the profiling report as a ncu-rep<\/code> file, download it to your local machine, and open it there.<\/p>\n\n\n\nFixing the Code<\/h2>\n\n\n\n
The reason for high number of transactions per request in this code is due to how elements are being indexed.<\/p>\n\n\n\n
... C[idx][idy] = A[idx][idy] + B[idx][idy]\n<\/code><\/pre>\n\n\n\nThe index built with threadIdx.x<\/code> (i.e. idx<\/code>) should appear in the last subscript for coalesced access across a warp; instead it appears in the first subscript. While either method can give correct results, they are not the same from a performance perspective. This arrangement results in each thread in a warp accessing data in a \u201ccolumn\u201d in memory, rather than a \u201crow\u201d (i.e. adjacent). We can fix this by modifying our kernel code as follows:<\/p>\n\n\n\n__global__ void matrix_add_2D(const arr_t* __restrict__ A, const arr_t* __restrict__ B, arr_t* __restrict__ C, const size_t sw, const size_t sh)\n{\n size_t idx = threadIdx.x + blockDim.x * (size_t) blockIdx.x;\n size_t idy = threadIdx.y + blockDim.y * (size_t) blockIdx.y;\n if((idx < sh) && (idy < sw))\n C[idy][idx] = A[idy][idx] + B[idy][idx];\n}\n<\/code><\/pre>\n\n\n\nRecompiling and profiling again, we see:<\/p>\n\n\n\n
$ ncu --metrics l1tex__t_sectors_pipe_lsu_mem_global_op_ld.sum,l1tex__t_requests_pipe_lsu_mem_global_op_ld.sum .\/example.o \n==PROF== Connected to process 4192288 (\/afs\/hep.wisc.edu\/user\/athete\/public\/gpu-profiling\/example.o)\n==PROF== Profiling \"matrix_add_2D\" - 0: 0%....50%....100% - 4 passes\nSuccess!\n==PROF== Disconnected from process 4192288\n[4192288] example.o@127.0.0.1\n matrix_add_2D(const unsigned int (*)[1024], const unsigned int (*)[1024], unsigned int (*)[1024], unsigned long, unsigned long) (32, 32, 1)x(32, 32, 1), Context 1, Stream 7, Device 0, CC 7.5\n Section: Command line profiler metrics\n ----------------------------------------------- ----------- ------------\n Metric Name Metric Unit Metric Value\n ----------------------------------------------- ----------- ------------\n l1tex__t_requests_pipe_lsu_mem_global_op_ld.sum 65,536\n l1tex__t_sectors_pipe_lsu_mem_global_op_ld.sum sector 262,144\n ----------------------------------------------- ----------- ------------\n<\/code><\/pre>\n\n\n\nNow the ratio of the metrics is 4:1 (transactions per request), indicating the desired transaction size of 32 bytes is achieved, and the efficiency of loads (and stores) is substantially improved over the previous case.<\/p>\n\n\n\n
\n\n\n\nProfiling MadGraph_aMC@NLO<\/h1>\n\n\n\n
MadGraph_amC@NLO (MG5aMC) is a code generator, written in Python, which allows the generation of the code for a chosen physics process in many languages. There is an ongoing effort to reengineer the Madgraph5 event generator, and port the matrix element calculation from the current Fortran version to faster implementations using C++ based programming models on CPUs and GPUs. For a complete overview of the MadGraph4GPU project<\/a> and its codebase<\/a>, see the highlighted links.<\/p>\n\n\n\nProfiling MadGraph Kernels<\/h2>\n\n\n\n
We can use the NVIDIA NSight software suite to profile GPU kernels in the CUDA version. First, we must clone the code repository, and build a CUDA-enabled standalone MadGraph application to profile. The profiling is performed on a g g > t t g g<\/code> process.<\/p>\n\n\n\nCreate a shell script that builds the MadGraph application and runs the profiler. This shell script is submitted to a HTCondor GPU node and run inside the singularity container we built earlier. An example of a such a script is given below.<\/p>\n\n\n\n
origdir=$(pwd)\ngit clone https:\/\/github.com\/madgraph5\/madgraph4gpu.git\ncd madgraph4gpu\nexport MADGRAPH4GPU_HOME=`pwd`\ncd epochX\/cudacpp\/gg_ttgg.sa\/SubProcesses\/P1_Sigma_sm_gg_ttxgg\nsed -i 's\/###CUFLAGS+= --maxrregcount 128\/CUFLAGS+= --maxrregcount 255\/' makefile\nmake cleanall\n# Builds standalone MG application\nmake BACKEND=cuda FPTYPE=d HELINL=0\n# Run profiler on standlone app \n# Args: {num_GPU_threads} {num_GPU_regs_per_thread} {num_iterations}\nncu --set full -o $origdir\/report .\/check_cuda.exe -p 216 256 1\n<\/code><\/pre>\n\n\n\nThen, we can use the job submission script shown before<\/a> to submit a profiling job.<\/p>\n\n\n\nIf the job runs successfully, you should see three log files in your working directory in addition to a .ncu-rep<\/code> report file that is tranferred from the HTCondor node after the job ends. Download the report file to your local computer, or another machine with a working Nsight Compute GUI for analy<\/p>\n\n\n\nAnalysis with NSight Compute<\/h2>\n\n\n\n
On opening the report, you should see a screen that looks something like this:<\/p>\n\n\n\n![\"NVIDIA](\"https:\/\/pages.hep.wisc.edu\/cms\/comp\/nsight-home-screen.png\")
<\/figure>\n\n\n\nThis screen lists all the kernels that were profiled by the application, along with a couple of superficial metrics like duration, compute throughput, memory througput, number of registers used per thread, launch configuration, and so on. To see more kernel-specific details, click on the row with that kernel.<\/p>\n\n\n\n
A much more detailed report screen should appear with multiple sections.<\/p>\n\n\n\n![\"An](\"https:\/\/pages.hep.wisc.edu\/cms\/comp\/nsight-report-page.png\")
<\/figure>\n\n\n\nIt is useful to begin by analyzing the data in the GPU Speed on Light section. This tab provides a high-level overview of the throughput for compute and memory resources of the GPU. For each unit, the throughput reports the achieved percentage of utilization with respect to the theoretical maximum. A Roofline plot for the kernel should also appear, which indicates how efficently single and double precision operations are performed by the GPU.<\/p>\n\n\n\n
At the end of each section, warnings and profiling advice indicate potential performance optimization you can perform to improve kernel behaviour across a wide range of GPU units. For example, in the kernel chosen for this example, memory is being more heavily utilized than compute, which is indicative of poor memory accesses or insufficient utilization of compute units within the GPU. The compute and memory workload analysis tabs are helpful to pinpoint the exact cause for this warning. The memory workload analysis tab also includes a schematic that shows the various logical and physical memory accesses that the kernel performs over its lifetime.<\/p>\n\n\n\n
Scheduler and Warp Stall tabs can be used to determine if we are using the appropriate kernel launch configuration that prioritizes efficient memory and compute utilization. Sometimes, memory and compute bottlenecks in the code can cause the GPU to sit idle for longer than needed. The source section can identify the exact line of SASS as well as the kernel source code that is the cause of the bottleneck. Either click on the Source tab, or navigate to the Source Counters section in the main profiling report to identify the top five warp stall locations in the source code. To obtain meaningful data from this section, compile the kernel with the -lineinfo<\/code> option.<\/p>\n\n\n\n![\"NVIDIA](\"https:\/\/pages.hep.wisc.edu\/cms\/comp\/nsight-source-page.png\")
<\/figure>\n\n\n\nFurther Reading<\/h2>\n\n\n\n
This is by no means an exhaustive guide to learning to use NSight or profiling tools. It is impossible for a single guide to cover all possible results you will encounter or edge cases that you might have to work around. Profiling is incredibly \u201chacky\u201d and requires a lot of time and patience to understand the inner workings of the GPU and the code to meaningfully interpret the data and suggestions provided by the profiler. Listed below are some more resources that are a good starting point once you\u2019re more comfortable with CUDA and GPU hardware.<\/p>\n\n\n\n
\n- Official NVIDIA Profiling Guide<\/a><\/li>\n\n\n\n
- A Hands-On NSight training exercise<\/a><\/li>\n<\/ul>\n\n\n\n
\n\n\n\nAppendix<\/h1>\n\n\n\n
Paste the contents below into a file with the extension .sif<\/code>. This creates a singularity definition file that you can use to build a singularity image to run GPU jobs on the CHTC cluster.<\/p>\n\n\n\nBootstrap: docker\nFrom: centos:7\n\n%post\n# Install required dependencies\nyum -y update\nyum -y install wget sudo cmake which git\nyum -y install centos-release-scl\n\n\n# Install GCC 10.2 and related packages\nyum -y install devtoolset-10-gcc devtoolset-10-gcc-c++ devtoolset-10-gcc-gfortran\n\n# TODO: Update Drivers and CUDA links\n# Download and install NVIDIA drivers\nwget https:\/\/us.download.nvidia.com\/XFree86\/Linux-x86_64\/525.116.04\/NVIDIA-Linux-x86_64-525.116.04.run\nchmod +x NVIDIA-Linux-x86_64-525.116.04.run\n.\/NVIDIA-Linux-x86_64-525.116.04.run -s --no-kernel-module\n\n# Download and install CUDA 12.1\nwget https:\/\/developer.download.nvidia.com\/compute\/cuda\/12.1.1\/local_installers\/cuda_12.1.1_530.30.02_linux.run\nchmod +x cuda_12.1.1_530.30.02_linux.run\n.\/cuda_12.1.1_530.30.02_linux.run --silent --toolkit --override\n\n# Clean up\nrm NVIDIA-Linux-x86_64-525.116.04.run cuda_12.1.1_530.30.02_linux.run\n\n%environment\n# Set environment variables for NVIDIA drivers and CUDA\n\nscl enable devtoolset-10 'echo \"GCC 10.2 enabled\"'\nexport PATH=\"\/opt\/rh\/devtoolset-10\/root\/bin:${PATH}\"\nexport PATH=\"\/usr\/local\/nvidia\/bin:${PATH}\"\nexport LD_LIBRARY_PATH=\"\/usr\/local\/nvidia\/lib64:${LD_LIBRARY_PATH}\"\nexport PATH=\"\/usr\/local\/cuda-12.1\/bin:${PATH}\"\nexport LD_LIBRARY_PATH=\"\/usr\/local\/cuda-12.1\/lib64:${LD_LIBRARY_PATH}\"\nexport FC=$(which gfortran)\nexport CC=$(which gcc)\nexport CXX=$(which g++)\n<\/code><\/pre>\n","protected":false},"excerpt":{"rendered":"Setting up your working environment Local Installation Since we\u2019re using GPUs on CHTC nodes, no software installation on your own device is required. (n.b. This is not entirely true \u2014 NSight Compute, the profiling tool we\u2019re going to use has a GUI that cannot be used on HPC systems directly. One has to either download …<\/p>\n","protected":false},"author":3,"featured_media":0,"parent":27,"menu_order":0,"comment_status":"closed","ping_status":"closed","template":"","meta":{"_uw_migration_status":"in-progress","_uw_gutenberg_post_content_before_migration":"","_uw_seo_meta_title":"","_uw_seo_meta_description":"","_uw_seo_twitter_card_type":"summary_large_image","_uw_seo_meta_image":"","_uw_seo_meta_image_url":"","_uw_seo_meta_image_sizes":[],"_uw_seo_custom_meta_tags":[],"footnotes":""},"class_list":["post-109","page","type-page","status-publish","hentry"],"_links":{"self":[{"href":"https:\/\/www.hep.wisc.edu\/cms-comp\/wp-json\/wp\/v2\/pages\/109","targetHints":{"allow":["GET"]}}],"collection":[{"href":"https:\/\/www.hep.wisc.edu\/cms-comp\/wp-json\/wp\/v2\/pages"}],"about":[{"href":"https:\/\/www.hep.wisc.edu\/cms-comp\/wp-json\/wp\/v2\/types\/page"}],"author":[{"embeddable":true,"href":"https:\/\/www.hep.wisc.edu\/cms-comp\/wp-json\/wp\/v2\/users\/3"}],"replies":[{"embeddable":true,"href":"https:\/\/www.hep.wisc.edu\/cms-comp\/wp-json\/wp\/v2\/comments?post=109"}],"version-history":[{"count":2,"href":"https:\/\/www.hep.wisc.edu\/cms-comp\/wp-json\/wp\/v2\/pages\/109\/revisions"}],"predecessor-version":[{"id":114,"href":"https:\/\/www.hep.wisc.edu\/cms-comp\/wp-json\/wp\/v2\/pages\/109\/revisions\/114"}],"up":[{"embeddable":true,"href":"https:\/\/www.hep.wisc.edu\/cms-comp\/wp-json\/wp\/v2\/pages\/27"}],"wp:attachment":[{"href":"https:\/\/www.hep.wisc.edu\/cms-comp\/wp-json\/wp\/v2\/media?parent=109"}],"curies":[{"name":"wp","href":"https:\/\/api.w.org\/{rel}","templated":true}]}}