-
Notifications
You must be signed in to change notification settings - Fork 1
/
Copy pathchapter5_6_bitmap.cu
30 lines (30 loc) · 1.01 KB
/
chapter5_6_bitmap.cu
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
#include <cuda_runtime.h>
#include <device_launch_parameters.h>
#include <math.h>
#include "cpu_bitmap.h"
#define DIM 512
#define PI 3.14159
__global__ void kernel(unsigned char* ptr) {
int x = threadIdx.x + blockIdx.x * blockDim.x;
int y = threadIdx.y + blockIdx.y * blockDim.y;
int offset = x + y * blockDim.x * gridDim.x;
__shared__ float shared[16][16];
const float period = 128.0f;
shared[threadIdx.x][threadIdx.y] = 255 * (sinf(x * 2.0f * PI / period) + 1.0f) *
(sinf(y * 2.0f * PI / period) + 1.0f) / 4.0f;
__syncthreads();
ptr[offset * 4 + 0] = 0;
ptr[offset * 4 + 1] = shared[15- threadIdx.x][15 - threadIdx.y];
ptr[offset * 4 + 2] = 0;
ptr[offset * 4 + 3] = 255;
}
int main(void) {
CPUBitmap bitmap(DIM, DIM);
unsigned char* dev_bitmap;
cudaMalloc((void**)&dev_bitmap, bitmap.image_size());
dim3 grids(DIM / 16, DIM / 16);
dim3 threads(16, 16);
kernel << <grids, threads >> > (dev_bitmap);
cudaMemcpy(bitmap.get_ptr(), dev_bitmap, bitmap.image_size(),cudaMemcpyDeviceToHost);
bitmap.display_and_exit();
}