【发布时间】：2011-10-05 00:52:11
【问题描述】：
我正在使用 CUDA 在 GPU (Fermi) 上进行有限差分计算（模板计算）。当我使用 CUDA 分析器测试我的代码时，我发现占用的是0.333。在我订购计算并将占用率增加到0.677 后，内核的执行时间没有减少而是增加了。换句话说，当占用率增加1/3 时，性能会有所下降。
我的问题是：
内核的性能是否依赖于计算而不考虑占用？
【问题讨论】：
标签： performance cuda
【发布时间】：2011-10-05 00:52:11
【问题描述】：
答案是“取决于”，这取决于您的工作负载的特征以及您如何定义性能。一般来说，如果您的瓶颈是数学吞吐量，那么您通常可以使用较低的占用率（12.5%-33%），但如果您的瓶颈是内存，那么您通常需要更高的占用率（66% 或更高）。这只是经验法则，不是绝对规则。大多数内核位于中间的某个位置，但也有两个极端的例外。
占用率是内核中一次可以处于活动状态的最大线程数（受每个线程或其他资源的寄存器计数限制）除以 GPU 在不受其他资源限制时可以激活的最大线程数。活跃意味着线程已分配硬件资源并可用于调度，而不是它有任何指令在给定时钟周期执行。
为一个线程发出指令i后，该线程的指令i+1可能无法立即运行，如果它依赖于指令的结果我。如果该指令是数学指令，则结果将在几个时钟周期内可用。如果它是一个内存加载指令，它可能是 100 个周期。 GPU 不会等待，而是会从其他满足依赖关系的线程发出指令。
因此，如果您主要从事数学运算，则只需要几个（在 GPU 术语中很少；在 CPU 上会被认为有很多）线程来隐藏数学指令的几个延迟周期，这样您就可以侥幸逃脱入住率低。但是如果你有大量的内存流量，你需要更多的线程来确保它们中的一些在每个周期都准备好执行，因为每个线程都会花费大量时间“休眠”等待内存操作完成。
如果您为增加占用率所做的算法更改也增加了每个线程上完成的工作量，并且如果您已经有足够的线程来保持 GPU 忙碌，那么更改只会减慢您的速度。增加占用率只会提高性能，直到您有足够的线程来保持 GPU 忙碌。
【讨论】：
Jesse Hall 已经回答了你的问题，所以我将限制自己来补充他的回答。
为了最大化算法性能，占用并不是唯一需要考虑的品质因数，它通常与执行时间一致。我建议看一看 Vasily Volkov 的指导性 GTC2010 演示：
Better Performance at Lower Occupancy
下面，我提供了一个简单的示例，灵感来自上述演示文稿的第二部分。
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>
#define BLOCKSIZE 512
//#define DEBUG
/*******************/
/* iDivUp FUNCTION */
/*******************/
int iDivUp(int a, int b) { return ((a % b) != 0) ? (a / b + 1) : (a / b); }
/********************/
/* CUDA ERROR CHECK */
/********************/
#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
inline void gpuAssert(cudaError_t code, char *file, int line, bool abort=true)
{
if (code != cudaSuccess)
{
fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
if (abort) exit(code);
}
}
/***********************************************/
/* MEMCPY1 - EACH THREAD COPIES ONE FLOAT ONLY */
/***********************************************/
__global__ void memcpy1(float *src, float *dst, unsigned int N)
{
const int tid = threadIdx.x + blockIdx.x * blockDim.x;
if (tid < N) {
float a0 = src[tid];
dst[tid] = a0;
}
}
/*******************************************/
/* MEMCPY2 - EACH THREAD COPIES TWO FLOATS */
/*******************************************/
__global__ void memcpy2(float *src, float *dst, unsigned int N)
{
const int tid = threadIdx.x + blockIdx.x * (2 * blockDim.x);
if (tid < N) {
float a0 = src[tid];
float a1 = src[tid + blockDim.x];
dst[tid] = a0;
dst[tid + blockDim.x] = a1;
}
}
/********************************************/
/* MEMCPY4 - EACH THREAD COPIES FOUR FLOATS */
/********************************************/
__global__ void memcpy4(float *src, float *dst, unsigned int N)
{
const int tid = threadIdx.x + blockIdx.x * (4 * blockDim.x);
if (tid < N) {
float a0 = src[tid];
float a1 = src[tid + blockDim.x];
float a2 = src[tid + 2 * blockDim.x];
float a3 = src[tid + 3 * blockDim.x];
dst[tid] = a0;
dst[tid + blockDim.x] = a1;
dst[tid + 2 * blockDim.x] = a2;
dst[tid + 3 * blockDim.x] = a3;
}
}
/***********************************************/
/* MEMCPY4_2 - EACH THREAD COPIES FOUR FLOATS2 */
/***********************************************/
__global__ void memcpy4_2(float2 *src, float2 *dst, unsigned int N)
{
const int tid = threadIdx.x + blockIdx.x * (4 * blockDim.x);
if (tid < N/2) {
float2 a0 = src[tid];
float2 a1 = src[tid + blockDim.x];
float2 a2 = src[tid + 2 * blockDim.x];
float2 a3 = src[tid + 3 * blockDim.x];
dst[tid] = a0;
dst[tid + blockDim.x] = a1;
dst[tid + 2 * blockDim.x] = a2;
dst[tid + 3 * blockDim.x] = a3;
}
}
/********/
/* MAIN */
/********/
void main()
{
const int N = 131072;
const int N_iter = 20;
// --- Setting host data and memory space for result
float* h_vect = (float*)malloc(N*sizeof(float));
float* h_result = (float*)malloc(N*sizeof(float));
for (int i=0; i<N; i++) h_vect[i] = i;
// --- Setting device data and memory space for result
float* d_src; gpuErrchk(cudaMalloc((void**)&d_src, N*sizeof(float)));
float* d_dest1; gpuErrchk(cudaMalloc((void**)&d_dest1, N*sizeof(float)));
float* d_dest2; gpuErrchk(cudaMalloc((void**)&d_dest2, N*sizeof(float)));
float* d_dest4; gpuErrchk(cudaMalloc((void**)&d_dest4, N*sizeof(float)));
float* d_dest4_2; gpuErrchk(cudaMalloc((void**)&d_dest4_2, N*sizeof(float)));
gpuErrchk(cudaMemcpy(d_src, h_vect, N*sizeof(float), cudaMemcpyHostToDevice));
// --- Warmup
for (int i=0; i<N_iter; i++) memcpy1<<<iDivUp(N,BLOCKSIZE), BLOCKSIZE>>>(d_src, d_dest1, N);
// --- Creating events for timing
float time;
cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);
/***********/
/* MEMCPY1 */
/***********/
cudaEventRecord(start, 0);
for (int i=0; i<N_iter; i++) {
memcpy1<<<iDivUp(N,BLOCKSIZE), BLOCKSIZE>>>(d_src, d_dest1, N);
#ifdef DEGUB
gpuErrchk(cudaPeekAtLastError());
gpuErrchk(cudaDeviceSynchronize());
#endif
}
cudaEventRecord(stop, 0);
cudaEventSynchronize(stop);
cudaEventElapsedTime(&time, start, stop);
printf("GB/s = %f\n", (1.e-6)*(float)(N*N_iter*sizeof(float))/time);
gpuErrchk(cudaMemcpy(h_result, d_dest1, N*sizeof(int), cudaMemcpyDeviceToHost));
for (int i=0; i<N; i++) if(h_result[i] != h_vect[i]) { printf("Error at i=%i! Host = %i; Device = %i\n", i, h_vect[i], h_result[i]); return; }
/***********/
/* MEMCPY2 */
/***********/
cudaEventRecord(start, 0);
for (int i=0; i<N_iter; i++) {
memcpy2<<<iDivUp(N/2,BLOCKSIZE), BLOCKSIZE>>>(d_src, d_dest2, N);
#ifdef DEGUB
gpuErrchk(cudaPeekAtLastError());
gpuErrchk(cudaDeviceSynchronize());
#endif
}
cudaEventRecord(stop, 0);
cudaEventSynchronize(stop);
cudaEventElapsedTime(&time, start, stop);
printf("GB/s = %f\n", (1.e-6)*(float)(N*N_iter*sizeof(float))/time);
gpuErrchk(cudaMemcpy(h_result, d_dest2, N*sizeof(int), cudaMemcpyDeviceToHost));
for (int i=0; i<N; i++) if(h_result[i] != h_vect[i]) { printf("Error at i=%i! Host = %i; Device = %i\n", i, h_vect[i], h_result[i]); return; }
/***********/
/* MEMCPY4 */
/***********/
cudaEventRecord(start, 0);
for (int i=0; i<N_iter; i++) {
memcpy4<<<iDivUp(N/4,BLOCKSIZE), BLOCKSIZE>>>(d_src, d_dest4, N);
#ifdef DEGUB
gpuErrchk(cudaPeekAtLastError());
gpuErrchk(cudaDeviceSynchronize());
#endif
}
cudaEventRecord(stop, 0);
cudaEventSynchronize(stop);
cudaEventElapsedTime(&time, start, stop);
printf("GB/s = %f\n", (1.e-6)*(float)(N*N_iter*sizeof(float))/time);
gpuErrchk(cudaMemcpy(h_result, d_dest4, N*sizeof(int), cudaMemcpyDeviceToHost));
for (int i=0; i<N; i++) if(h_result[i] != h_vect[i]) { printf("Error at i=%i! Host = %i; Device = %i\n", i, h_vect[i], h_result[i]); return; }
/*************/
/* MEMCPY4_2 */
/*************/
cudaEventRecord(start, 0);
for (int i=0; i<N_iter; i++) {
memcpy4_2<<<iDivUp(N/8,BLOCKSIZE), BLOCKSIZE>>>((float2*)d_src, (float2*)d_dest4_2, N);
#ifdef DEGUB
gpuErrchk(cudaPeekAtLastError());
gpuErrchk(cudaDeviceSynchronize());
#endif
}
cudaEventRecord(stop, 0);
cudaEventSynchronize(stop);
cudaEventElapsedTime(&time, start, stop);
printf("GB/s = %f\n", (1.e-6)*(float)(N*N_iter*sizeof(float))/time);
gpuErrchk(cudaMemcpy(h_result, d_dest4_2, N*sizeof(int), cudaMemcpyDeviceToHost));
for (int i=0; i<N; i++) if(h_result[i] != h_vect[i]) { printf("Error at i=%i! Host = %i; Device = %i\n", i, h_vect[i], h_result[i]); return; }
cudaDeviceReset();
}
以下是上述代码在 GeForce GT540M 和 Kepler K20c 上运行时的性能。
块大小 32
GT540M K20c Tesla C2050
memcpy1 2.3GB/s 13% 28.1GB/s 18% 14.9GB/s 12%
memcpy2 4.4GB/s 13% 41.1GB/s 18% 24.8GB/s 13%
memcpy4 7.5GB/s 13% 54.8GB/s 18% 34.6GB/s 13%
memcpy4_2 11.2GB/2 14% 68.8GB/s 18% 44.0GB7s 14%
块大小 64
GT540M K20c Tesla C2050
memcpy1 4.6GB/s 27% 44.1GB/s 36% 26.1GB/s 26%
memcpy2 8.1GB/s 27% 57.1GB/s 36% 35.7GB/s 26%
memcpy4 11.4GB/s 27% 63.2GB/s 36% 43.5GB/s 26%
memcpy4_2 12.6GB/s 27% 72.8GB/s 36% 49.7GB/s 27%
块大小 128
GT540M K20c Tesla C2050
memcpy1 8.0GB/s 52% 60.6GB/s 78% 36.1GB/s 52%
memcpy2 11.6GB/2 52% 61.6GB/s 78% 44.8GB/s 52%
memcpy4 12.4GB/2 52% 62.2GB/s 78% 48.3GB/s 52%
memcpy4_2 12.5GB/s 52% 61.9GB/s 78% 49.5GB7s 52%
块大小 256
GT540M K20c Tesla C2050
memcpy1 10.6GB/s 80% 61.2GB/s 74% 42.0GB/s 77%
memcpy2 12.3GB/s 80% 66.2GB/s 74% 48.2GB/s 77%
memcpy4 12.4GB/s 80% 66.4GB/s 74% 45.5GB/s 77%
memcpy4_2 12.6GB/s 70% 72.6GB/s 74% 50.8GB/s 77%
块大小 512
GT540M K20c Tesla C2050
memcpy1 10.3GB/s 80% 54.5GB/s 75% 41.6GB/s 75%
memcpy2 12.2GB/s 80% 67.1GB/s 75% 47.7GB/s 75%
memcpy4 12.4GB/s 80% 67.9GB/s 75% 46.9GB/s 75%
memcpy4_2 12.5GB/s 55% 70.1GB/s 75% 48.3GB/s 75%
以上结果表明，如果您正确地利用 指令级并行 (ILP)为每个线程分配更多工作以隐藏延迟。
【讨论】：