CUDA（19）之PTX优化原理

article/2025/9/13 11:46:42

摘要

本文主要讲述CUDA中的PTX的原理实现和分析。

1. 不作优化的代码实现

Nvidia GTX 1050, CUDA 8.0测试代码如下：

#include "cuda_runtime.h"
#include "device_launch_parameters.h"#include <stdio.h>__global__ void gpu(int *d_ptr, int length) {int elemID = blockIdx.x * blockDim.x + threadIdx.x;for (int innerloops = 0; innerloops < 100000; innerloops++) {if (elemID < length) {//unsigned int laneid;d_ptr[elemID] = elemID % 32;}}
}void valid(int *h_ptr, int length) {for (int elemID = 0; elemID<length; elemID++) {h_ptr[elemID] = elemID % 32;}
}int main(int argc, char **argv) {const int N = 1000;int *d_ptr;cudaMalloc(&d_ptr, N * sizeof(int));int *h_ptr;cudaMallocHost(&h_ptr, N * sizeof(int));//start timingfloat time_elapsed = 0;cudaEvent_t start, stop;cudaEventCreate(&start);cudaEventCreate(&stop);cudaEventRecord(start, 0);// GPU kernel without PTXdim3 cudaBlockSize(256, 1, 1);dim3 cudaGridSize((N + cudaBlockSize.x - 1) / cudaBlockSize.x, 1, 1);gpu << <cudaGridSize, cudaBlockSize >> >(d_ptr, N);cudaGetLastError();cudaDeviceSynchronize();// Finish timingcudaEventRecord(stop, 0);cudaEventSynchronize(start);cudaEventSynchronize(stop);cudaEventElapsedTime(&time_elapsed, start, stop);// Printprintf("Time Used on GPU:%f(ms)\n", time_elapsed);// CPU (results for validate)valid(h_ptr, N);int *h_d_ptr;cudaMallocHost(&h_d_ptr, N * sizeof(int));cudaMemcpy(h_d_ptr, d_ptr, N * sizeof(int), cudaMemcpyDeviceToHost);bool bValid = true;for (int i = 0; i<N && bValid; i++) {if (h_ptr[i] != h_d_ptr[i]) {bValid = false;}}printf("Test %s.\n", bValid ? "Successful" : "Failed");cudaFree(d_ptr);cudaFreeHost(h_ptr);cudaFreeHost(h_d_ptr);return bValid ? EXIT_SUCCESS : EXIT_FAILURE;
}

Nvidia GTX 1050, CUDA 8.0测试结果如下：

Nvidia 780Ti, CUDA 7.5测试代码如下：

#include <stdio.h>
#include <assert.h>
#include <cuda_runtime.h>
#include <helper_functions.h>
#include <helper_cuda.h>__global__ void gpu(int *d_ptr, int length){int elemID = blockIdx.x * blockDim.x + threadIdx.x;for(int innerloops = 0; innerloops < 100000; innerloops++){if (elemID < length){//unsigned int laneid;d_ptr[elemID] = elemID % 32;}}
}void valid(int *h_ptr, int length){for (int elemID=0; elemID<length; elemID++){h_ptr[elemID] = elemID % 32;}
}int main(int argc, char **argv){const int N = 1000;int *d_ptr;checkCudaErrors(cudaMalloc(&d_ptr, N * sizeof(int)));int *h_ptr;checkCudaErrors(cudaMallocHost(&h_ptr, N * sizeof(int)));//start timingfloat time_elapsed=0;cudaEvent_t start,stop;cudaEventCreate(&start);    cudaEventCreate(&stop);cudaEventRecord( start,0);// GPU kernel without PTXdim3 cudaBlockSize(256,1,1);dim3 cudaGridSize((N + cudaBlockSize.x - 1) / cudaBlockSize.x, 1, 1);gpu<<<cudaGridSize, cudaBlockSize>>>(d_ptr, N);checkCudaErrors(cudaGetLastError());checkCudaErrors(cudaDeviceSynchronize());// Finish timingcudaEventRecord(stop,0);    cudaEventSynchronize(start);    cudaEventSynchronize(stop);   cudaEventElapsedTime(&time_elapsed,start,stop);// Printprintf("Time Used on GPU:%f(ms)\n",time_elapsed);// CPU (results for validate)valid(h_ptr, N);int *h_d_ptr;checkCudaErrors(cudaMallocHost(&h_d_ptr, N *sizeof(int)));checkCudaErrors(cudaMemcpy(h_d_ptr, d_ptr, N *sizeof(int), cudaMemcpyDeviceToHost));bool bValid = true;for (int i=0; i<N && bValid; i++){if (h_ptr[i] != h_d_ptr[i]){bValid = false;}}printf("Test %s.\n", bValid ? "Successful" : "Failed");checkCudaErrors(cudaFree(d_ptr));checkCudaErrors(cudaFreeHost(h_ptr));checkCudaErrors(cudaFreeHost(h_d_ptr));return bValid ? EXIT_SUCCESS: EXIT_FAILURE;
}

Nvidia 780Ti, CUDA 7.5测试结果如下：

2. PTX

PTX主要特点是预计算的索引值放到L1缓存，实现内存操作的预取优化，达到本案例中性能的大幅提升。

Nvidia 780Ti, CUDA 7.5测试代码如下：

#include <stdio.h>
#include <assert.h>
#include <cuda_runtime.h>
#include <helper_functions.h>
#include <helper_cuda.h>__global__ void gpu_ptx(int *d_ptr, int length){int elemID = blockIdx.x * blockDim.x + threadIdx.x;for(int innerloops = 0; innerloops < 100000; innerloops++){if (elemID < length){unsigned int laneid;asm("mov.u32 %0, %%laneid;" : "=r"(laneid)); // 索引缓存d_ptr[elemID] = laneid;}}
}void valid(int *h_ptr, int length){for (int elemID=0; elemID<length; elemID++){h_ptr[elemID] = elemID % 32;}
}int main(int argc, char **argv){const int N = 1000;int *d_ptr;checkCudaErrors(cudaMalloc(&d_ptr, N * sizeof(int)));int *h_ptr;checkCudaErrors(cudaMallocHost(&h_ptr, N * sizeof(int)));//start timingfloat time_elapsed=0;cudaEvent_t start,stop;cudaEventCreate(&start);    cudaEventCreate(&stop);cudaEventRecord( start,0);// GPU kernel using PTXdim3 cudaBlockSize(256,1,1);dim3 cudaGridSize((N + cudaBlockSize.x - 1) / cudaBlockSize.x, 1, 1);gpu_ptx<<<cudaGridSize, cudaBlockSize>>>(d_ptr, N);checkCudaErrors(cudaGetLastError());checkCudaErrors(cudaDeviceSynchronize());// Finish timingcudaEventRecord(stop,0);    cudaEventSynchronize(start);    cudaEventSynchronize(stop);   cudaEventElapsedTime(&time_elapsed,start,stop);// Printprintf("Time Used on GPU:%f(ms)\n",time_elapsed);// CPU (results for validate)valid(h_ptr, N);int *h_d_ptr;checkCudaErrors(cudaMallocHost(&h_d_ptr, N *sizeof(int)));checkCudaErrors(cudaMemcpy(h_d_ptr, d_ptr, N *sizeof(int), cudaMemcpyDeviceToHost));bool bValid = true;for (int i=0; i<N && bValid; i++){if (h_ptr[i] != h_d_ptr[i]){bValid = false;}}printf("Test %s.\n", bValid ? "Successful" : "Failed");checkCudaErrors(cudaFree(d_ptr));checkCudaErrors(cudaFreeHost(h_ptr));checkCudaErrors(cudaFreeHost(h_d_ptr));return bValid ? EXIT_SUCCESS: EXIT_FAILURE;
}

Nvidia 780Ti, CUDA 7.5测试结果如下：

Nvidia GTX 1050, CUDA 8.0测试代码如下：

#include "cuda_runtime.h"
#include "device_launch_parameters.h"#include <stdio.h>__global__ void gpu_ptx(int *d_ptr, int length) {int elemID = blockIdx.x * blockDim.x + threadIdx.x;for (int innerloops = 0; innerloops < 100000; innerloops++) {if (elemID < length) {unsigned int laneid;asm("mov.u32 %0, %%laneid;" : "=r"(laneid)); // 索引缓存d_ptr[elemID] = laneid;}}
}void valid(int *h_ptr, int length) {for (int elemID = 0; elemID<length; elemID++) {h_ptr[elemID] = elemID % 32;}
}int main(int argc, char **argv) {const int N = 1000;int *d_ptr;cudaMalloc(&d_ptr, N * sizeof(int));int *h_ptr;cudaMallocHost(&h_ptr, N * sizeof(int));//start timingfloat time_elapsed = 0;cudaEvent_t start, stop;cudaEventCreate(&start);cudaEventCreate(&stop);cudaEventRecord(start, 0);// GPU kernel using PTXdim3 cudaBlockSize(256, 1, 1);dim3 cudaGridSize((N + cudaBlockSize.x - 1) / cudaBlockSize.x, 1, 1);gpu_ptx << <cudaGridSize, cudaBlockSize >> >(d_ptr, N);cudaGetLastError();cudaDeviceSynchronize();// Finish timingcudaEventRecord(stop, 0);cudaEventSynchronize(start);cudaEventSynchronize(stop);cudaEventElapsedTime(&time_elapsed, start, stop);// Printprintf("Time Used on GPU:%f(ms)\n", time_elapsed);// CPU (results for validate)valid(h_ptr, N);int *h_d_ptr;cudaMallocHost(&h_d_ptr, N * sizeof(int));cudaMemcpy(h_d_ptr, d_ptr, N * sizeof(int), cudaMemcpyDeviceToHost);bool bValid = true;for (int i = 0; i<N && bValid; i++) {if (h_ptr[i] != h_d_ptr[i]) {bValid = false;}}printf("Test %s.\n", bValid ? "Successful" : "Failed");cudaFree(d_ptr);cudaFreeHost(h_ptr);cudaFreeHost(h_d_ptr);return bValid ? EXIT_SUCCESS : EXIT_FAILURE;
}

Nvidia GTX 1050, CUDA 8.0测试结果如下：