Pasted

#include <cuda_runtime.h>
 
static void GPUHandleError( cudaError_t err, const char *file, const int line ) {
    if (err != cudaSuccess) {
        printf( "%s in %s at line %d\n", cudaGetErrorString( err ), file, line );
        exit( EXIT_FAILURE );
    }
}
 
#define GPU_HANDLE_ERROR( err ) (GPUHandleError( err, __FILE__, __LINE__ ))
 
// 1D to 2D to 1D
unsigned int idx = blockIdx.x * blockDim.x + threadIdx.x; // idx < lx * ly
unsigned int x = idx / sizeY;
unsigned int y = idx % sizeY;
 
// 2D to 1D with Blocks
unsigned int x = blockIdx.x;
unsigned int y = blockIdx.y;
unsigned int offset = x + y * gridDim.x;
 
// 2D to 1D with Blocks and Threads
unsigned int x = threadIdx.x + blockIdx.x * blockDim.x;
unsigned int y = threadIdx.y + blockIdx.y * BlockDim.y;
unsigned int offset = x + y * gridDim.x * gridDim.x;
 
// Shared memory
__shared__ float cache[threadsPerBlock];
unsigned int tid = threadIdx.x + blockIdx.x * blockDim.x;
unsigned int cacheIndex = threadIdx.x;
unsigned int stride = blockDim.x * gridDim.x;
__syncthreads();
 
 
// Other hand
cudaSetDevice(0);
 
unsigned int minGridSize, BLOCK_KERNEL;
cudaOccupancyMaxPotentialBlockSize(&minGridSize, &BLOCK_KERNEL, kernel, 0, size);
unsigned int GRID_KERNEL = (size + BLOCK_KERNEL - 1) / BLOCK_KERNEL; 
 
 
cudaEvent_t start, stop;
cudaEventCreate(&start); cudaEventCreate(&stop);
 
double *d_var, *h_var;
size_t size = lx * ly * sizeof(double);
 
cudaMalloc((void **) &d_var, size);
 
cudaMemcpy(d_var, h_var, size, cudaMemcpyHostToDevice);
cudaDeviceSynchronize();
 
cudaEventRecord(start, 0);
cudaDeviceSynchronize();
 
kernel<<<GRID_KERNEL, BLOCK_KERNEL>>> (...);
cudaDeviceSynchronize();
cudaGetLastError();
 
cudaEventRecord(&stop, 0);
cudaDeviceSynchronize();
 
float timer; // result in ms. /1000 to convert to s.
cudaEventElapsedTime(&timer, start, stop);
 
cudaEventDestroy(start); cudaEventDestroy(stop);
 
cudaMemcpy(h_var, d_var, size, cudaMemcpyDeviceToHost);
cudaDeviceSynchronize();
 
cudaFree(d_var);
cudaDeviceReset();