Basics

// Minimal CUDA program
#include <cuda_runtime.h>
#include <cstdio>

__global__ void hello() {
    printf("Hello from thread %d\n", threadIdx.x);
}

int main() {
    hello<<<1, 4>>>();          // 1 block, 4 threads
    cudaDeviceSynchronize();   // wait for GPU to finish
    return 0;
}

__host__ __device__ float clamp(float v, float lo, float hi) {
    return fminf(fmaxf(v, lo), hi);   // works on CPU and GPU
}

__constant__ float kPi = 3.14159265f;   // lives in constant cache
__device__ int gCounter = 0;            // global device variable

__global__ void demo() {
    __shared__ float tile[256];          // block-local fast memory
    tile[threadIdx.x] = kPi * threadIdx.x;
}

kernel<<<gridDim, blockDim, sharedBytes, stream>>>(args...);

dim3 grid(128, 1, 1);
dim3 block(256, 1, 1);
myKernel<<<grid, block, 0, myStream>>>(d_data, n);

float4 a = make_float4(1.0f, 2.0f, 3.0f, 4.0f);
float dot = a.x*a.x + a.y*a.y + a.z*a.z + a.w*a.w;

#include <cuda_runtime.h>

int devCount;
cudaGetDeviceCount(&devCount);

cudaDeviceProp prop;
cudaGetDeviceProperties(&prop, 0);   // device 0

printf("Name:          %s\n",  prop.name);
printf("SM count:      %d\n",  prop.multiProcessorCount);
printf("Max threads/block: %d\n", prop.maxThreadsPerBlock);
printf("Shared mem/block:  %zu bytes\n", prop.sharedMemPerBlock);
printf("Warp size:     %d\n",  prop.warpSize);            // always 32
printf("Global mem:    %zu MB\n", prop.totalGlobalMem >> 20);
printf("Compute cap:   %d.%d\n", prop.major, prop.minor);

cudaSetDevice(0);        // select GPU 0 for this host thread
int dev;
cudaGetDevice(&dev);     // query current device

// Peer access (multi-GPU)
int canAccess;
cudaDeviceCanAccessPeer(&canAccess, 0, 1);  // can GPU 0 access GPU 1?
if (canAccess) cudaDeviceEnablePeerAccess(1, 0);