Skip to content
Search

CUDA:Example

lastmod 2018-08-17created 2018-02-012 edits

Simple example 

#include <cuda.h>
#include <stdio.h>

static void HandleError(cudaError_t err, const char * file, int line)
{
    if (err != cudaSuccess) {
        printf("%s in %s at line %d\n", cudaGetErrorString(err), file, line);
        exit(EXIT_FAILURE);
    }
}

#define HANDLE_ERROR(err) (HandleError(err, __FILE__, __LINE__))

static int const TEST_BLOCK_SIZE = 10;

__global__ void __add__(int * a, int * b, int * c)
{
    int tid = blockIdx.x;
    if (tid < TEST_BLOCK_SIZE) {
        c[tid] = a[tid] + b[tid];
    }
}

bool checkCuda()
{
    int a[TEST_BLOCK_SIZE], b[TEST_BLOCK_SIZE], c[TEST_BLOCK_SIZE];
    int * dev_a, * dev_b, * dev_c;

    // fill the arrays 'a' and 'b' on the CPU.
    for (int i = 0; i < TEST_BLOCK_SIZE; i++) {
        a[i] = i;
        b[i] = i * 2;
    }

    // allocate the memory on the GPU.
    HANDLE_ERROR(cudaMalloc((void**) &dev_a, TEST_BLOCK_SIZE * sizeof(int)));
    HANDLE_ERROR(cudaMalloc((void**) &dev_b, TEST_BLOCK_SIZE * sizeof(int)));
    HANDLE_ERROR(cudaMalloc((void**) &dev_c, TEST_BLOCK_SIZE * sizeof(int)));

    // copy the arrays 'a' and 'b' to the GPU.
    HANDLE_ERROR(cudaMemcpy(dev_a, a, TEST_BLOCK_SIZE * sizeof(int), cudaMemcpyHostToDevice));
    HANDLE_ERROR(cudaMemcpy(dev_b, b, TEST_BLOCK_SIZE * sizeof(int), cudaMemcpyHostToDevice));

    // Run CUDA method.
    __add__<<<TEST_BLOCK_SIZE, 1>>>(dev_a, dev_b, dev_c);

    // copy the array 'c' back from the GPU to the CPU.
    HANDLE_ERROR(cudaMemcpy(c, dev_c, TEST_BLOCK_SIZE * sizeof(int), cudaMemcpyDeviceToHost));

    // check the results.
    bool result = true;
    for (int i = 0; i < TEST_BLOCK_SIZE; i++) {
        if (c[i] != i * 3) {
            result = false;
            break;
        }
    }

    // free the memory allocated on the GPU
    cudaFree(dev_a);
    cudaFree(dev_b);
    cudaFree(dev_c);

    return result;
}

Array example

This example reads a texture from an image into a matrix on the GPU. 

cudaArray* cu_array;
texture<float, 2> tex;

// 행렬 할당
cudaMallocArray(&cu_array, cudaCreateChannelDesc<float>(), width, height);

// 이미지 데이터를 행렬로 복사
cudaMemcpy(cu_array, image, width*height, cudaMemcpyHostToDevice);

// 행렬을 텍스처에 연결한다.
cudaBindTexture(tex, cu_array);

// 커널을 실행한다
dim3 blockDim(16, 16, 1);
dim3 gridDim(width / blockDim.x, height / blockDim.y, 1);
kernel<<< gridDim, blockDim, 0 >>>(d_odata, width, height);
cudaUnbindTexture(tex);

__global__ void kernel(float* odata, int height, int width)
{
   unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;
   unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;
   float c = texfetch(tex, x, y);
   odata[y*width+x] = c;
}

CUDA Error 

void printCudaError(char const * message)
{
    cudaError_t err = cudaGetLastError();
    std::string format;

    char const * name = cudaGetErrorName(err);
    char const * body = cudaGetErrorString(err);

    if (message == NULL) {
        format = "CUDA [%d/%s] %s\n";
        fprintf(stderr, format.c_str(), err, name, body);
    } else {
        format = "%s [%d/%s] %s\n";
        fprintf(stderr, format.c_str(), message, err, name, body);
    }
}

CUDA Library

This summarizes how to write a library with CUDA and then LINK it with GCC.

First, write the three files below. Note that you must not call __device__ or __global__ functions directly. 

// gpu.cu
#include <cuda.h>
#include <stdio.h>
extern "C" void checkErroreCuda()
{
    cudaError_t err = cudaGetLastError();
    if (err != cudaSuccess) {
        printf("%s\n", cudaGetErrorString(err));
    } else {
        printf("No error.\n");
    }
}

// gpu.h
extern "C" {
void checkErroreCuda();
}

// main.cpp
#include "gpu.h"
int main()
{
    checkErroreCuda();
    return 0;
}

Compile as follows. 

nvcc --cudart static -lib -ccbin g++ -o libgpu.a gpu.cu
g++ -c main.cpp
g++ main.o libgpu.a -L/usr/local/cuda/lib64 -lcuda -lcudart_static

Device Query

Among the CUDA sample projects there is one called deviceQuery. Build and run it to see information about the device.

The code for checking CUDA device information is below. (It was used in the CMake file of the Caffe project.) 

#include <cstdio>

int main()
{
  int count = 0;

  if (cudaSuccess != cudaGetDeviceCount(&count)) {
    return -1;
  }

  if (count == 0) {
    return -1;
  }

  for (int device = 0; device < count; ++device) {
    cudaDeviceProp prop;
    if (cudaSuccess == cudaGetDeviceProperties(&prop, device)) {
      std::printf("%d.%d", prop.major, prop.minor);
    }
  }
  return 0;
}

See also

sitemap