Publish tokenizer_kernel kernel and performance card

df4a67f verified 3 days ago

7.89 kB

	#include <cuda_runtime.h>
	#include <stdio.h>
	#include <ctype.h>

	#define MAX_TOKENS 1024

	// Enhanced token types
	typedef enum {
	TOKEN_IDENTIFIER = 0,
	TOKEN_NUMBER = 1,
	TOKEN_OPERATOR = 2,
	TOKEN_KEYWORD = 3,
	TOKEN_STRING = 4,
	TOKEN_COMMENT = 5,
	TOKEN_PREPROCESSOR = 6,
	TOKEN_PUNCTUATION = 7
	} TokenType;

	// Add token metadata
	typedef struct {
	TokenType type;
	int start_pos;
	int end_pos;
	int length;
	int line;
	int column;
	char lexeme[256];
	unsigned int hash;
	} EnhancedTokenGPU;

	typedef struct {
	int type; // 0: identifier, 1: number, 2: operator
	int start_pos;
	int end_pos;
	int length;
	} TokenGPU;

	// Add shared memory optimization
	__shared__ char shared_input[1024];
	__shared__ int shared_token_count;

	// Enhanced tokenization kernel with better pattern matching
	__global__ void enhanced_tokenize_kernel(const char* input, size_t input_length,
	EnhancedTokenGPU* tokens, int* token_count,
	bool enable_comments, bool enable_preprocessing) {
	int idx = blockIdx.x * blockDim.x + threadIdx.x;
	if (idx >= input_length) return;

	// Load chunk into shared memory
	int local_idx = threadIdx.x;
	if (local_idx < 1024 && idx < input_length) {
	shared_input[local_idx] = input[idx];
	}
	__syncthreads();

	// Enhanced token detection with more patterns
	if (idx > 0 && (isalnum(shared_input[local_idx-1]) && isalnum(shared_input[local_idx]))) return;

	int tcount = atomicAdd(token_count, 0);
	if (tcount >= MAX_TOKENS) return;

	EnhancedTokenGPU token;
	token.start_pos = idx;
	token.hash = 0;

	// Calculate line and column
	int line = 1, column = 1;
	for (int i = 0; i < idx; i++) {
	if (input[i] == '\n') {
	line++;
	column = 1;
	} else {
	column++;
	}
	}
	token.line = line;
	token.column = column;

	// Enhanced pattern matching
	if (isalpha(shared_input[local_idx]) \|\| shared_input[local_idx] == '_') {
	// Handle identifiers and keywords
	int end = local_idx;
	while (end < 1024 && (isalnum(shared_input[end]) \|\| shared_input[end] == '_')) {
	token.hash = token.hash * 31 + shared_input[end];
	end++;
	}
	token.type = TOKEN_IDENTIFIER;
	token.end_pos = idx + (end - local_idx) - 1;
	token.length = end - local_idx;
	}
	// ... Add more token pattern matching ...

	// Store token if valid
	if (token.length > 0) {
	int new_count = atomicAdd(token_count, 1);
	if (new_count < MAX_TOKENS) {
	tokens[new_count] = token;
	}
	}
	}

	__global__ void tokenize_kernel(const char* input, size_t input_length, TokenGPU* tokens, int* token_count) {
	int idx = blockIdx.x * blockDim.x + threadIdx.x;
	if (idx >= input_length) return;

	// Skip if not at token boundary
	if (idx > 0 && (isalnum(input[idx-1]) && isalnum(input[idx]))) return;

	int tcount = *token_count;
	if (tcount >= MAX_TOKENS) return;

	if (isalpha(input[idx])) {
	int end = idx;
	while (end < input_length && isalnum(input[end])) end++;
	int new_count = atomicAdd(token_count, 1);
	if (new_count < MAX_TOKENS) {
	tokens[new_count].type = 0;
	tokens[new_count].start_pos = idx;
	tokens[new_count].end_pos = end - 1;
	tokens[new_count].length = end - idx;
	}
	} else if (isdigit(input[idx])) {
	int end = idx;
	while (end < input_length && isdigit(input[end])) end++;
	int new_count = atomicAdd(token_count, 1);
	if (new_count < MAX_TOKENS) {
	tokens[new_count].type = 1;
	tokens[new_count].start_pos = idx;
	tokens[new_count].end_pos = end - 1;
	tokens[new_count].length = end - idx;
	}
	} else if (input[idx] == '+' \|\| input[idx] == '-' \|\| input[idx] == '*' \|\| input[idx] == '/') {
	int new_count = atomicAdd(token_count, 1);
	if (new_count < MAX_TOKENS) {
	tokens[new_count].type = 2;
	tokens[new_count].start_pos = idx;
	tokens[new_count].end_pos = idx;
	tokens[new_count].length = 1;
	}
	}
	}

	__global__ void matrix_mult_kernel(float* a, float* b, float* c, int m, int n, int k) {
	int row = blockIdx.y * blockDim.y + threadIdx.y;
	int col = blockIdx.x * blockDim.x + threadIdx.x;
	if (row < m && col < n) {
	float sum = 0.0f;
	for (int i = 0; i < k; i++) {
	sum += a[row * k + i] * b[i * n + col];
	}
	c[row * n + col] = sum;
	}
	}

	// Add parallel matrix operations
	__global__ void enhanced_matrix_mult_kernel(float* a, float* b, float* c,
	int m, int n, int k,
	bool use_shared_memory) {
	// ... existing matrix multiplication code ...

	// Add shared memory optimization
	__shared__ float shared_a[16][16];
	__shared__ float shared_b[16][16];

	// ... implement block matrix multiplication ...
	}

	// Add new CUDA utilities
	void initialize_cuda_context(void) {
	cudaFree(0); // Force context initialization
	}

	void optimize_kernel_launch(dim3* blocks, dim3* threads, size_t shared_memory_size) {
	int device;
	cudaGetDevice(&device);
	cudaDeviceProp props;
	cudaGetDeviceProperties(&props, device);

	// Optimize launch configuration based on device properties
	// ... implementation ...
	}

	void cuda_tokenize(const char* input, TokenGPU* tokens, int* token_count) {
	size_t input_length = strlen(input);
	char* d_input;
	TokenGPU* d_tokens;
	int* d_token_count;

	cudaMalloc(&d_input, input_length + 1);
	cudaMalloc(&d_tokens, MAX_TOKENS * sizeof(TokenGPU));
	cudaMalloc(&d_token_count, sizeof(int));
	cudaMemcpy(d_input, input, input_length + 1, cudaMemcpyHostToDevice);
	cudaMemset(d_token_count, 0, sizeof(int));

	int threads = 256;
	int blocks = (input_length + threads - 1) / threads;
	tokenize_kernel<<<blocks, threads>>>(d_input, input_length, d_tokens, d_token_count);

	cudaMemcpy(token_count, d_token_count, sizeof(int), cudaMemcpyDeviceToHost);
	cudaMemcpy(tokens, d_tokens, token_count sizeof(TokenGPU), cudaMemcpyDeviceToHost);

	cudaFree(d_input);
	cudaFree(d_tokens);
	cudaFree(d_token_count);
	}

	void cuda_matrix_mult(float* a, float* b, float* c, int m, int n, int k) {
	float d_a, d_b, *d_c;
	cudaMalloc(&d_a, m * k * sizeof(float));
	cudaMalloc(&d_b, k * n * sizeof(float));
	cudaMalloc(&d_c, m * n * sizeof(float));
	cudaMemcpy(d_a, a, m * k * sizeof(float), cudaMemcpyHostToDevice);
	cudaMemcpy(d_b, b, k * n * sizeof(float), cudaMemcpyHostToDevice);

	dim3 threads(16, 16);
	dim3 blocks((n + threads.x - 1) / threads.x, (m + threads.y - 1) / threads.y);
	matrix_mult_kernel<<<blocks, threads>>>(d_a, d_b, d_c, m, n, k);

	cudaMemcpy(c, d_c, m * n * sizeof(float), cudaMemcpyDeviceToHost);
	cudaFree(d_a); cudaFree(d_b); cudaFree(d_c);
	}
	void print_tokens(TokenGPU* tokens, int token_count) {
	for (int i = 0; i < token_count; i++) {
	printf("Token %d: Type %d, Start %d, End %d, Length %d\n",
	i, tokens[i].type, tokens[i].start_pos, tokens[i].end_pos, tokens[i].length);
	}
	}
	int main() {
	const char* input = "int a = 5 + 3;";
	TokenGPU tokens[MAX_TOKENS];
	int token_count;

	cuda_tokenize(input, tokens, &token_count);
	print_tokens(tokens, token_count);

	float a[6] = {1, 2, 3, 4, 5, 6};
	float b[6] = {7, 8, 9, 10, 11, 12};
	float c[4] = {0};

	cuda_matrix_mult(a, b, c, 2, 3, 2);
	for (int i = 0; i < 4; i++) {
	printf("%f ", c[i]);
	}
	printf("\n");

	return 0;
	}
	// Compile with nvcc -o kernel kernel.cu