Skip to content

Commit 8aa85ec

Browse files
committed
added unit 2 lecture quizzes and moved some things around
1 parent 77dfb0d commit 8aa85ec

File tree

5 files changed

+166
-0
lines changed

5 files changed

+166
-0
lines changed

LectureQuizzes/2_18.cu

+23
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,23 @@
1+
#include <stdio.h>
2+
3+
#define NUM_BLOCKS 16
4+
#define BLOCK_WIDTH 1
5+
6+
__global__ void hello()
7+
{
8+
printf("Hello world! I'm a thread in block %d\n", blockIdx.x);
9+
}
10+
11+
12+
int main(int argc,char **argv)
13+
{
14+
// launch the kernel
15+
hello<<<NUM_BLOCKS, BLOCK_WIDTH>>>();
16+
17+
// force the printf()s to flush
18+
cudaDeviceSynchronize();
19+
20+
printf("That's all!\n");
21+
22+
return 0;
23+
}

LectureQuizzes/2_37.cu

+79
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,79 @@
1+
#include <stdio.h>
2+
#include "gputimer.h"
3+
4+
#define NUM_THREADS 1000000
5+
#define ARRAY_SIZE 100
6+
bool atomic = true;
7+
8+
#define BLOCK_WIDTH 1000
9+
10+
void print_array(int *array, int size)
11+
{
12+
printf("{ ");
13+
for (int i = 0; i < size; i++) { printf("%d ", array[i]); }
14+
printf("}\n");
15+
}
16+
17+
__global__ void increment_naive(int *g)
18+
{
19+
// which thread is this?
20+
int i = blockIdx.x * blockDim.x + threadIdx.x;
21+
22+
// each thread to increment consecutive elements, wrapping at ARRAY_SIZE
23+
i = i % ARRAY_SIZE;
24+
g[i] = g[i] + 1;
25+
}
26+
27+
__global__ void increment_atomic(int *g)
28+
{
29+
// which thread is this?
30+
int i = blockIdx.x * blockDim.x + threadIdx.x;
31+
32+
// each thread to increment consecutive elements, wrapping at ARRAY_SIZE
33+
i = i % ARRAY_SIZE;
34+
atomicAdd(& g[i], 1);
35+
}
36+
37+
int main(int argc,char **argv)
38+
{
39+
GpuTimer timer;
40+
if (atomic) {
41+
printf("atomic %d total threads in %d blocks writing into %d array elements\n",
42+
NUM_THREADS, NUM_THREADS / BLOCK_WIDTH, ARRAY_SIZE);
43+
} else {
44+
printf("%d total threads in %d blocks writing into %d array elements\n",
45+
NUM_THREADS, NUM_THREADS / BLOCK_WIDTH, ARRAY_SIZE);
46+
}
47+
48+
// declare and allocate host memory
49+
int h_array[ARRAY_SIZE];
50+
const int ARRAY_BYTES = ARRAY_SIZE * sizeof(int);
51+
52+
// declare, allocate, and zero out GPU memory
53+
int * d_array;
54+
cudaMalloc((void **) &d_array, ARRAY_BYTES);
55+
cudaMemset((void *) d_array, 0, ARRAY_BYTES);
56+
57+
// launch the kernel - comment out one of these
58+
timer.Start();
59+
60+
// Instructions: This program is needed for the next quiz
61+
// uncomment increment_naive to measure speed and accuracy
62+
// of non-atomic increments or uncomment increment_atomic to
63+
// measure speed and accuracy of atomic icrements
64+
if (atomic) {
65+
increment_atomic<<<NUM_THREADS/BLOCK_WIDTH, BLOCK_WIDTH>>>(d_array);
66+
} else {
67+
increment_naive<<<NUM_THREADS/BLOCK_WIDTH, BLOCK_WIDTH>>>(d_array);
68+
}
69+
timer.Stop();
70+
71+
// copy back the array of sums from GPU and print
72+
cudaMemcpy(h_array, d_array, ARRAY_BYTES, cudaMemcpyDeviceToHost);
73+
// print_array(h_array, ARRAY_SIZE);
74+
printf("Time elapsed = %g ms\n", timer.Elapsed());
75+
76+
// free GPU memory allocation and exit
77+
cudaFree(d_array);
78+
return 0;
79+
}

LectureQuizzes/gputimer.h

+40
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,40 @@
1+
#ifndef __GPU_TIMER_H__
2+
#define __GPU_TIMER_H__
3+
4+
struct GpuTimer
5+
{
6+
cudaEvent_t start;
7+
cudaEvent_t stop;
8+
9+
GpuTimer()
10+
{
11+
cudaEventCreate(&start);
12+
cudaEventCreate(&stop);
13+
}
14+
15+
~GpuTimer()
16+
{
17+
cudaEventDestroy(start);
18+
cudaEventDestroy(stop);
19+
}
20+
21+
void Start()
22+
{
23+
cudaEventRecord(start, 0);
24+
}
25+
26+
void Stop()
27+
{
28+
cudaEventRecord(stop, 0);
29+
}
30+
31+
float Elapsed()
32+
{
33+
float elapsed;
34+
cudaEventSynchronize(stop);
35+
cudaEventElapsedTime(&elapsed, start, stop);
36+
return elapsed;
37+
}
38+
};
39+
40+
#endif /* __GPU_TIMER_H__ */

README.md

+23
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,23 @@
1+
These assignments require OpenCV as a prerequisite. The easiest way to install is probably with conda.
2+
3+
To install OpenCV in a conda environment.
4+
```
5+
conda create -n cs344 -y
6+
conda activate cs344
7+
conda install -y -c anaconda opencv
8+
```
9+
10+
To build
11+
```
12+
cd assignments
13+
mkdir build
14+
cd build
15+
cmake ..
16+
make
17+
```
18+
The binaries will then be contained within `assignments/bin`.
19+
20+
# HW1 Passing Instructions
21+
From the `HW1` directory.
22+
23+
Run `../bin/HW1 HW1/cinque_terre_small.jpg'`

assignments/README.md

+1
Original file line numberDiff line numberDiff line change
@@ -9,6 +9,7 @@ conda install -y -c anaconda opencv
99

1010
To build
1111
```
12+
cd assignments
1213
mkdir build
1314
cd build
1415
cmake ..

0 commit comments

Comments
 (0)