CUDA C編程入門-不同的grid和block大小對CUDA內建的變數的影響，以及如何確定Thread ID

標籤：style   blog   http   color   os   io   strong   for   

測試環境
　　測試的GPU平台為GTX660M，計算能力為3.0
首先介紹一下GPU提供的函數：

int printf(const char *format[, arg, ...]);

　　從核函數格式化輸出到主機，只支援計算能力在2.x及以上的裝置。行為與標準的C相似。這裡我們用於輸出內建變數的值。
核函數調用方式

kernel<<<Dg, Db, Ns, S >>>

　　其中，Dg，類型為dim3，表示grid的每一維大小，計算能力為1.x的Dg.z等於1。表示載入的block數量為Dg各維的乘積。
Db也是Dim3類型，表示為每個block的每一維帶下，表示每個block中線程的數量，等於Db各維的乘積。
Ns為每個Block的共用記憶體的大小，單位為位元組。是一個可選的參數，預設為0。
S類型為cudaStream_t，表示相關的流。是一個可選的參數，預設為0。
CUDA內建的變數有：

• gridDim，類型為dim3，dim3預設初始化為1，表示grid的每一維的大小
• blockDim，類型為dim3，表示block的每一維的大小
• blockIdx，類型為uint3，表示為grid中block的索引
• threadIdx，類型為uint3，表示為block中thread索引
• warpSize，類型為int，表示warp的大小

內建變數的限制
從計算能力3.0的相關硬體描述可以找到以下限制：

• grid維度為3維，每個維度大小限制為：x為2^31-1，y和z為65535
• block維度為3維，每一維度大小限制為：x和y為1024，z為64
• 每個block中線程的最大數為1024，即blockDim中x、y和z的乘積不能超過1024
• warp大小為32

下面我們先搞清楚不同的grid和block大小對CUDA內建的變數的影響
代碼如下，grid為（2，3，1），而block大小為（2，2，1）：

#include "cuda_runtime.h"#include "device_launch_parameters.h"#include <stdio.h>__global__ void MyKernel(){    printf("(%d,%d,%d) (%d, %d, %d) (%d, %d, %d) (%d, %d, %d) %d\n",        gridDim.x, gridDim.y, gridDim.z,         blockDim.x, blockDim.y, blockDim.z,        blockIdx.x, blockIdx.y, blockIdx.z,        threadIdx.x, threadIdx.y, threadIdx.z,        warpSize);}int main(){    cudaError_t cudaStatus;    // Choose which GPU to run on, change this on a multi-GPU system.    cudaStatus = cudaSetDevice(0);    if (cudaStatus != cudaSuccess) {        fprintf(stderr, "cudaSetDevice failed! Do you have a CUDA-capable GPU installed?");    }    printf("gridDim | blockDim | blockIdx | threadIdx | warpSize\n");    dim3 gridSize(2, 3, 1);    dim3 blockSize(2, 2, 1);    // Launch a kernel on the GPU with one thread for each element.    MyKernel<<<gridSize, blockSize>>>();    // Check for any errors launching the kernel    cudaStatus = cudaGetLastError();    if (cudaStatus != cudaSuccess) {        fprintf(stderr, "addKernel launch failed: %s\n",         cudaGetErrorString(cudaStatus));    }    // cudaDeviceSynchronize waits for the kernel to finish, and returns    // any errors encountered during the launch.    cudaStatus = cudaDeviceSynchronize();    if (cudaStatus != cudaSuccess) {        fprintf(stderr, "cudaDeviceSynchronize returned error code %d after launching addKernel!\n", cudaStatus);    }    // cudaDeviceReset must be called before exiting in order for profiling and    // tracing tools such as Nsight and Visual Profiler to show complete traces.    cudaStatus = cudaDeviceReset();    if (cudaStatus != cudaSuccess) {        fprintf(stderr, "cudaDeviceReset failed!");        return 1;    }    return 0;}

注意：需要設定CUDA代碼產生的選項：

執行的結果如下:

　　說明調用核函數輸入的第一個參數設定gridDim，第二個設定blockDim，然後gridDim影響blockIdx的索引值，blockDim影響threadIdx索引值，warpSize大小為32。