example-c

Author: roxacarv

Codes/produtointerno.c

@@ -0,0 +1,57 @@
+#include "../common/book.h"
+
+#define imin(a,b) (a<b?a:b)
+
+const int N = 33 * 1024;
+const int threadsPerBlock = 256;
+const int blocksPerGrid = imin(32, (N+threadsPerBlock-1) / threadsPerBlock);
+
+__global__ void dot(float *a, float *b, float *c)
+{
+	__shared__ float cache[threadsPerBlock];
+	int tid = threadIdx.x + blockIdx.x * blockDim.x;
+	int cacheIndex = threadIdx.x;
+	float temp = 0;
+	while (tid < N) {
+		temp += a[tid] * b[tid];
+		tid += blockDim.x * gridDim.x;
+	}
+
+	cache[cacheIndex] = temp;
+	int i = blockDim.x/2;
+	while (i != 0)
+	{
+		if (cacheIndex < i)
+			cache[cacheIndex] += cache[cacheIndex + 1];
+		__syncthreads();
+		i /= 2;
+	} 
+}
+
+int main(void)
+{
+	float *a, *b, c, *partial_c;
+	float *dev_a, *dev_b, *dev_partial_c;
+	a = new float[N];
+	b = new float[N];
+	partial_c = new float[blocksPerGrid];
+
+	cudaMalloc((void**)&dev_a, N*sizeof(float));
+	cudaMalloc((void**)&dev_b, N*sizeof(float));
+	cudaMalloc((void**)&dev_partial_c, blocksPerGrid*sizeof(float));
+
+	for (int i = 0; i < N; i++) {
+		a[i] = i;
+		b[i] = i * 2;
+	}
+
+	cudaMemcpy(dev_a, a, N*sizeof(float), cudaMemcpyHostToDevice);
+	cudaMemcpy(dev_b, b, N*sizeof(float), cudaMemcpyHostToDevice);
+
+	dot<<<blocksPerGrid, threadsPerBlock>>>(dev_a, dev_b, dev_partial_c);
+
+	cudaMemcpy(partial_c, dev_partial_c, blocksPerGrid*sizeof(float), cudaMemcpyDeviceToHost);
+	c = 0;
+	for (int i = 0; i < blocksPerGrid; i++)
+		c += partial_c[i];
+}

Codes/raytracingexample.c

@@ -0,0 +1,93 @@
+#include "cuda.h"
+#include "../common/book.h"
+#include "../common/cpu_bitmap.h"
+
+#define rnd(x) (x*rand() / RAND_MAX)
+#define SPHERES 20
+
+#define INF 2e10f
+
+struct Sphere {
+	float r, g, b;
+	float radius;
+	float x, y, z;
+	__device__ float hit(float ox, float oy, float *n) {
+		float dx = ox - x;
+		float dy = oy - y;
+		if(dx*dx + dy*dy < radius*radius) {
+			float dz = sqrtf(radius*radius - dx*dx - dy*dy);
+			*n = dz / sqrtf(radius*radius);
+			return dz + z;
+		}
+		return -INF;
+	}
+};
+
+__global__ void kernel(unsigned char *ptr)
+{
+	int x = threadIdx.x + blockIdx.x * blockDim.x;
+	int x = threadIdx.y + blockIdx.y * blockDim.y;
+	int offset = x + y * blockDim.x * gridDim.x;
+	float ox = (x-DIM/2);
+	float oy = (y-DIM/2);
+
+	float r = 0, g = 0, b = 0;
+	float maxz = -INF;
+	for(int i = 0; i < SPHERES; i++)
+	{
+		float n;
+		float t = s[i].hit(ox, oy, &n);
+		if(t > maxz) {
+			float fscale = n;
+			r = s[i].r * fscale;
+			r = s[i].g * fscale;
+			r = s[i].b * fscale;
+			maxz = t;
+		}
+	}
+
+	ptr[offset*4 + 0] = (int)(r*255);
+	ptr[offset*4 + 1] = (int)(g*255);
+	ptr[offset*4 + 2] = (int)(b*255);
+	ptr[offset*4 + 3] = 255;
+}
+
+int main(void) {
+	DataBlock data;
+	cudaEvent_t start, stop;
+	cudaEventCreate(&start);
+	cudaEventCreate(&stop);
+	cudaEventRecord(start, 0);
+
+	CPUBitmap bitmap(DIM, DIM, &data);
+	unsigned char *dev_bitmap;
+	Sphere *s;
+
+	cudaMalloc((void**)&dev_bitmap, bitmap.image_size());
+	cudaMalloc((void**)&s, sizeof(Sphere) * SPHERES);
+
+	Sphere *temp_s = (Sphere*)malloc(sizeof(Sphere) * SPHERES);
+	for(int i = 0; i < SPHERES; i++)
+	{
+		temp_s[i].r = rnd(1.0f);
+		temp_s[i].g = rnd(1.0f);
+		temp_s[i].b = rnd(1.0f);
+		temp_s[i].x = rnd(1000.0f) - 500;
+		temp_s[i].y = rnd(1000.0f) - 500;
+		temp_s[i].z = rnd(1000.0f) - 500;
+		temp_s[i].radius = rnd(100.0f) + 20;
+	}
+
+	cudaMemcpyToSymbol(s, temp_s, sizeof(Sphere) * SPHERES);
+	free(temp_s);
+
+	dim3 grids(DIM/16, DIM/16);
+	dim3 threads(16, 16);
+	kernel<<<grids, threads>>>(s, dev_bitmap);
+
+	cudaMemcpy(bitmap.get_ptr(), dev_bitmap, bitmap.image_size(), cudaMemcpyDeviceToHost);
+	bitmap.display_and_exit();
+
+	cudaFree(dev_bitmap);
+	cudaFree(s);
+}