University:
Boston University
Course:
MET CS 201 | Introduction to Programming
Academic year:

2024
Views:

52

Pages:

20
Author:

manogat3uex

GPU Computing with CUDA CUDA Libraries - Thrust 1 Outline of lecture ‣ CUDA Libraries ‣ What is Thrust? ‣ Features of thrust ‣ Best practices in thrust 2 CUDA Libraries ‣ NVIDIA has developed several libraries to abstract the user from CUDA Bell, Dalton, Olson. Towards AMG on GPU 3 CUDA Libraries ‣ CUDPP ‣ MAGMA ‣ IMSL ‣ VSIPL ‣ NVML ‣ CUPTI ‣ ..... 4 CUDA Libraries ‣ CUDPP ‣ MAGMA ‣ IMSL ‣ VSIPL ‣ NVML ‣ CUPTI ‣ ..... 4 Thrust - Introduction ‣ Template library for CUDA - Resembles C++ Standard Template Library (STL) - Collection of data parallel primitives ‣ Objectives - Programmer productivity - Encourage generic programming - High performance - Interoperability ‣ Comes with CUDA 4.0 5 Thrust - Introduction ‣ Containers ‐ thrust::host_vector ‐ thrust::device_vector ‣ Algorithms ‐ thrust::sort() ‐ thrust::reduce() ‐ thrust::inclusive_scan() ‣ http://code.google.com/p/thrust/ ‣ Slides from Nathan Bell and Jared Hoberock - NVIDIA 6 Thrust - Containers ‣ Thrust provides two vector containers - host_vector: resides on CPU - device_vector: resides on GPU ‣ Hides cudaMalloc and cudaMemcpy // allocate host vector with two elements thrust::host_vector h_vec(2); // copy host vector to device thrust::device_vector d_vec = h_vec; // manipulate device values from the host d_vec[0] = 13; d_vec[1] = 27; std::cout << "sum: " << d_vec[0] + d_vec[1] << std::endl; // vector memory automatically released w/ free() or cudaFree() 7 Thrust - Iterators ‣ Iterators can be thought as pointers to array elements - They carry other information // allocate device vector thrust::device_vector d_vec(4); d_vec.begin(); // returns iterator at first element of d_vec d_vec.end() // returns iterator one past the last element of d_vec // [begin, end] pair defines a sequence of 4 elements d_vec.begin() d_vec.end() 8 Thrust - Iterators ‣ Use iterators like pointers // allocate device vector thrust::device_vector d_vec(4); thrust::device_vector::iterator begin = d_vec.begin(); *begin = 13; // same as d_vec[0] = 13; int temp = *begin; // same as temp = d_vec[0]; begin++; // advance iterator one position *begin = 25; // same as d_vec[1] = 25; 9 Thrust - Iterators ‣ Important: keep track of your memory space // initialize random values on host thrust::host_vector h_vec(1000); thrust::generate(h_vec.begin(), h_vec.end(), rand); // copy values to device thrust::device_vector d_vec = h_vec; // compute sum on host int h_sum = thrust::reduce(h_vec.begin(), h_vec.end()); // compute sum on device int d_sum = thrust::reduce(d_vec.begin(), d_vec.end()); 10 Thrust - Iterators ‣ Convertible to raw pointers // allocate device vector thrust::device_vector d_vec(4); // obtain raw pointer to device vector’s memory int * ptr = thrust::raw_pointer_cast(&d_vec[0]); // use ptr in a CUDA C kernel my_kernel<<>>(N, ptr); // Note: ptr cannot be dereferenced on the host! 11 Thrust - Iterators ‣ Wrap raw pointers to use in thrust int N = 10; // raw pointer to device memory int * raw_ptr; cudaMalloc((void **) &raw_ptr, N * sizeof(int)); // wrap raw pointer with a device_ptr thrust::device_ptr dev_ptr(raw_ptr); // use device_ptr in thrust algorithms thrust::fill(dev_ptr, dev_ptr + N, (int) 0); // access device memory through device_ptr dev_ptr[0] = 1; // free memory cudaFree(raw_ptr); 12 Thrust - Algorithms ‣ Standard algorithms - Reductions - Transformations - Prefix sums - Sorting ‣ Many have straight analog to STL ‣ You can use your own user defined types 13 Thrust - Algorithms ‣ Reductions #include // declare storage device_vector i_vec = ... device_vector f_vec = ... // sum of integers (equivalent calls) reduce(i_vec.begin(), i_vec.end()); reduce(i_vec.begin(), i_vec.end(), 0, plus()); // sum of floats (equivalent calls) reduce(f_vec.begin(), f_vec.end()); reduce(f_vec.begin(), f_vec.end(), 0.0f, plus()); // maximum of integers reduce(i_vec.begin(), i_vec.end(), 0, maximum()); 14 Thrust - Algorithms ‣ Thrust comes with lots of important built in transformations #include #include #include #include #include #include #include #include //allocate three device_vectors with 10 elements thrust::device_vectorX(10); thrust::device_vector Y(10); thrust::device_vectorZ(10); //initialize X to 0, 1, 2, 3,.... thrust::sequence(X.begin(),X.end()); //compute Y=‐X thrust::transform(X.begin(),X.end(),Y.begin(),thrust::negate()); //fill Z with two s thrust::fill(Z.begin(),Z.end(),2); //compute Y = X mod 2 thrust::transform(X.begin(),X.end(),Z.begin(),Y.begin(),thrust::modulus()); //replace all the ones in Y with tens thrust::replace(Y.begin(),Y.end(),1,10); //print Y thrust::copy(Y.begin(),Y.end(),std::ostream_iterator(std::cout,"\n")); return 0; 15 Thrust - Algorithms ‣ Thrust can do general types and operators using functors struct negate_float2 { __host__ __device__ float2 operator()(float2 a) { return make_float2(‐a.x, ‐a.y); } }; // declare storage device_vector input = ... device_vector output = ... // create functor negate_float2 func; // negate vectors transform(input.begin(), input.end(), output.begin(), func); 16 Thrust - Algorithms // compare x component of two float2 structures struct compare_float2 { __host__ __device__ bool operator()(float2 a, float2 b) { return a.x < b.x; } }; // declare storage device_vector vec = ... // create comparison functor compare_float2 comp; // sort elements by x component sort(vec.begin(), vec.end(), comp); 17 Thrust - Algorithms ‣ Prefix sums #include int data [6] = {1,0,2,2,1,3}; thrust::inclusive_scan(data, data + 6, data);//in‐place scan //data is now {1,1,3,5,6,9} data[2] = data[0] + data[1] + data[2] #include int data [6] = {1,0,2,2,1,3}; thrust::exclusive_scan(data, data + 6, data);//in‐place scan //data is now {0,1,1,3,5,6} data[2] = data[0] + data[1] 18 Thrust - Algorithms ‣ Sorting #include . . . const int N=6; int A [N] = {1,4,2,8,5,7}; thrust::sort(A,A+N); // A is now {1,2,4,5,7,8} 19