GPU 编程 ¶

for (int i = 0; i < N; ++i) {
    C[i] = A[i] + B[i];
}

void CPUFunction() {
    printf("This function is defined to run on the CPU.\n");
}

__global__ void GPUFunction() {
    printf("This function is defined to run on the GPU.\n");
}

int main() {
    CPUFunction();
    GPUFunction<<<1, 1>>>();
    cudaDeviceSynchronize();
}

$ nvcc -arch=sm_70 -o out code.cu -run

i = threadIdx.x + blockIdx.x * blockDim.x

__global__ void kernel(int *a, int N) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    int stride = gridDim.x * blockDim.x;
    for (int i = idx; i < N; i += stride) {
        // a[i] ...
    }
}

dim3 threadsPerBlock(16, 16, 1);
dim3 numberOfBlocks(16, 16, 1);
kernel<<<numberOfBlocks, threadsPerBlock>>>(...);

__global__ void kernel(int **a, int N1, int N2) {
    int idxx = threadIdx.x + blockIdx.x * blockDim.x;
    int stridex = gridDim.x * blockDim.x;
    int idxy = threadIdx.y + blockIdx.y * blockDim.y;
    int stridey = gridDim.y * blockDim.y;
    for (int i = idxx; i < N1; i += stridex) {
        for (int j = idxy; j < N2; j += stridey) {
            // a[i][j] ...
        }
    }
}

cudaError_t err;
err = cudaMallocManaged(&a, N);
if (err != cudaSuccess) {
    printf("Error: %s\n", cudaGetErrorString(err));
}

kernel<<<1, -1>>>();
cudaError_t err;
err = cudaGetLastError;
if (err != cudaSuccess) {
    printf("Error: %s\n", cudaGetErrorString(err));
}

#include <assert.h>
inline cudaError_t checkCuda(cudaError_t result) {
    if (result != cudaSuccess) {
        fprintf(stderr, "CUDA Runtime Error: %s\n", cudaGetErrorString(result));
        assert(result == cudaSuccess);
    }
    return result;
}

int main() {
    ...
    kernel<<<..., ...>>>(...);
    checkCuda(cudaGetLastError());
    checkCuda(cudaDeviceSynchronize());
    ...
}

int *a;
size_t size = n * sizeof(int);
a = (int*)malloc(size) -> cudaMallocManaged(&a, size)

$ nsys profile --stats=true ./out

int deviceId;
cudaGetDevice(&deviceId);
cudaDeviceProp props;
cudaGetDeviceProperties(&props, deviceId);

int deviceId;
int numberOfSMs;

cudaGetDevice(&deviceId);
cudaDeviceGetAttribute(&numberOfSMs, cudaDevAttrMultiProcessorCount, deviceId);

size_t threadsPerBlock = 256;
size_t numberOfBlocks = 32 * numberOfSMs;

int deviceId;
cudaGetDevice(&deviceId);

cudaMemPrefetchAsync(pointerToSomeUMData, size, deviceId);
...
cudaMemPrefetchAsync(pointerToSomeUMData, size, cudaCpuDeviceId);

int *host_a, *device_a;        // Define host-specific and device-specific arrays.
cudaMalloc(&device_a, size);   // `device_a` is immediately available on the GPU.
cudaMallocHost(&host_a, size); // `host_a` is immediately available on CPU, and is page-locked, or pinned.

initializeOnHost(host_a, N);   // No CPU page faulting since memory is already allocated on the host.

// `cudaMemcpy` takes the destination, source, size, and a CUDA-provided variable for the direction of the copy.
cudaMemcpy(device_a, host_a, size, cudaMemcpyHostToDevice);

kernel<<<blocks, threads, 0, someStream>>>(device_a, N);

// `cudaMemcpy` can also copy data from device to host.
cudaMemcpy(host_a, device_a, size, cudaMemcpyDeviceToHost);

verifyOnHost(host_a, N);

cudaFree(device_a);
cudaFreeHost(host_a);          // Free pinned memory like this.

cudaStream_t stream;
cudaStreamCreate(&stream);
kernel<<<number_of_blocks, threads_per_block, 0, stream>>>();
cudaStreamDestroy(stream);

int N = 2<<24;
int size = N * sizeof(int);

int *host_array;
int *device_array;

cudaMallocHost(&host_array, size);               // Pinned host memory allocation.
cudaMalloc(&device_array, size);                 // Allocation directly on the active GPU device.

initializeData(host_array, N);                   // Assume this application needs to initialize on the host.

const int numberOfSegments = 4;                  // This example demonstrates slicing the work into 4 segments.
int segmentN = N / numberOfSegments;             // A value for a segment's worth of `N` is needed.
size_t segmentSize = size / numberOfSegments;    // A value for a segment's worth of `size` is needed.

// For each of the 4 segments...
for (int i = 0; i < numberOfSegments; ++i) {
    // Calculate the index where this particular segment should operate within the larger arrays.
    segmentOffset = i * segmentN;

    // Create a stream for this segment's worth of copy and work.
    cudaStream_t stream;
    cudaStreamCreate(&stream);

    // Asynchronously copy segment's worth of pinned host memory to device over non-default stream.
    cudaMemcpyAsync(&device_array[segmentOffset],  // Take care to access correct location in array.
                    &host_array[segmentOffset],    // Take care to access correct location in array.
                    segmentSize,                   // Only copy a segment's worth of memory.
                    cudaMemcpyHostToDevice,
                    stream);                       // Provide optional argument for non-default stream.

    // Execute segment's worth of work over same non-default stream as memory copy.
    kernel<<<number_of_blocks, threads_per_block, 0, stream>>>(&device_array[segmentOffset], segmentN);

    // `cudaStreamDestroy` will return immediately (is non-blocking), but will not actually destroy stream until
    // all stream operations are complete.
    cudaStreamDestroy(stream);
}