93
edits
Changes
→Assignment 3
for (int j = 0; j < height; j++) {
for (int k = 0; k < width; k++) {
std::cout << std::setw(4) << (int)img[idx(k, j, w, h, i)];
}
std::cout << std::endl;
}
std::cout << "Trying to read " << filename << std::endl;
h = cimg.height();
int size = w * h * cimg.spectrum();
for (int i = 0; i < size; i++) {
img[i] = cimg[i];
__global__ void rot90(floatPX_TYPE* src, floatPX_TYPE* dst, int src_w, int src_h, int z) {
int k = blockIdx.x * blockDim.x + threadIdx.x;
// Allocate device memory for src and dst
std::cout << "Allocating device memory ..." << std::endl;
cudaMalloc((void**)&d_src, w * h * sizeof(floatPX_TYPE) * 3); cudaMalloc((void**)&d_dst, w * h * sizeof(floatPX_TYPE) * 3);
// Copy h_src to d_src
std::cout << "Copying source image to device ..." << std::endl;
cudaMemcpy(d_src, h_src, w * h * sizeof(floatPX_TYPE) * 3, cudaMemcpyHostToDevice);
// Launch grid 3 times (one grid per colour channel)
// Copy d_dst to h_dst
std::cout << "Copying rotated image to host ..." << std::endl;
cudaMemcpy(h_dst, d_dst, w * h * sizeof(floatPX_TYPE) * 3, cudaMemcpyDeviceToHost);
// Dealocate memory
// Allocate device memory for src and dst
std::cout << "Allocating device memory ..." << std::endl;
cudaMalloc((void**)&d_src, w * h * sizeof(floatPX_TYPE) * 3); cudaMalloc((void**)&d_dst, w * h * sizeof(floatPX_TYPE) * 3);
// Copy h_src to d_src
std::cout << "Copying source image to device ..." << std::endl;
cudaMemcpy(d_src, h_src, w * h * sizeof(floatPX_TYPE) * 3, cudaMemcpyHostToDevice);
// Rotate image 6 x 2 times, copying result back to host each time
// Copy d_dst to h_dst
std::cout << "Copying result to host ..." << std::endl;
cudaMemcpy(h_dst, d_dst, w * h * sizeof(floatPX_TYPE) * 3, cudaMemcpyDeviceToHost);
// Rotate again
// Copy d_src to h_src
cudaMemcpy(h_src, d_src, w * h * sizeof(floatPX_TYPE) * 3, cudaMemcpyDeviceToHost);
std::cout << "Copying result to host ..." << std::endl;
}
;Timeline Results
For each run, I list the 4 operations that took the most amount of time. For a tiny image, allocating memory source and destination variables on the device took the longest amount of time, but still, it took well under half a second. cudaMalloc It took the same short amount of time for every case however. Initializing the CImg variable from the .jpg file quickly became the biggest issue. This operation is CPU bound, and is dependent on the logic of ImageMagick. Copying the rotated image back to the host (cudaMemcpy) starts to become a hot spot as well between the large and huge sized image is a noticeable increase.
[[File:Summary-2.png]]
=== Assignment 3 ===
==== "Register" Index ====
For my first attempt at optimization, I thought maybe, just maybe, the index calculations were being performed from within global memory:
<nowiki>
__global__ void rot90(PX_TYPE* src, PX_TYPE* dst, int src_w, int src_h, int z) {
int k = blockIdx.x * blockDim.x + threadIdx.x;
int j = blockIdx.y * blockDim.y + threadIdx.y;
if (k < src_w && j < src_h)
dst[(src_h - 1 - j) + k * src_h + src_w * src_h * z] = src[threadIdx.x + threadIdx.y * src_w + src_w * src_h * z];
}</nowiki>
So I declared two register variables and determined the indexes prior:
<nowiki>
__global__ void rot90(PX_TYPE* src, PX_TYPE* dst, int src_w, int src_h, int z) {
int k = blockIdx.x * blockDim.x + threadIdx.x;
int j = blockIdx.y * blockDim.y + threadIdx.y;
int d = (src_h - 1 - j) + k * src_h + src_w * src_h * z;
int s = threadIdx.x + threadIdx.y * src_w + src_w * src_h * z;
if (k < src_w && j < src_h)
dst[d] = src[s];
}</nowiki>
This only had a slightly negative effect. Although, such a small difference may have been due to the luck of the run:
[[File:Summary-4.png]]
==== Unsigned Char vs. Float ====
The first real improvement came from changing PX_TYPE from float back to unsigned char, as used in the serial version. Unsigned char is good enough for all .jpg colour values (255). GPUs are designed to perform operations on floating point numbers, however, we are not performing any calculations outside of the indexing. The performance of the kernel was the same for float or unsigned char. We copy the source image to device once, and back to the host 12 times, making size relevant.
{| class="wikitable"
|+Size Comparison
|-
|
|Unsigned Char
|Float
|-
|Tiny_Shay.jpg
|1.93 KB
|7.73 KB
|-
|Medium_Shay.jpg
|5.71 MB
|22.8 MB
|-
|Large_Shay.jpg
|22.8 MB
|91.4 MB
|-
|Huge_Shay.jpg
|91.4 MB
|365 MB
|}
This saves almost one second worth of latency for the largest file, bringing cudaMemcpy down to about the same time as the kernel execution:
[[File:Summary-5.png]]
==== Shared Memory ====
I could not think of how to utilize shared memory for this application. No calculations are being performed. Copying to shared memory would be an additional operation, as one write to global memory is required either way. By copying a small chunk of the source image to shared memory to improve read time, the indexing logic would no longer work.
==== Constant Memory ====
Utilizing constant memory for the source image was something I wanted to try. The largest unsigned char file of 91.4 MB seemed affordable, and we do not write to it.
Since it's required to use a constant value when declaring the size of the host variable, I needed to define the size of the largest file and use that for all files:
<nowiki>#define SRC_MAX_SIZE 95883264
__constant__ PX_TYPE d_src[SRC_MAX_SIZE];</nowiki>
Copy the actual number of elements over:
<nowiki>// Copy h_src to d_src
std::cout << "Copying source image to device ..." << std::endl;
cudaMemcpyToSymbol(d_src, h_src, w * h * sizeof(PX_TYPE) * 3);</nowiki>
Compiling gave an error saying 91MB is too much memory to use:
<nowiki>CUDACOMPILE : ptxas error : File uses too much global constant data (0x5b71000 bytes, 0x10000 max)</nowiki>
The only example file that fit and compiled was Tiny_Shay.jpg, which there is no point in improving.