In [0]:
!/usr/local/cuda/bin/nvcc --version
In [0]:
!pip install git+git://github.com/andreinechaev/nvcc4jupyter.git
In [0]:
%load_ext nvcc_plugin
In [0]:
%%cu
#include <iostream>
int main() {
std::cout << "Hello world!";
return 0;
}
Out[0]:
In [0]:
%%cu
#include<stdio.h>
#include<stdlib.h>
#include<iostream>
#define DIM 1000
struct cuComplex {
float r;
float i;
__device__ cuComplex( float a, float b ) : r(a), i(b) {}
__device__ float magnitude2( void ) {
return r * r + i * i;
}
__device__ cuComplex operator*(const cuComplex& a) {
return cuComplex(r*a.r - i*a.i, i*a.r + r*a.i);
}
__device__ cuComplex operator+(const cuComplex& a) {
return cuComplex(r+a.r, i+a.i);
}
};
__device__ int julia( int x, int y ) {
const float scale = 1.5;
float jx = scale * (float)(DIM/2 - x)/(DIM/2);
float jy = scale * (float)(DIM/2 - y)/(DIM/2);
cuComplex c(-0.8, 0.156);
cuComplex a(jx, jy);
int i = 0;
for (i=0; i<200; i++) {
a = a * a + c;
if (a.magnitude2() > 1000)
return 0;
}
return 1;
}
__global__ void kernel( int *ptr ) {
// map from threadIdx/BlockIdx to pixel position
int x = blockIdx.x;
int y = blockIdx.y;
int offset = x + y * gridDim.x;
// now calculate the value at that position
int juliaValue = julia( x, y );
ptr[offset] = juliaValue;
}
int main( void ) {
int bitmap[DIM*DIM];
int *dev_bitmap;
cudaMalloc( (void**)&dev_bitmap, DIM*DIM*sizeof(int) );
dim3 grid(DIM,DIM);
kernel<<<grid,1>>>( dev_bitmap );
cudaMemcpy( bitmap, dev_bitmap, DIM*DIM*sizeof(int), cudaMemcpyDeviceToHost );
cudaFree( dev_bitmap );
//for (int i=0; i<DIM*DIM; i++) printf("%d ", bitmap[i]);
FILE *myfile;
myfile=fopen("/tmp/test.bin","wb");
if (!myfile) {
printf("Unable to open file!");
return 1;
}
fwrite(bitmap, sizeof(int), DIM*DIM, myfile);
fclose(myfile);
std::cout << "Success!";
}
Out[0]:
In [0]:
import matplotlib.pyplot as plt
import numpy as np
bitmap = np.fromfile('/tmp/test.bin', dtype=np.int32)
fig = plt.figure(figsize=(8,8))
ax = fig.add_subplot(111)
ax.imshow(a.T, origin="lower")
Out[0]:
In [0]:
!rm /tmp/test.bin